#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的加速比