【问题标题】:GPU for loops: avoid warp divergence & implicit syncthreadsGPU for loops:避免扭曲发散和隐式同步线程
【发布时间】:2013-02-04 11:25:15
【问题描述】:

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

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

问题:

  1. 如果我通过乘法来避免“最大”操作中的条件,这将避免扭曲发散,对吗? (参见下面的示例代码)
  2. (1.) 中提到的额外乘法运算是值得的,对吗? - 即比任何形式的经线发散都要快。
  3. 导致扭曲发散(所有线程的一组指令)的相同机制可以在 for 循环(与非 GPU 计算中的“#pragma omp for”语句非常相似)。因此,在一个线程检查另一个线程保存的值之前,我不需要在 for 循环之后对 warp 进行“同步线程”调用,对吗? (这是因为“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 上实施我的算法之前,我正在研究算法的各个方面,以确保它值得实施。所以请多多包涵..

【问题讨论】:

    标签: c++ c cuda gpu gpgpu


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

    【讨论】:

    • 不确定我是否理解您对 2) 的回答。我试图获得最大的 float 值。所以你说“max.f32”是一个可以并行计算的函数,即会产生扭曲发散惩罚吗?
    • @MatthewParks 是的,对不起,我把它搞混了。对于浮点数,指令为 max{.ftz}.f32。
    • @MatthewParks 我不确定你在这里问什么:“所以你说“max.f32”是一个可以并行计算的函数,即不会产生扭曲发散惩罚? ”。您是在谈论由于数组长度不同而导致的分歧吗?这与最大值计算无关。或者你是在问分歧是否会做类似if (a &gt; b &amp;&amp; a &gt; c) max_data[threadIdx.x][j] = a; 的事情?
    • 我不是在谈论数组长度的差异——我理解这是一个单独的问题。我确实在问“if”语句的分歧。
    • 所以基本上,架构可以通过使用称为“谓词”的东西一步完成“最大”和“最小”评估(即没有条件,没有分歧),因此没有分歧惩罚。
    猜你喜欢
    • 2014-12-09
    • 2015-12-10
    • 2020-02-15
    • 1970-01-01
    • 1970-01-01
    • 2021-04-04
    • 1970-01-01
    • 1970-01-01
    • 1970-01-01
    相关资源
    最近更新 更多