c++ - GPU 循环 : avoid warp divergence & implicit syncthreads

Question

我的情况:warp 中的每个线程都在其自己完全独立且不同的数据数组上运行。所有线程循环遍历它们的数据数组。每个线程的循环迭代次数不同。 (这会产生成本，我知道)。

在for循环中，每个线程计算完三个 float 后需要保存最大值。在 for 循环之后，warp 中的线程将通过检查仅由它们在 warp 中的“相邻线程”计算的最大值(由奇偶校验确定)来“通信”。

问题:

1. 如果我通过乘法避免“最大”运算中的条件，这将避免扭曲发散，对吧？ (参见下面的示例代码)
2. (1.) 中提到的额外乘法运算是值得的，对吧？ - 即比任何类型的曲速发散都快得多。
3. 导致 warp 发散(所有线程的一组指令)的相同机制可以在 for 循环结束时用作隐式“线程屏障”(对于 warp)(与非 GPU 计算中的“#pragma omp for”语句的方式非常相似)。因此，在一个线程检查另一个线程保存的值之前，我不需要在 for 循环之后对 warp 进行“syncthreads”调用，对吗？ (这是因为“synthreads”仅适用于“整个 GPU”，即 inter-warp 和 inter-MP，对吧？)

示例代码:

__shared__ int N_per_data;  // loaded from host
__shared__ float ** data;  //loaded from host
data = new float*[num_threads_in_warp];
for (int j = 0; j < num_threads_in_warp; ++j)
     data[j] = new float[N_per_data[j]];

// the values of jagged matrix "data" are loaded from host.


__shared__  float **max_data = new float*[num_threads_in_warp];
for (int j = 0; j < num_threads_in_warp; ++j)
     max_data[j] = new float[N_per_data[j]];

for (uint j = 0; j <  N_per_data[threadIdx.x]; ++j)
{
   const float a = f(data[threadIdx.x][j]);
   const float b = g(data[threadIdx.x][j]);
   const float c = h(data[threadIdx.x][j]);

  const int cond_a = (a > b)  &&  (a > c);
  const int cond_b = (b > a)  && (b > c);
  const int cond_c = (c > a)  && (c > b);

  // avoid if-statements.  question (1) and (2)
  max_data[threadIdx.x][j] =   conda_a * a  +  cond_b * b  +  cond_c * c; 
}



 // Question (3):
// No "syncthreads"  necessary in next line:

// access data of your mate at some magic positions (assume it exists):
float my_neighbors_max_at_7 = max_data[threadIdx.x + pow(-1,(threadIdx.x % 2) == 1) ][7]; 

在 GPU 上实现我的算法之前，我正在研究算法的各个方面，以确保值得为实现付出努力。所以请多多包涵..

Answer 1

1. 是
2. 我的猜测是否定的 - 取决于您如何使用 ifs 编写其他版本。
   编译器可能会使用谓词来屏蔽不需要的写入，在这种情况下，不会有真正的线程分歧，只有少数执行但屏蔽了写入指令。
   你应该让编译器发挥它的魔力，并比较两个版本的反编译代码，以确定哪个是更好的解决方案。
   在您计算最大有符号整数 d = a > b 的特定情况下？ a : b 转换为一条 PTX ISA 指令 max.s32，因此实际上没有必要让它变得像您所做的那样复杂...只需将最大值计算到一个临时变量中并执行一次无条件写入。
3. 是的，但是 synthreads 屏障是 block 内屏障，不是 block 间屏障，当然也不是 mp 间屏障。