cuda的block thread wrap 同步与数据处理
2018-01-18 19:37
666 查看
cuda的block thread wrap 同步与数据处理
原创 2017年01月07日 15:45:57833主要涉及三个函数 和原子同步指令1 __syncthreads(); 使得同一个block之间线程间同步,达到相同的执行点后再往后执行,同时使得修改的全局以及共享内存对block内的线程可见
2 __threadfence(); 该线程在该语句前对全局存储器或共享存储器的访问已经全部完成,执行结果对grid中的所有线程可见。3 __thradfence_block(); 该线程在该语句前对全局存储器或者共享存储器的访问已经全部完成,执行结果对block中的所有线程可见。
注意CUDA中一个GRID包含多个BLOCK。每个BLOCK包含多个THREAD。另外还有一个概念叫做wrap。WRAP是线程束的意思,也就是GPU实际执行运算的时候是以wrap单位的。比如wrap=32或者wrap=16.假设你的wrap是16,每次运行一定是
4000
16个线程一起运行。即使你使用了一个block里面的1个thread.GPU也会凑足16个thread,只是这些thread处于不活跃状态,此时就浪费了15个thread线程的资源。http://www.myexception.cn/cuda/1931284.html
wrap了解清楚以后再看内存同步解释
1 同一个wrap内,对于global/shared 的内存操作是 该wrap内的所有线程即时可见的。 例如某个线程修改了一个内存,那么wrap内的其他线程访问的这个内存是修改后的。
2同一个block内。对于global/shared的内存操作。block内的线程不一定是即时可见的。例如一个block内有 32个线程。32个线程属于两个wrap wrapA wrapB。wrapA修改了某个内存,wrapB里面访问不一定是修改后的(wrapA访问的一定是修改后的,见第一条)。因此需要使用__threadfence_block() _syncthread() __threadfence()来实现让修改数据对其他线程可见。
3同一个grid内,对于global/shard的内存操作。其他线程不一定是可见的必须使用__threadfence()来实现让修改数据对其他线程可见。其他两个函数只对于block有效。
4原子操作。atomic的操作修改对于所有的线程都是即时可见的。
比如使用多个block对一个向量进行求和运算。向量长为30。分为三个blocka[0]-------a[29]首先 每个block分别计算10个数相加。block0 计算 a[0]-a[9] block1 计算 a[10]-a[19] block2 计算 a[20]-a[29]。计算的结果放到a[blockIdx.x]里面。由最后一个线程计算 a[0]+a[1]+a[2]得到最后结果。如何知道是最后一个block呢?我们用一个count来计数。每个block的线程0计算完毕的时候。使用原子加 atomicadd使得count+1.如果count等于block的大小。那么说明是最后一个block了。可以由他来执行最后的a[0]+a[1]+a[2]了此时就需要用到_threadfence()确保。a[1]和a[2]的结果对于block0的最后计算线程是可见的。
if(threadIdx.x == 0) { result[blockIdx.x] = partialSum;执行fence保证a[1]和a[2]的结果对于后续的 float totalSum = calclateTotalSum(rusult);是可见的。
__threadfence();
unsigned int value = atomicInc(&count ,gridDim.x);
isLastBlockDone = (value ==griddim.x-1)
}
__syncthreads();
if(isLastBlockDone)
{
float totalSum = calclateTotalSum(rusult);
if(threadIdx.x ==0)
{
result[0] = totalSum;
count = 0;
}
}
还有一个函数就走__threadfence_system()。这个就是保证所有的内存读写对于所有线程已完全可见了。包括主机线程等等
关于内存可以参见 http://www.mamicode.com/info-detail-226119.html
相关文章推荐
- cuda的block thread wrap 同步与数据处理
- 手动Handler与Thread处理UI数据同步问题
- 如何处理集群、分布架构的数据同步问题
- springmvc+spring线程池处理http并发请求数据同步控制问题
- 大量和批量的数据可以用数据库的同步机制处理
- Java如何处理多线程的数据同步问题
- CUDA之Thread、Wrap执行详解
- CUDA Thread Block:transpose
- java多线程处理导入数据拆分List集合,同步处理插入数据
- 测试基于传统复制下主从同步数据不一致处理
- 关于gluster数据不同步的处理
- CUDA编程——GPU架构,由sp,sm,thread,block,grid,warp说起
- 数据同步处理(以sql server 2000为例)
- linux shell同步数据并处理
- cuda之thread,block,gird详解
- Cuda 学习教程(五):GPU架构-Sp,sm,thread,block,grid,warp
- gpu/cuda-01-grid/block/thread
- Java如何利用synchronized处理多线程的数据同步问题
- cudaStreamSynchronize vs CudaDeviceSynchronize vs cudaThreadSynchronize CUDA中的屏障同步
- ijkplayer播放器的read_thread数据接受线程在处理音频的时候如果首个数据包没有发现audio stream 以后就不会有音频数据