可以使用__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); 
} 

就是一般的情况下，其采取相同的块和线程参数连续两个内核可与中间合并__syncthreads()是真的吗？ （我的真实情况是6个连续的非平凡内核，有很多设置和拆卸开销）。

Answer 1

最简单，最一般的答案是否定的。我只需要找到范例为了支持这个范例而打破的一个例子。让我们提醒自己：

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

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

比方说，我们有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 

对于第二个步骤，在经典平行减少的方式，框零和的输出向量的左手元件，和块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()。这后两种情况下可能有“embarassingly并行”问题，这表现出高度的独立性特点。