GPU for loops:避免扭曲发散和隐式同步线程
GPU for loops: avoid warp divergence & implicit syncthreads
我的情况是:经纱中的每个线程都有自己完全独立的&不同的数据阵列。所有线程都在其数据数组上循环。每个线程的循环迭代次数不同。(我知道这会产生成本)。
在for循环中,每个线程需要在计算三个浮点值后保存最大值。在for循环之后,warp中的线程将通过检查warp中仅由其"相邻线程"计算的最大值(由奇偶校验确定)来"通信"。
问题:
- 如果我通过乘法来避免"max"运算中的条件句,这将避免扭曲发散,对吧?(请参阅下面的示例代码)
- (1.)中提到的额外乘法运算是值得的,对吧?-即比任何种类的翘曲发散都快得多
- 导致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上实现我的算法之前,我正在研究算法的各个方面,以确保它值得实现。所以请耐心等待。
- 是
- 我的猜测是否定的——这取决于你将如何写另一个带有if的版本
编译器可能会使用谓词来屏蔽不需要的写入,在这种情况下,不会有真正的线程分歧,只有几个执行但屏蔽的写入指令
你应该让编译器发挥它的魔力,比较两个版本的反编译代码,以确定什么是更好的解决方案
在计算最大有符号整数d=a>b的特殊情况下?a:b翻译成一条PTX ISA指令max.s32,所以真的没有必要让它像你做的那样复杂。。。只需将最大值计算到一个临时变量中,然后进行一次无条件写入 - 是的,但synthreads屏障是一个块内屏障,不是块间屏障,当然也不是mp间屏障
相关文章:
- 线程过程中的线程同步问题
- 具有不断变化的资源量的 C++ 多线程同步
- C++服务器上的线程同步
- 由 std::shared_ptr 向量指向的数据之间的线程同步
- STD ::原子与静态变量用于线程同步
- 使用 Visual c++ 进行多线程同步不起作用
- 依靠网络 I/O 在C++中提供跨线程同步
- 与上一个线程具有相同ID的线程同步
- 高精度线程同步
- 将两个线程同步到同一个计时器
- 如何将递归函数的线程与子线程同步
- 线程同步 101.
- qt 中的线程同步
- 线程同步和成员函数
- 布尔变量的线程同步
- 线程同步与boost::condition_variable
- 与线程同步
- 如何跨线程同步变量?C++
- 线程同步:等待两个bool变量
- 多线程同步