#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间同步，一般只此一种方法，别无分号。
祝您好运~
修改了部分措辞，更为准确一些。
不客气的， 我修改了6#的部分措辞，请以修改后的为准。
祝您编码顺利~
嗯
这也是我碰到的问题啊
担心进出kernel的开销比较大
看来也是没其它办法了
system
10
您也不妨实际测试一下，能满足要求的话，不妨直接用就好了。
祝您编码顺利~
system
11
嗯
目前的效果也是很不错的
对IPP有30的加速比