CUDA syncthreads() 阻塞其他线程

Question

#define TS 32
int num_devices = 0;

__global__ void shared_kernel(float* A, float* B, float* C, int M, int N, int K) {

   int global_col = blockDim.x * blockIdx.x + threadIdx.x;
   int global_row = blockDim.y * blockIdx.y + threadIdx.y;
   int local_col  = threadIdx.x;
   int local_row  = threadIdx.y;
   if (global_row >= M || global_col >= N) return;


   __shared__ float Asub[TS][TS];
   __shared__ float Bsub[TS][TS];
 
   const int num_tiles = K / TS;

   float acc = 0;
 
   for(int t = 0; t < num_tiles; t++){
       const int t_row = TS * t + local_row;
       const int t_col = TS * t + local_col;
       Asub[local_row][local_col] = A[global_row * K + t_col];
       Bsub[local_row][local_col] = B[t_row * N + global_col];


       __syncthreads();
       printf("[DEBUG] first sync threads, global_row: %d, global_col: %d\n", global_row, global_col);
 

       for (int k = 0; k < K; ++k) {
         acc += Asub[local_row][k] * Bsub[k][local_col];
       }
 

       __syncthreads();
       printf("[DEBUG] second sync threads, global_row: %d, global_col: %d\n", global_row, global_col);
   }

   C[global_row * N + global_col] = acc;
}

static float *a_d, *b_d, *c_d;

void mat_mul(float *A, float *B, float *C, int M, int N, int K) {
 cudaMemcpy(a_d, A, M * K * sizeof(float), cudaMemcpyHostToDevice);
 cudaMemcpy(b_d, B, K * N * sizeof(float), cudaMemcpyHostToDevice);

 dim3 blockDim(TS, TS);
 dim3 gridDim(M/TS, N/TS);
 shared_kernel<<<gridDim, blockDim>>>(a_d, b_d, c_d, M, N, K); 

 cudaMemcpy(C, c_d, M * N * sizeof(float), cudaMemcpyDeviceToHost);

 cudaDeviceSynchronize();
}

void mat_mul_init(float *A, float *B, float *C, int M, int N, int K) {

 cudaGetDeviceCount(&num_devices); 
 cudaSetDevice(0);

 cudaMalloc(&a_d, M * K * sizeof(float));
 cudaMalloc(&b_d, K * N * sizeof(float));
 cudaMalloc(&c_d, M * N * sizeof(float));
}

上面的例子是共享内存的矩阵乘法。我运行上面的内核有 dim3 blockDim(TS, TS) 和 dim3 gridDim(M/TS, N/TS) 和 M, N, K = 128.

我在启动内核后检查 float * C 的值为零。此外，我发现在第一个 __syncthreads() 之后只有少数 global_row 被打印（从 37 到 81），并且在第二个 __syncthreads() 之后没有 printf DEBUG 消息。

我怀疑是 __syncthreads() 引起了问题，但我不知道如何解决。我的代码和其他网站的其他矩阵乘法代码几乎一样。

你能给我一些解决方法的提示吗？

Answer 1

任何时候您在使用 CUDA 代码时遇到问题，我建议您使用 proper CUDA error checking 和运行您的代码以及 compute-sanitizer 或 cuda-memcheck。对于这种类型的分析，如果您不使用内核printf.

会更容易

如果你这样做，你会看到这样的输出：

========= Invalid __shared__ read of size 4
=========     at 0x000002f0 in shared_kernel(float*, float*, float*, int, int, int)
=========     by thread (0,2,0) in block (0,1,0)
=========     Address 0x00002000 is out of bounds
=========     Saved host backtrace up to driver entry point at kernel launch time
...  (and more output)

因此，我们可以看出您的内核正在进行无效的 __shared__ 读取操作。您的内核中发生了什么？您可以使用方法 here 来识别特定的代码行。然而这是一个相当简单的内核，只有一行是从共享内存读取的，它在这里：

   for (int k = 0; k < K; ++k) {
     acc += Asub[local_row][k] * Bsub[k][local_col];  // shared reads here

快速检查会发现，如果您让此循环遍历 K=128 范围，那么您将在此处索引出界：

   for (int k = 0; k < K; ++k) {
     acc += Asub[local_row][k] * Bsub[k][local_col];
                            ^         ^

当 k 大于 31 时，因为这会超出您的共享数组维度：

#define TS 32

__shared__ float Asub[TS][TS];
__shared__ float Bsub[TS][TS];

我不会费心为你写一个固定的kernel/code，因为正如你已经指出的，这个主题在许多其他地方都有涉及，并且已经在[中提供了一个规范的例子=24=].

FWIW，如果我将您的 for 循环更改为：

   for (int k = 0; k < TS; ++k) {

然后运行-时间错误对我来说消失了。 cuda-memcheck 报告没有错误。

CUDA syncthreads() 阻塞其他线程

CUDA syncthreads() block other threads

cuda