Leistung der statischen gegenüber der dynamischen CUDA Shared-Memory-Allokation

Question

Ich habe 2 Kernel, die genau dasselbe tun. Einer von ihnen allokiert den gemeinsam genutzten Speicher statisch, während der andere den Speicher zur Laufzeit dynamisch allokiert. Ich verwende den gemeinsam genutzten Speicher als 2D-Array. Für die dynamische Zuweisung habe ich ein Makro, das den Speicherort berechnet. Die Ergebnisse, die von den 2 Kernels erzeugt werden, sind jedoch genau gleich. Die Zeitmessungsergebnisse, die ich von beiden Kernels erhalten habe, unterscheiden sich um das 3-Fache! Die statische Speicherzuordnung ist viel schneller. Es tut mir leid, dass ich meinen Code nicht posten kann. Kann jemand eine Begründung dafür geben?

Answer 1

Ich habe keine Beweise dafür, dass die statische gemeinsame Speicherzuweisung schneller ist als die dynamische gemeinsame Speicherzuweisung. Wie oben in den Kommentaren bestätigt wurde, wäre es unmöglich, Ihre Frage ohne Replikator zu beantworten. Zumindest im Fall des unten stehenden Codes sind die Zeitmessungen für denselben Kernel, wenn er mit statischen oder dynamischen gemeinsamen Speicherzuweisungen ausgeführt wird, genau gleich:

#include 
#include 

#define BLOCK_SIZE 512

/********************/
/* CUDA ERROR CHECK */
/********************/
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, char *file, int line, bool abort=true)
{
    if (code != cudaSuccess) 
    {
        fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
        if (abort) exit(code);
    }
}

/***********************************/
/* SHARED MEMORY STATIC ALLOCATION */
/***********************************/
__global__ void kernel_static_memory_allocation(int *d_inout, int N)
{
    __shared__ int s[BLOCK_SIZE];

    const int tid   = threadIdx.x;
    const int i     = blockIdx.x * blockDim.x + threadIdx.x;

    if (i < N) {

        s[tid] = d_inout[i];
        __syncthreads();

        s[tid] = s[tid] * s[tid];
        __syncthreads();

        d_inout[i] = s[tid];
    }
}

/************************************/
/* SHARED MEMORY DYNAMIC ALLOCATION */
/************************************/
__global__ void kernel_dynamic_memory_allocation(int *d_inout, int N)
{
    extern __shared__ int s[];

    const int tid   = threadIdx.x;
    const int i     = blockIdx.x * blockDim.x + threadIdx.x;

    if (i < N) {

        s[tid] = d_inout[i];
        __syncthreads();

        s[tid] = s[tid] * s[tid];
        __syncthreads();

        d_inout[i] = s[tid];
    }
}

/********/
/* MAIN */
/********/
int main(void)
{
    int N = 1000000;

    int* a = (int*)malloc(N*sizeof(int));

    for (int i = 0; i < N; i++) { a[i] = i; }

    int *d_inout; gpuErrchk(cudaMalloc(&d_inout, N * sizeof(int))); 

    int n_blocks = N/BLOCK_SIZE + (N%BLOCK_SIZE == 0 ? 0:1);

    gpuErrchk(cudaMemcpy(d_inout, a, N*sizeof(int), cudaMemcpyHostToDevice));

    float time;
    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);
    cudaEventRecord(start, 0);  
    kernel_static_memory_allocation<<>>(d_inout, N);
    gpuErrchk(cudaPeekAtLastError());
    gpuErrchk(cudaDeviceSynchronize());
    cudaEventRecord(stop, 0);
    cudaEventSynchronize(stop);
    cudaEventElapsedTime(&time, start, stop);
    printf("Statische Allokation - vergangene Zeit:  %3.3f ms \n", time);

    cudaEventRecord(start, 0);  
    kernel_dynamic_memory_allocation<<>>(d_inout, N);
    gpuErrchk(cudaPeekAtLastError());
    gpuErrchk(cudaDeviceSynchronize());
    cudaEventRecord(stop, 0);
    cudaEventSynchronize(stop);
    cudaEventElapsedTime(&time, start, stop);
    printf("Dynamische Allokation - vergangene Zeit:  %3.3f ms \n", time);

}

Der mögliche Grund dafür ist, dass die disassemblierten Codes für die beiden Kerne vollkommen identisch sind und sich nicht ändern, selbst wenn int N = 1000000; durch int N = rand(); ersetzt wird.