核同步

Question

我对Cuda编程很陌生，我正在实现经典的Floyd算法。该算法由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-loop中的内核必须连续执行，这就解释了为什么我在每次内核执行之后编写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将同步。