关于atomicAdd()函数

#include <stdio.h>
#include <time.h>
#include “cuda_runtime.h”
#define N 2
device int a;
global void cudatest2(int b){
int threadIndex=blockIdx.x;
atomicAdd( &a , threadIndex+2);
__threadfence();
while(b<N){
printf(“*** block=%d,a=%d \n",threadIndex,a);
b++;
}
b=0;
atomicAdd( &a , threadIndex+2);
__threadfence();
while(b<N){
printf("
* block=%d,a=%d **\n”,threadIndex,a);
b++;
}
}
int main(){
clock_t start = clock();
int dest = 0;
cudatest2<<<2,1>>>(dest);
cudaThreadSynchronize();
printf(“Time %1.2f sec\n”,double(clock()-start)/CLOCKS_PER_SEC);
return 0;
}
[attach]3068[/attach]

各位大神,求助啊,怎么会出现这个结果呢?我觉得两个Block的线程打印的结果应该相同啊。

LZ您好,两个block之间的先后顺序是不保证的,并且原子写入也不不保证和普通读取一致。
如果您将两个线程放在同一个block里面,并将原“__threadfence()”改为“__syncthreads()”,那么两个线程的最后打印结果应该是统一的。

祝您编码愉快~

谢谢ice版主您的回复,看来不同的Block之间的threads是不能同步的。

LZ您好,不客气的,要完成block之间的同步,只能是结束kernel。此时如果有后续的逻辑需要继续计算,可以重启kernel。

祝您编码愉快~

ice版主您好,通过结束kernal函数实现block之间的同步,以前见过,可要如果大量数据运算的话,频繁重启kernal函数会不会增加代码的整体运行时间啊,而且好多时候是上下文相关的,重启核函数,使得本来可以用局部的变量改成全局变量,这样运算也会增加相当多的运行时间,有没有什么别的好方法呢?

LZ您好,启动kernel的开销并不大,但是如您所说,一些中间变量必须先保存出来,再读进去,这个会影响效能的。

不过这也是没有办法的事情,一个grid中的所有blocks间同步,一般只此一种方法,别无分号。

祝您好运~

修改了部分措辞,更为准确一些。

好的,谢谢ice版主,这么晚还回复我们的问题。

不客气的, 我修改了6#的部分措辞,请以修改后的为准。

祝您编码顺利~


这也是我碰到的问题啊
担心进出kernel的开销比较大
看来也是没其它办法了

您也不妨实际测试一下,能满足要求的话,不妨直接用就好了。

祝您编码顺利~


目前的效果也是很不错的
对IPP有30的加速比