码迷,mamicode.com
首页 > 其他好文 > 详细

cuda基础---cuda通信机制

时间:2020-06-07 12:45:31      阅读:61      评论:0      收藏:0      [点我收藏+]

标签:tom   cad   通过   访问   表达式   grid   text   alc   声明   

1:同步函数

1)__syncthreads()

  实现线程块中的线程同步----保证线程块中所有线程执行到同一位置;

  只有当整个线程块都走向相同分支时,才能在条件语句中使用__syncthreads)(;

  一个warp内的线程不用同步;

  __syncthreads()调用花费时间周期,应尽量避免/节约使用。

 

2)memory fence 函数

  保证数据通信的可靠性,并不要求所有线程都运行到同一位置;

  多个线程间可以正确的操作共享数据,实现grid/block内的线程间通信。

(1)__threadfence()

  该函数调用后,该线程在此语句前对全局存储器或共享存储器的访问已经全部完成,且结果对grid内所有线程可见。

(2)__threadfence_block()

  该函数调用后,该线程在此语句前对全局存储器或共享存储器的访问已经全部完成,且结果对block内所有线程可见。

(3)实例---N个元素求和

 1 __device__ unsigned int count = 0;
 2 __shared__ bool isLastBlockDone;
 3 
 4 //每个block首先计算出数组的一个子序列的和,再将结果存储到全局内存
 5 //所有block完成后,由最后一个块完成求和操作
 6 
 7 __global__ void sum(const float* array,unsigned int N,float* result)
 8 {
 9 //每个block对输入数组的一个subset求和
10 float partialSum = calcultePartialSum(array,N);
11 if(threadIdx.x == 0)
12 {
13 //0号线程存储局部和partialSum并负责写回全局存储器
14 result[blockIdx.x] = partialSum;
15 //线程0要保证它的结果对所有其他的线程可见
16 __threadfence();
17 //每个block的0号线程负责计算结束后做一次标记,count++
18 unsigned int value = atomicInc(&count,gridDim.x);
19 //每个block的0号线程来判断该块是否是最后一个计算结束的快
20 isLastBlckDone = (value == gridDim.x-1);
21 }
22 //做一次同步,保证每个线程读到正确的isLastBlockDone值
23 __syncthreads();
24 
25 if(isLastBlockDone)
26 {
27 //最后的那个block负责局部和的球和
28 float totalSum = calculateTotalSum(result);
29 if(threadIdx.x == 0)
30 {//last block的线程0存储急速那和并将其协会全局存储器,count 置0来保证下一个kernel的正常计数
31 result[0] = totalSum;
32 count = 0;
33 }
34 }
35 }

3)kernel间通信

  通过global memory实现--pinned memory

4)GPU与CPU线程同步

  A:主机端代码使用cudaThreadSynchronize()----确定所有设备端线程均已运行结束

 

2:volatile关键字

  存在于全局变量或共享存储器中的变量通过此关键字声明为敏感变量。每次对该变量的引用都会被编译成一次真实的内存读指令。

 

3:ATOM操作

  当多个线程同时访问全局或共享存储器的同一位置时,保证每个线程能够实现对共享可写数据的互斥操作--在一个操作完成之前,其他任何线程都无法访问此地址。

 1 __global__ void testAtom(int* g_odata)
 2 {
 3 const unsigned int tid = blockDim.x*blockIdx.x+threadIdx.x;
 4 
 5 //算术原子操作
 6 atomicAdd(&g_odata[0],10);
 7 atomicMax(&g_odata[0],tid);
 8 atomicInc((unsigned int*)&g_odata[0],17);
 9 atomicCAS(&g_odata[0],tid-1,tid);
10 
11 //位原子操作
12 atomicAnd(&g_odata[8],2*tid+7);
13 atomicOr(&g_odata[8],1<<tid);
14 atomicXor(&g_odata[8],tid);        
15 }

A:图像亮度直方图统计

 

4:VOTE操作

  由下面两个函数实现此指令

A:int __all(int predicate);

  一个warp中所有的threads的判断表达式结果为真则返回1;否则返回0

B:int __any(int predicate);

  一个warp中只要存在thread的判断表达式结果为真,则返回1,否则返回0.

cuda基础---cuda通信机制

标签:tom   cad   通过   访问   表达式   grid   text   alc   声明   

原文地址:https://www.cnblogs.com/pengtangtang/p/13060003.html

(0)
(0)
   
举报
评论 一句话评论(0
登录后才能评论!
© 2014 mamicode.com 版权所有  联系我们:gaon5@hotmail.com
迷上了代码!