CUDA通信机制

分类:  GPU/CUDA/OpenCL2009-08-25 21:57  1479人阅读  评论(3)  收藏  举报

CUDA体系架构支持怎么样的通信方法呢?往下看喽@_@

1.       __syncthreads()

Block内的线程同步。Block内所有线程都执行到这一位置(BAR指令),先到的要等后来的,到齐了后再继续后面的任务。执行结果对block内所有线程可见

2.       Memory Fence

l  __threadfence()

Grid内的线程同步。保证该语句前的,grid中所有线程发出的访存指令(对global memory / shared memory)都已结束。执行结果对grid内所有线程可见

l  __threadfence_block()

Block内的线程同步。保证该语句前的,block中所有线程发出的访存指令(对global memory / shared memory)都已结束。执行结果对block内所有线程可见

3.       Mapped pinned memory

CUDA2.2中引入mapped memory,允许多个GPU设备从内核程序中直接访问同一块pinned memory。CPU/GPUs共用一块内存,所以特别要注意做好同步。

4.       cudaThreadSynchronize()

GPU与CPU同步。Kernel启动后控制权异步返回,该函数用以确定所有设备端线程均已运行结束。

5.       volatile

__global__ void myKernel(int* result, int* myArray)

{

     int tid = threadIdx.x;

     int ref1 = myArray[tid] * 1;

     myArray[tid + 1] = 2;

     int ref2 = myArray[tid] * 1;

     result[tid] = ref1 * ref2;

}

上面这段程序,myArray[]初值为1。这段程序跑出来的结果result[]全是1,说明计算ref2时myArray[tid]并如有如约变成2(由tid-1修改),程序员是想用改变后的值2来着。为什么会出现这种情况呢?

是这样,在计算ref1时从显存取出myArray[tid],然后在计算ref2时compiler不会再去显存取一次而是去复用刚刚取到的结果,myArray[tid+1]=2这句操作呢是要写回显存,所以喽^o^ 可以通过加上volatile关键字来解决。

__global__ void myKernel(int* result, volatile int* myArray)

volatile限定符将变量声明为敏感变量,compiler认为其它线程可能随时会修改变量的值,因此每次对该变量的引用都会被编译成一次真实的访存指令。注意,在上面的代码中即使将myArray声明为valatile仍不能确保ref的值是2,这是因为线程tid可能在myArray[tid]被赋为2之前就已经进行了读操作,怎么办呢?做一个同步就好喽。

6.       atomic function

当多个线程同时访问global memory / shared memory同一位置时,原子函数来确保对这个32-/64-bit字的read-modify-write原子操作。保证每个线程能够实现对共享可写数据的互斥操作:在一个操作完成前,其它任何线程都无法访问此地址。

7.       vote

CUDA2.0引入的新特性,1.2计算能力的硬件支持。Vote指令作用域是一个warp。

int __all(int predicate); //warp中所有线程的判断表达式结果为真,则返回1,否则0

int __any(int predicate); //warp中只要一个thread的判断表达式结果为真,则返回1,否则0

  • 1
    点赞
  • 0
    收藏
    觉得还不错? 一键收藏
  • 0
    评论
评论
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值