c++ - 可以使用 __syncthreads() 合并单独的 CUDA 内核吗？

Question

假设我有这个玩具代码:

#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); 
}

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

(我的真实情况是 6 个连续的非平凡内核，它们有很多设置和拆卸开销)。

Answer 1

最简单、最普遍的答案是否定的。我只需要找到一个范式被打破的例子来支持这一点。让我们提醒自己:

1. __syncthreads() 是 block 级执行屏障，但不是设备范围的执行屏障。唯一定义的设备范围执行障碍是内核启动(假设我们正在讨论将内核发布到同一流中，以便顺序执行)。

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

假设我们有两个函数:

1. 反转 vector 的元素
2. 对 vector 元素求和

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

为了简单起见，我们假设我们只有(需要)两个线程 block 。对于第一步， block 0 将 vector 的左侧复制到右侧(顺序颠倒)， block 1 从右到左复制:

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

对于第二步，在经典的并行归约方式中， block 0 对输出 vector 的左侧元素求和， block 1 对右侧元素求和:

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

只要第一个函数在 kernel1 中发布，第二个函数在 kernel2 中发布，在 kernel1 之后进入同一个流，这一切都可以正常工作。对于每个内核， block 0 是否在 block 1 之前执行都没有关系，反之亦然。

如果我们组合这些操作以便我们有一个内核，并且 block 0 将 vector 的前半部分复制/反转到输出 vector 的后半部分，然后执行 __syncthreads() ，然后对输出 vector 的前半部分求和，事情很可能会中断。如果 block 0 在 block 1 之前执行，那么第一步会很好( vector 的复制/反转)但第二步将对尚未填充的输出数组一半进行操作，因为block 1 还没有开始执行。计算出的总和将是错误的。

在不尝试给出正式证据的情况下，我们可以看到在上述情况下，数据从一个 block 的“域”移动到另一个 block 的“域”，我们冒着破坏事物的风险，因为之前的设备范围同步(内核启动)对于正确性是必要的。但是，如果我们可以限制 block 的“域”，以便后续操作消耗的任何数据仅由该 block 中的先前操作产生，那么 __syncthreads() 可能足以让这个策略正确。 (前面的愚蠢示例可以很容易地重新设计以实现这一点，只需让 block 0 负责输出 vector 的前半部分，从而从 后半部分复制输入 vector ，反之亦然。)

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