在kernel里有这样一段代码 int i = threadIdx.x + blockDim.x*blockIdx.x ;
if(dev_pos_e[i]>=nc||dev_pos_e[i]<0)
{ int nnp;
nnp=np[0]-1;
dev_x_e[i]=dev_x_e[nnp];
dev_vx_e[i]=dev_vx_e[nnp];
dev_vy_e[i]=dev_vy_e[nnp];
dev_vz_e[i]=dev_vz_e[nnp];
atomicSub(&nnp,1);
}
请问:
变量nnp是存储在各个线程的local memory里面的吗?也就是说每个符合条件的线程在执行这段代码时都会从 nnp=np[0]-1算起吗?如果是的话,那么后面的原子操作是不是就没意义了?
如果我想所有线程访问同一个nnp,是不是就要将其在函数外设置为 device int nnp?那么上述代码中的 nnp=np[0]-1
(即nnp的初始化)是不是就会产生线程的访问冲突?应该怎么解决呢?而后面的原子操作是有意义的?
对与您的代码:
int nnp;
nnp = np[0] - 1;
atomicSub(&nnp, 1);
我看不出您使用atomicSub这里有何用途。因为该变量是私有的。
请楼主确定你要问的问题,和,你发的代码是匹配的?你发错了??
呵呵,感谢版主回复!应该是我的问题没有阐述清楚,不过看了您的回复让我我大概确定了,代码里的nnp是私有变量,如果想让atomicSub(&nnp,1)起作用,nnp必须存储在全局存储器,对吧?
楼主您好,您没能成功看懂我的建议。
atomic*()的作用是什么?是为了用它们而用它们么?还是因为您的算法需要?
例如,1个block中的多个线程都需要对同一个shared memory中的变量分别进行+1操作,那么此时您需要让这多个线程安全的访问该变量,即读取旧值-增加1,回写新值的这3步不能中途被其他也要执行这三步的其他线程给打断。此时您才需要。
对于您的简单的线程私有的变量,您需要对他+1, 您直接对它+1即可。无需atomicAdd。
您觉得呢?
如您所言,我的确是因为算法需要所以nnp必须为所有线程访问,也就是说nnp不能为私有变量,所以必须要用到原子函数的,经版主提醒我方确定代码中的nnp是私有变量(我一直以为代码中的nnp是存储在全局存储器中的,所以才会使用原子函数),是我的问题描述不清楚,不过非常感谢版主的细心提醒和指点!版主真的很负责任!:loveliness: