CUDA：线程内的变量声明——有重叠吗？

Question

首先，我的问题措辞不正确；我认为最好使用 NVidia 的 CUDA C 编程指南中的示例进行询问。

在第3.2.3节（共享内存）中，下面给出了使用共享内存进行矩阵乘法的代码--希望我可以复制到这里。

__global__ void MatMulKernel(Matrix A, Matrix B, Matrix C)
{
// Block row and column
int blockRow = blockIdx.y;
int blockCol = blockIdx.x;

// Each thread block computes one sub-matrix Csub of C
Matrix Csub = GetSubMatrix(C, blockRow, blockCol);

// Each thread computes one element of Csub
// by accumulating results into Cvalue
float Cvalue = 0;

// Thread row and column within Csub
int row = threadIdx.y;
int col = threadIdx.x;

// Loop over all the sub-matrices of A and B that are
// required to compute Csub
// Multiply each pair of sub-matrices together
// and accumulate the results
for (int m = 0; m < (A.width / BLOCK_SIZE); ++m) {

    // Get sub-matrix Asub of A
    Matrix Asub = GetSubMatrix(A, blockRow, m);

    // Get sub-matrix Bsub of B
    Matrix Bsub = GetSubMatrix(B, m, blockCol);

    // Shared memory used to store Asub and Bsub respectively
    __shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
    __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];

    // Load Asub and Bsub from device memory to shared memory
    // Each thread loads one element of each sub-matrix
    As[row][col] = GetElement(Asub, row, col);
    Bs[row][col] = GetElement(Bsub, row, col);

    // Synchronize to make sure the sub-matrices are loaded
    // before starting the computation
    __syncthreads();

    // Multiply Asub and Bsub together
    for (int e = 0; e < BLOCK_SIZE; ++e)
        Cvalue += As[row][e] * Bs[e][col];

    // Synchronize to make sure that the preceding
    // computation is done before loading two new
    // sub-matrices of A and B in the next iteration
    __syncthreads();
}

// Write Csub to device memory
// Each thread writes one element
SetElement(Csub, row, col, Cvalue);
}

在第 7 行：Matrix Csub = GetSubMatrix(C, blockRow, blockCol)，每个线程都会执行该语句吗？这不会抵消使用共享内存来减少全局内存访问量的全部意义吗？我的印象是我在这里缺少一些基本的东西..

此外，当然还有更好的方式来表达这个问题。我就是不知道怎么办！

谢谢，

扎基尔

Answer 1

每个线程同时执行相同的指令（或处于空闲状态），所以每个线程都进入GetSubMatrix是的。每个线程需要一些项目。因此，如果有 N 线程和 3N 项要复制，则每个线程将复制 3 个。

例如，如果我正在复制一个向量，我可能会执行以下操作

float from* = ???;
float to*   = ???;
int   num   = ???;
int   thread = threadIdx.x + threadIdx.y*blockDim.x ...; // A linear index
int   num_threads = blockDim.x * blockDim.y * blockDim.z;
for(int i=threadIdx.x; i < num; i+= num_threads) {
     to[i] = from[i];
}

每个线程都参与一次复制一位。顺便说一句：如果您能够设法让所有线程复制一系列连续的元素，您将获得额外的复制速度。