可以使用__syncthreads(）合并单独的CUDA内核

Can separate CUDA kernels be merged using __syncthreads()?

本文关键字：单独 CUDA 合并 内核 syncthreads 可以使      更新时间：2023-10-16

假设我有这个玩具代码：

#define N (1024*1024)
#define M (1000000)
__global__ void cudakernel1(float *buf)
{
   int i = threadIdx.x + blockIdx.x * blockDim.x;
   buf[i] = 1.0f * i / N;
   for(int j = 0; j < M; j++)
      buf[i] *= buf[i];
}
__global__ void cudakernel2(float *buf)
{
   int i = threadIdx.x + blockIdx.x * blockDim.x;
   for(int j = 0; j < M; j++)
      buf[i] += buf[i];
}
int main()
{
   float data[N];
   float *d_data;
   cudaMalloc(&d_data, N * sizeof(float));
   cudakernel1<<<N/256, 256>>>(d_data);
   cudakernel2<<<N/256, 256>>>(d_data);
   cudaMemcpy(data, d_data, N * sizeof(float), cudaMemcpyDeviceToHost);
   cudaFree(d_data); 
}

我可以这样合并两个内核吗：

#define N (1024*1024)
#define M (1000000)
__global__ void cudakernel1_plus_2(float *buf)
{
   int i = threadIdx.x + blockIdx.x * blockDim.x;
   buf[i] = 1.0f * i / N;
   for(int j = 0; j < M; j++)
      buf[i] *= buf[i];
   __syncthreads();
   for(int j = 0; j < M; j++)
      buf[i] += buf[i];
}
int main()
{
   float data[N];
   float *d_data;
   cudaMalloc(&d_data, N * sizeof(float));
   cudakernel1_plus_2<<<N/256, 256>>>(d_data);
   cudaMemcpy(data, d_data, N * sizeof(float), cudaMemcpyDeviceToHost);
   cudaFree(d_data); 
}

采用相同块和线程参数的两个连续内核可以与中间__syncthreads()合并的一般情况是真的吗？

（我的真实案例是6个连续的非琐碎内核，它们有大量的设置和拆卸开销）。

最简单、最普遍的答案是否定的。我只需要找到一个范例来支持它。让我们提醒自己：

  1. __syncthreads()是块级执行屏障，但不是设备范围的执行屏障。唯一定义的设备范围执行障碍是内核启动（假设我们谈论的是将内核发布到同一个流中，以进行顺序执行）。

  2. 特定内核启动的线程块可以以任何顺序执行。

假设我们有两个功能：

  1. 反转矢量的元素
  2. 求和矢量元素

让我们假设向量反转不是一个就地操作（输出与输入不同），并且每个线程块处理向量的块大小的块，读取元素并存储到输出向量中的适当位置。

为了简单起见，我们假设我们只需要两个线程块。对于第一步，块0将矢量的左手边复制到右手边（颠倒顺序），块1从右向左复制：

1 2 3 4 5 6 7 8
|blk 0 |blk 1  |
      | /
       X
      /| 
     v |  v
8 7 6 5 4 3 2 1

对于第二步，以经典的并行归约方式，块0对输出向量的左手元素求和，块1对右手元素求和：

8 7 6 5 4 3 2 1
    /     /
  blk0    blk1
   26      10

只要第一个函数在kernel1中发布，第二个函数在kernel2中发布，进入kernel1之后的同一流，这一切都会起作用。对于每个内核，块0是否在块1之前执行并不重要，反之亦然。

如果我们将这些操作组合起来，使我们有一个单独的内核，并且块0将向量的前半部分复制/反转为输出向量的后半部分，然后执行__syncthreads()，然后对输出向量的前一半求和，那么事情很可能会破裂。如果块0在块1之前执行，则第一步将很好（复制/反转矢量），但第二步将对尚未填充的输出数组的一半进行操作，因为块1尚未开始执行。计算出的总和是错误的。

在不试图给出正式证明的情况下，我们可以看到，在上述情况下，数据从一个块的"域"移动到另一个块"域"，我们面临着破坏事物的风险，因为之前的设备范围同步（内核启动）是正确性所必需的。然而，如果我们可以限制一个块的"域"，使后续操作消耗的任何数据仅由该块中的前一个操作产生，则__syncthreads()可能足以正确地允许该策略。（前面的愚蠢例子可以很容易地进行重新设计，只需让块0负责输出向量的前半部分，从而从输入向量的后半部分复制，反之亦然。）

最后，如果我们将数据范围限制为单个线程，那么我们甚至可以在不使用__syncthreads()的情况下进行这样的组合。后两种情况可能具有"令人尴尬的平行"问题的特征，表现出高度的独立性。

相关文章：