一个求max的方法,结果不稳定,n个thread中对应值与register中的一个变量比较 如果大则将该变量赋值到register,小则不变,经过测试,获得的最大值不稳定,请问我该怎么保持结果的稳定呢!!
万分感谢!
代码
global max(float *dev_max){
int ind=blockIdx.x;
dev_max[0]=max(ind,dev_max[0]);
}
原因分析:针对一个100线程运算的最后阶段,a线程取得的dev_max是60,a线程的ind 是99,所以a线程的结果应该是99;
与此同时,b线程取得的dev_max是55,a线程的ind 是98,所以a线程的结果应该是98;
但是由于a线程先于b线程完成,所以导致结果是98,请问各位大大,有没有较好的办法解决这个问题。
是不是还要加什么呢!
楼主您好,
首先感谢您的午夜来访,
但需要指出的是,贵代码存在严重的竞态, 此BUG将导致您无法得到正确结果。
解决方案:
请您将原来您的block中的直接max判断改成多线程安全的reduction实现。
例如还是您这个100个数的求最大值,您可以考虑这样做:
global void phystry_max(float *dev_max)
{
int ind = blockIdx.x;
shared float buffer[128];
if (ind < 100) buffer[ind] = dev_max[ind]; else buffer[ind] = dev_max[0];
__syncthreads();
for (int i = 64; i > 0; i >>= 1)
{
if (ind < i) buffer[ind] = max(buffer[ind], buffer[ind + i]);
__syncthreads();
}
dev_max[0] = buffer[0];
}
这个范例代码的整个读取,判断最大值,写入的过程是有序可控的,不存在竞争状态。
因此可以正确判断这100个数的最大值。
请您参考此代码给出您的自己的正确实现(注意,手写无编译验证,仅供参考)。
感谢您的午夜来访,祝您午夜愉快。
EDIT: 修正笔误一处。
多谢 !!感激!!!可是按您这个方法不就成了顺序执行了吗?
大神 代码很漂亮!
LZ您好:
2#横扫斑竹的代码是并行规约的,请您仔细看下。
祝您编码顺利~
花了大半个小时 终于是看懂了!!
唉!!小白还是小白啊。。
谢谢两位版主的指点!
恭喜您成功获得大量经验值!
并行规约是公认的CUDA编程的入门标杆,您现在已经走在正确的康庄大道上,祝您好运~
感谢深夜造访论坛,祝您晚安~