c++ - 内核同步

Question

我是 Cuda 编程的新手，我正在实现经典的 Floyd APSP 算法。该算法由 3 个嵌套循环组成，两个内循环中的所有代码都可以并行执行。

作为我代码的主要部分，这里是内核代码:

__global__ void dfloyd(double *dM, size_t k, size_t n)
{
    unsigned int x = threadIdx.x + blockIdx.x * blockDim.x;
    unsigned int y = threadIdx.y + blockIdx.y * blockDim.y;
    unsigned int index = y * n + x;
    double d;

    if (x < n && y < n)
    {
        d=dM[x+k*n] + dM[k+y*n];
        if (d<dM[index])
            dM[index]=d;
    }
}

这是主函数中启动内核的部分(为了便于阅读，我省略了错误处理代码):

double *dM;
cudaMalloc((void **)&dM, sizeof_M);
cudaMemcpy(dM, hM, sizeof_M, cudaMemcpyHostToDevice);

int dimx = 32;
int dimy = 32;
dim3 block(dimx, dimy);
dim3 grid((n + block.x - 1) / block.x, (n + block.y - 1) / block.y);

for (size_t k=0; k<n; k++)
{
    dfloyd<<<grid, block>>>(dM, k, n);
    cudaDeviceSynchronize();
}

cudaMemcpy(hM, dM, sizeof_M, cudaMemcpyDeviceToHost);

[为了便于理解，dM是指设备端存储的距离矩阵，hM是主机端存储的，n是指节点的数量。]

k 循环内的内核必须串行执行，这解释了为什么我在每次内核执行后编写 cudaDeviceSynchronize() 指令。 但是，我注意到将此同步指令放在循环的外部会导致相同的结果。

现在，我的问题。做下面两段代码

for (size_t k=0; k<n; k++)
{
    dfloyd<<<grid, block>>>(dM, k, n);
    cudaDeviceSynchronize();
}

和

for (size_t k=0; k<n; k++)
{
    dfloyd<<<grid, block>>>(dM, k, n);
}
cudaDeviceSynchronize();

是否等价？

Answer 1

它们不等价，但会给出相同的结果。第一个会让主机在每次内核调用后等待，直到内核返回，而另一个会让主机只等待一次。 也许令人困惑的部分是它为什么有效；在 CUDA 中，保证连续执行同一流(在您的情况下为默认流)上的两个连续内核调用。

性能方面，建议使用第二个版本，因为与主机同步会增加开销。

编辑:在那种特定情况下，您甚至不需要调用 cudaDeviceSynchronize()，因为 cudaMemcpy 会同步。