CUDA 将共享内存用于 Hadamard 产品

Question

我试图了解我是否正确使用和理解具有一维输入和矩阵输入的简单内核的共享内存。我正在做哈达玛积，即一维输入应该与矩阵输入的每一行进行逐元素点积。

const int DSIZE = 4;
const int block_size = 2;

__global__ void hdot(const float *A, const float *B, float *C, int ds){

    int idx_A = threadIdx.x + blockIdx.x * blockDim.x;

    int idx_B = threadIdx.x + blockIdx.x * blockDim.x;
    int idy_B = threadIdx.y + blockIdx.y * blockDim.y;

    __shared__ float sharedA[block_size];
    __shared__ float sharedB[block_size][block_size];


    if ((idx_B < ds) && (idy_B < ds)){
        sharedA[threadIdx.x] = A[idx_A];
        sharedB[threadIdx.y][threadIdx.x] = B[idx_B * ds + idy_B];
        __syncthreads();

        C[idx_B * ds + idy_B] = sharedA[threadIdx.x] * sharedB[threadIdx.y][threadIdx.x];
    }
}

我认为我正确定义了我的共享内存，但我发现我的执行时间没有变化。不过数值结果是准确的。

非常感谢！

Answer 1

有关您要比较的内容、数据大小等的更多详细信息......将会有所帮助。这里有一些你可以看看的事情

共享内存是L1 Cache

在硬件上，L1缓存和共享内存是相同的元素，区别只是你自己管理它（共享内存）还是让运行时为你管理它（缓存）

在您的代码中，即使您分配给全局内存，数组也可能在第一次加载时被完全拉入缓存，然后它与共享内存一样快。如果您处理小数组，它们只会从全局内存加载一次。

共享内存库冲突

共享内存库冲突会导致共享内存的加载和存储被序列化，这可能会导致非常糟糕的性能。您可以在这里阅读更多相关信息https://github.com/Kobzol/hardware-effects-gpu/blob/master/bank-conflicts/README.md