静态与动态 CUDA 共享内存分配的性能答案

【问题标题】：Performance of static versus dynamic CUDA shared memory allocation静态与动态 CUDA 共享内存分配的性能
【发布时间】：2011-07-07 18:45:11
【问题描述】：

我有2 内核做同样的事情。其中一个静态分配共享内存，而另一个在运行时动态分配内存。我将共享内存用作二维数组。所以对于动态分配，我有一个计算内存位置的宏。现在，2 内核生成的结果完全相同。然而，我从两个内核得到的计时结果是3 倍！静态内存分配要快得多。很抱歉，我无法发布任何代码。有人可以为此提供理由吗？

【问题讨论】：

您是否调查过 NVCC 使用 -ptx 选项生成的 PTX 代码？如果输出代码有明显差异，这可能有助于解释为什么一个内核更快。
可能与编译器优化有关。您可以尝试使用 -O0 选项编译这两个代码吗？
你确定你的动态大小计算和静态数组的大小一样吗？如果它更大（可能是错误的），那么您可能会减少内核的占用率。顺便说一句，要获得任何实际答案，您需要提供实际详细信息——一个代码示例。
allocating shared memory 的可能重复项

标签： cuda

【解决方案1】：

我没有证据表明静态共享内存分配比动态共享内存分配更快。正如上面的 cmets 所证明的那样，如果没有复制器，就不可能回答您的问题。至少在以下代码的情况下，同一个内核的时序，在使用静态或动态共享内存分配运行时，是完全相同的：

#include <cuda.h>
#include <stdio.h>

#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<<<n_blocks,BLOCK_SIZE>>>(d_inout, N);
    gpuErrchk(cudaPeekAtLastError());
    gpuErrchk(cudaDeviceSynchronize());
    cudaEventRecord(stop, 0);
    cudaEventSynchronize(stop);
    cudaEventElapsedTime(&time, start, stop);
    printf("Static allocation - elapsed time:  %3.3f ms \n", time);

    cudaEventRecord(start, 0);  
    kernel_dynamic_memory_allocation<<<n_blocks,BLOCK_SIZE,BLOCK_SIZE*sizeof(int)>>>(d_inout, N);
    gpuErrchk(cudaPeekAtLastError());
    gpuErrchk(cudaDeviceSynchronize());
    cudaEventRecord(stop, 0);
    cudaEventSynchronize(stop);
    cudaEventElapsedTime(&time, start, stop);
    printf("Dynamic allocation - elapsed time:  %3.3f ms \n", time);

}

可能的原因是两个内核的反汇编代码完全相同，即使将 int N = 1000000; 替换为 int N = rand(); 也不会改变。

【讨论】：