关于设置变量的问题

加速计算 CUDA
1

在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的初始化）是不是就会产生线程的访问冲突？应该怎么解决呢？而后面的原子操作是有意义的？

2

对与您的代码：
int nnp;
nnp = np[0] - 1;
atomicSub(&nnp, 1);

我看不出您使用atomicSub这里有何用途。因为该变量是私有的。

请楼主确定你要问的问题，和，你发的代码是匹配的？你发错了？？

3

呵呵，感谢版主回复！应该是我的问题没有阐述清楚，不过看了您的回复让我我大概确定了，代码里的nnp是私有变量，如果想让atomicSub(&nnp,1)起作用，nnp必须存储在全局存储器，对吧？

4

楼主您好，您没能成功看懂我的建议。

atomic*()的作用是什么？是为了用它们而用它们么？还是因为您的算法需要？

例如，1个block中的多个线程都需要对同一个shared memory中的变量分别进行+1操作，那么此时您需要让这多个线程安全的访问该变量，即读取旧值-增加1,回写新值的这3步不能中途被其他也要执行这三步的其他线程给打断。此时您才需要。

对于您的简单的线程私有的变量，您需要对他+1， 您直接对它+1即可。无需atomicAdd。

您觉得呢？

5

如您所言，我的确是因为算法需要所以nnp必须为所有线程访问，也就是说nnp不能为私有变量，所以必须要用到原子函数的，经版主提醒我方确定代码中的nnp是私有变量（我一直以为代码中的nnp是存储在全局存储器中的，所以才会使用原子函数），是我的问题描述不清楚，不过非常感谢版主的细心提醒和指点！版主真的很负责任！:loveliness:

6

您客气了。能够服务您是我的荣幸。