库达 : Is shared memory always helpful?

Question

看programming guide的时候，感觉共享内存总会提高性能，但好像不是。我有两个功能:

const int Ntimes=1;

__global__ void testgl(float *A, float *C, int numElements){

    int ti = threadIdx.x;
    int b0 = blockDim.x*blockIdx.x;

    if (b0+ti < numElements){
        for(int i=0;i<Ntimes;i++){
            A[b0+ti]=A[b0+ti]*A[b0+ti]*10-2*A[b0+ti]+1;
        }
        C[b0+ti] = A[b0+ti]*A[b0+ti];
    }
}


__global__ void testsh(float *A, float *C, int numElements){

    int ti = threadIdx.x;
    int b0 = blockDim.x*blockIdx.x;

    __shared__ float a[1024];

    if (b0+ti < numElements){
        a[ti]=A[b0+ti];
    }

    __syncthreads();

    if (b0+ti < numElements){
        for(int i=0;i<Ntimes;i++){
            a[ti]=a[ti]*a[ti]*10-2*a[ti]+1;
        }
        C[b0+ti] = a[ti]*a[ti];
    }
}

int main(void){

    int numElements = 500000;
    size_t size = numElements * sizeof(float);

    // Allocate the host input
    float *h_A = (float *)malloc(size);
    float *h_B = (float *)malloc(size);

    // Allocate the host output
    float *h_C = (float *)malloc(size);
    float *h_D = (float *)malloc(size);


    // Initialize the host input
    for (int i = 0; i < numElements; i++){
        h_A[i] = rand()/(float)RAND_MAX;
        h_B[i] = h_A[i];
    }

    // Allocate the device input
    float *d_A = NULL; cudaMalloc((void **)&d_A, size);
    float *d_B = NULL; cudaMalloc((void **)&d_B, size);
    float *d_C = NULL; cudaMalloc((void **)&d_C, size);
    float *d_D = NULL; cudaMalloc((void **)&d_D, size);


    //Copy to Device
    cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);  
    cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);


    // Launch the Vector Add CUDA Kernel
    int threadsPerBlock = 1024;
    int blocksPerGrid =(numElements + threadsPerBlock - 1) / threadsPerBlock;

    testgl<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_C, numElements);

    testsh<<<blocksPerGrid, threadsPerBlock>>>(d_B, d_D, numElements);

    // Copy the device resultto the host 
    cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);
    cudaMemcpy(h_D, d_D, size, cudaMemcpyDeviceToHost);


    // Free device global memory
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);
    cudaFree(d_D);

    // Free host memory
    free(h_A);
    free(h_B);
    free(h_C);
    free(h_D);

    // Reset the device and exit
    cudaDeviceReset();

    return 0;
}

如果Ntimes设置为1，testgl耗时49us，testsh耗时97us。如果Ntimes设置为100，testgl耗时9.7ms，testsh耗时8.9ms。

不知道为什么要长100多倍。

所以看起来共享内存只有在我们想在设备上做很多事情时才有用，对吗？

这里使用的显卡是GTX680。

提前致谢。

Answer 1

共享内存总能提高性能

那不是真的。这取决于算法。如果您在内核中有一个完美合并的内存访问，并且您只访问一次全局内存，它可能无济于事。但是如果你正在实现一个矩阵乘法，你需要保留部分和，那么它将很有用。

如果您在内核中多次访问同一内存位置，这也会有所帮助，在这种情况下它会有所帮助，因为共享内存延迟比全局内存少 100 倍，因为它是片上内存。

当您分析内核带宽受限时，这是一个思考是否有使用共享内存和提高性能的好地方。检查占用计算器以检查共享内存的使用是否会影响占用也是更好的策略。

只有当我们想在设备上做很多事情时，共享内存才有用 ?

部分是的。当我们想在设备中做很多事情时，共享内存会有所帮助。

在上述内核的情况下，当您在内核中多次访问全局内存时，它应该有所帮助。如果能提供完整的复现器来分析代码会很有帮助。了解您正在运行的卡的详细信息也将很有帮助。