Могут ли исполняемые блоки cuda прерываться во время выполнения?

Question

Меня беспокоит потенциальное состояние гонки в одном из моих ядер куды. Я работаю над симулятором N-Body для алгоритма Barnes Hunt Tree. Целью этого ядра является вычисление полной массы и центров масс для каждой ветви дерева. Я хочу «итерации» в обратном порядке в массиве контейнеров, потому что те, которые были назначены последними, в меньшей степени зависят от других дочерних контейнеров, также первые контейнеры в массиве, вероятно, зависят от более поздних контейнеров.

Я использую атомный счетчик, чтобы отслеживать, какие блоки начинают сначала, а первый блок обрабатывает первые несколько контейнеров и так далее. Я беспокоюсь о том, может ли выполнение блока временно приостановлено до тех пор, пока другие блоки не закончатся или что-то в этом роде? Это проблема, так как, скажем, первый блок начинается, а затем дает для других по какой-либо причине. В этом случае, если другие зависят от вычислений, выполненных первым блоком, они будут циклически индексироваться.

__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. Тем не менее, люди сделали исследования, что, например: Shucai Xiao and Wu-chun Feng, Inter-Block GPU связи посредством синхронизации Fast Барьерный

В вашем случае можно было бы просто сделать что-либо несколько вызовы ядра с каждым одного блока или, если вы дерзкими чтобы выполнить синхронизацию с помощью самопроизвольной (медленной) блокирующей атомной операции в глобальной памяти.

Для вашей основной проблемы лучшим решением может быть проверка кода с помощью cuda-memcheck.