c++ - cuda执行 block 可以在执行过程中被打断吗？

Question

我担心我的一个 cuda 内核中存在潜在的竞争条件。我正在为 Barnes Hunt Tree 算法开发一个 N 体模拟器。该内核的目的是计算树的每个分支的总质量和质心。我想在容器数组上以相反的顺序“迭代”，因为最后分配的容器最不可能依赖于其他子容器，数组中的第一个容器也可能依赖于后面的容器。

我正在使用一个原子计数器来跟踪哪些 block 先启动，第一个 block 处理前几个容器，依此类推。我担心的是一个 block 的执行是否可以暂时暂停，直到其他 block 完成或类似的事情？这是一个问题，因为说第一个 block 开始，然后出于任何原因为其他 block 让步。在这种情况下，如果其他人依赖于第一个 block 执行的计算，它们将无限循环。

__global__ void compute_mass_centers_kernel()
{
    int blockNum = atomicAdd(&dev::block_number, 1);
    int cindex = dev::ncontainers - blockNum * blockDim.x - 1 - threadIdx.x;
    if(cindex < 0)
        return;

    Container& c = dev::containers[cindex];
    int missing_ptrs[8];
    int missing = 0;

    float total_mass = 0.0f;
    double3 com = {0}; 
    for(int i = 0; i < 8; i++)
    {
        if(c[i] > 1)
        {
            Object& o = objat(c[i]);
            total_mass += o.m;
            com.x += (double)o.p.x * o.m;
            com.y += (double)o.p.y * o.m;
            com.z += (double)o.p.z * o.m;
        }
        else if(c[i] < 1)
        {
            missing_ptrs[missing++] = c[i];
        }
    }

    while(missing)
    {
        for(int i = 0; i < missing; i++)
        {
            Container& c2 = ctrat(missing_ptrs[i]);
            if(c2.total_mass >= 0.0f)
            {
                total_mass += c2.total_mass;
                com.x += (double)c2.center_of_mass.x * c2.total_mass;
                com.y += (double)c2.center_of_mass.y * c2.total_mass;
                com.z += (double)c2.center_of_mass.z * c2.total_mass;
                missing_ptrs[i--] = missing_ptrs[--missing];
            }
        }
    }

    c.center_of_mass.x = com.x / total_mass;
    c.center_of_mass.y = com.y / total_mass;
    c.center_of_mass.z = com.z / total_mass;
    c.total_mass = total_mass;
}

void compute_mass_centers()
{
    int threads, blocks;
    cudaOccupancyMaxPotentialBlockSize(&blocks, &threads, compute_mass_centers_kernel, 0, 0);
    cucheck();

    int ncontainers;
    cudaMemcpyFromSymbol(&ncontainers, dev::ncontainers, sizeof(int), 0, cudaMemcpyDeviceToHost);
    cucheck();

    blocks = (ncontainers + (threads - 1)) / threads;

    cudaMemcpyToSymbol(dev::block_number, &ZERO, sizeof(int), 0, cudaMemcpyHostToDevice);
    cucheck();

    compute_mass_centers_kernel<<< blocks, threads >>>();
    cucheck();
}

Answer 1

没有像 CUDA block 间同步这样的东西。尽管如此，人们已经对此进行了研究，例如:Shucai Xiao and Wu-chun Feng , block 间 GPU 通信 通过快速屏障同步

在您的情况下，可以简单地对每个 block 执行多个内核调用，或者如果您喜欢冒险，则可以在全局内存中进行一个自制的(缓慢的)阻塞原子操作以进行同步。

对于您的潜在问题，最好的解决方案可能是使用 cuda-memcheck 检查您的代码.