有没有办法在 cuda 的条件 while 循环中使用同步线程？答案

【问题标题】：Are there ways to use syncthreads in a conditional while loop in cuda?有没有办法在 cuda 的条件 while 循环中使用同步线程？
【发布时间】：2019-08-23 22:32:04
【问题描述】：

假设每个线程块存在 64 个线程。

每个线程运行以下代码：

int tid = threadIdx.x;

while (tid < 96) {
  // write data to shared memory
  __syncthreads();
  // read data from shared memory
  tid += 64;
}

在这种情况下，while 语句的第二次运行导致最后 32 个线程没有到达__syncthreads()，我认为这会导致未定义的行为。

或者，自从最后 32 个线程到达代码末尾后，__syncthreads() 是否正常工作？

我目前不知道如何解决这类问题。

我必须更改算法吗？

【问题讨论】：

"__syncthreads() 允许在条件代码中使用，但前提是条件在整个线程块中的计算结果相同，否则代码执行可能会挂起或产生意外的副作用。" (CUDA Programming Guide) 但是如果其他线程早点返回，TBH 感觉不像 UB。
感觉可能不像，但肯定是未定义的行为。

标签： cuda

【解决方案1】：

让所有线程循环到相同的（最大值）值，并将循环体（不包括__syncthreads()）封装到另一个有条件地检查循环计数器与实际结束值的情况：

#define round_up(x, m) ((x) + (m) - 1 - ((x) + (m) - 1) % (m))

int tid = threadIdx.x;
int maxtid = round_up(96, blockDim.x) - 1;

while (tid <= maxtid) { 
  if (tid < 96) {
    // write data to shared memory
  }
  __syncthreads();
  if (tid < 96) {
    // read data from shared memory
  }
  tid += 64;
}

【讨论】：

但这不会也可能导致线程退出外循环的死锁吗？
@talonmies 你还看到潜在的死锁吗？我想知道我是否在这里错过了什么。

【解决方案2】：

__syncthreads()

如果不是所有块线程都进入它，则会导致未定义的行为或死锁。

通过单个线程块扫描大于块大小的空间，

copyLimit = 8192; // assuming 8k is going to be copied
w = blockDim.x; // this will be running inside copyLimit, nLoop times to fill all of it, but masked by some if-else to not overflow it
nLoop = copyLimit/w + 1;
for(int i=0;i<nLoop;i++)
{
     // load from global to shared
     if(threadIdx.x+i*w<copyLimit)
        doLoad();
}
__syncthreads(); // only once! A loading doesn't need sync with a loading.

for(int i=0;i<nLoop;i++)
{

     if(threadIdx.x+i*w<copyLimit)
     {
        // compute, assuming its just embarrassingly parallel
     }
}
// can also compute here too depending on compute job, to use all pipelines

for(int i=0;i<nLoop;i++)
{
     // save to global from shared
     if(threadIdx.x+i*w<copyLimit)
        doStore();
}
__syncthreads(); // only once!
// so that you can use stored values by other threads
doSomeWork(sharedArray);

如果每个线程的循环周期数未知（例如处理不平衡的树），则有一个共享的活动计数器。

active=1; // start working
while(active>0)
{
     // work
     if(!isFinished())
         doWork();         // sets isFinished() if it has no other job
      // any syncthreads or syncwarp whatever you need to sync
      __syncthreads(); // is not undefined behavior



     // when thread finishes its job, its not active
     if(isFinished())
         activeList[threadIdx.x] = 0;

     // reduction in a shared array, to find total number of active threads
     // and broadcast it to all threads 
     active=reduceActiveThreads(); // includes its own syncthreads    
}
// all block threads exit here together, as soon as last thread completes its job

如果有 Volta+ 架构，您也可以尝试它的 warp 版本，以减少循环中空闲线程的丢失周期。（依靠它的独立线程调度）。即使没有 Volta，warp 减少也可以比共享数组减少更快。

如果所有线程的周期数相同但在编译时未知，那么找到所有线程的最大值就足够了。然后使用它为块的所有线程循环该次数不会产生未定义的行为。

int nLoop = findNumCycles(threadIdx.x, someParameters);
nLoop = reduceN(nLoop, threadIdx.x); // max(of all nLoop values)
for(int i=0;i<nLoop;i++)
{
     // can synchronize block now
     __syncthreads();
}
 // or here, only once, if there was only a loading from global into shared
 __syncthreads();

【讨论】：