cuda:共享内存总是有用的吗?

8ljdwjyq  于 2023-05-06  发布在  其他
关注(0)|答案(1)|浏览(103)

当我阅读编程指南时,我得到了共享内存总是会提高性能的感觉,但似乎不是。我有两个功能:

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
先谢谢你了。

sg2wtvxw

sg2wtvxw1#

共享内存将始终提高性能

那不是真的。这取决于算法。如果您在内核中有一个完美的合并内存访问,并且您只访问全局内存一次,那么它可能没有帮助。但是如果你正在实现假设一个矩阵乘法,你需要保存部分和,那么它将是有用的。
如果您在内核中多次访问同一个内存位置,这也会很有帮助,因为共享内存的延迟比全局内存小100倍,因为它是片上内存。
当你分析内核是带宽有限的,那么它是一个很好的地方去思考是否有一个使用共享内存和提高性能的范围。检查占用率计算器以检查共享内存的使用是否会影响占用率也是更好的策略。

只有当我们想在设备中做很多事情时,共享内存才有帮助

部分是。当我们想在设备中做很多事情时,共享内存会有所帮助。
在上面的内核中,由于您在内核中多次访问全局内存,因此它应该有所帮助。如果您能提供完整的复制器来分析代码,这将是很有帮助的。此外,了解您正在运行的卡的详细信息也会有所帮助。

相关问题