执行期间可以中断执行块吗？

我担心我的一个cuda内核中潜在的竞争条件。我正在研究Barnes Hunt Tree算法的N-Body模拟器。此内核的目的是计算树的每个分支的总质量和质量中心。我想在容器数组上以相反顺序“迭代”，因为那些分配的最后一个最不可能依赖于其他子容器，数组中的第一个容器也可能依赖于后来的容器。执行期间可以中断执行块吗？

我正在使用原子计数器来跟踪哪些块首先启动，并且第一个块处理前几个容器等等。我担心的是，一个区块的执行可能暂时中止，直到其他区块完成或类似的事情？这是一个问题，因为说第一个街区开始，然后出于任何原因而屈服于其他街区。在这种情况下，如果其他依赖于第一个块执行的计算，他们将无限循环。

__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(); 
}

来源

2016-02-20 chasep255

我真的不是100％肯定它是什么，你真的想知道这里，因为它似乎连接到代码甩了你在内，但没有块调度，执行，或在CUDA运行时模型退休秩序的保证。任何依赖预定执行顺序的代码都依赖于未定义的行为。 – talonmies

没有像CUDA块间同步那样的东西。尽管如此，人们已经做了研究，例如：Shucai Xiao and Wu-chun Feng，块间GPU通信通过快速的栅栏同步

你的情况，人们可以简单地做或者几个内核，每一个块调用或者如果你是冒险精神在全局内存中自制（缓慢）阻塞原子操作来同步。

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

来源

2016-02-20 15:21:01 Ax3l

总是值得指出的是，在符合CUDA编程模型的同时，在块之间没有任何依赖关系。代码现在可以工作，但没有保证它明天会起作用。 – Jez

为什么我得到一个倒票？有趣的是，我想我已经很清楚了，在CUDA中没有像句子一中的块间同步那样的东西。第2段是启动多个内核而不是一个的调试可能性，第3段是他正在寻找的实际工具，但似乎并不知道它存在。 – Ax3l

执行期间可以中断执行块吗？

回答

相关问题