CUDA入门学习(三):共享内存与线程同步
来源:互联网 发布:淘宝优惠券套现 编辑:程序博客网 时间:2024/05/21 07:10
共享内存实际上是可受用户控制的一级缓存。每个SM中的一级缓存与共享内存共享一个64KB的内存段在开普勒架构的设备中,根据应用程序的需要,每个线程块可以配置为16KB的一级缓存或共享内存。而在费米架构的设备中,可以根据喜好选择16KB或者48KB的一级缓存或者共享内存。早期费米架构中只有固定的16KB共享内存而没有一级缓存。共享内存的延迟极低,大约有1.5TB/s的带宽,远远高于全局内存的190GB/s,但是它的速度只有寄存器的十分之一。只有当数据重复利用,全局内存合并,或线程之间由数据共享时使用共享内存才更合适,否则,数据直接从全局内存加载到寄存器性能会更好。申请共享内存后,其内容在每一个用到的block被复制一遍,使得在每个block内,每一个thread都可以访问和操作这块内存,而无法访问其他block内的共享内存。这种机制就使得一个block之内的所有线程可以互相交流和合作。
共享内存是基于存储体切换的架构(bank-switched architecture)。无论有多少个线程发起操作,每个存储体每个周期只执行一次操作。因此,如果线程束中的每个线程访问一个存储体,那么所有的线程操作都可以在一个周期内同时执行。此时无须顺序地访问,因为每个线程访问的存储体在共享内存中都是独立的,互不影响。此外,如果线程束中的所有线程同时访问相同地址的存储体时,就会想常量内存一样触发一个广播机制到线程束中的每一个线程。然而如果有其他的访问方式,存储体冲突将不同程度地得到解决。这意味着,线程访问共享内存需要排队等候,当一个线程访问时,其它线程都将阻塞。因此储存体应尽可能避免冲突。
这里有必要对第四张图说明下。所有线程访问同一存储体,如果是读的话会触发广播机制,但是如果是写的话会冲突,这将导致对同一存储体顺序进行32次访问操作,这里注意箭头朝向,该图是参考博文https://segmentfault.com/a/1190000007533157,但是可能博主没考虑到读写的区别吧,如有问题,欢迎指正修改。
我们来看个向量点乘的例子吧。
#include "../common/book.h"#define imin(a,b) (a<b?a:b)const int N = 33 * 1024;const int threadsPerBlock = 256;const int blocksPerGrid = imin( 32, (N+threadsPerBlock-1) / threadsPerBlock );__global__ void dot( float *a, float *b, float *c ) { __shared__ float cache[threadsPerBlock]; int tid = threadIdx.x + blockIdx.x * blockDim.x; int cacheIndex = threadIdx.x; float temp = 0; while (tid < N) { temp += a[tid] * b[tid]; tid += blockDim.x * gridDim.x; } // set the cache values cache[cacheIndex] = temp; // synchronize threads in this block __syncthreads(); // for reductions, threadsPerBlock must be a power of 2 // because of the following code int i = blockDim.x/2; while (i != 0) { if (cacheIndex < i) cache[cacheIndex] += cache[cacheIndex + i]; __syncthreads(); i /= 2; } if (cacheIndex == 0) c[blockIdx.x] = cache[0];}int main( void ) { float *a, *b, c, *partial_c; float *dev_a, *dev_b, *dev_partial_c; // allocate memory on the cpu side a = (float*)malloc( N*sizeof(float) ); b = (float*)malloc( N*sizeof(float) ); partial_c = (float*)malloc( blocksPerGrid*sizeof(float) ); // allocate the memory on the GPU HANDLE_ERROR( cudaMalloc( (void**)&dev_a, N*sizeof(float) ) ); HANDLE_ERROR( cudaMalloc( (void**)&dev_b, N*sizeof(float) ) ); HANDLE_ERROR( cudaMalloc( (void**)&dev_partial_c, blocksPerGrid*sizeof(float) ) ); // fill in the host memory with data for (int i=0; i<N; i++) { a[i] = i; b[i] = i*2; } // copy the arrays 'a' and 'b' to the GPU HANDLE_ERROR( cudaMemcpy( dev_a, a, N*sizeof(float), cudaMemcpyHostToDevice ) ); HANDLE_ERROR( cudaMemcpy( dev_b, b, N*sizeof(float), cudaMemcpyHostToDevice ) ); dot<<<blocksPerGrid,threadsPerBlock>>>( dev_a, dev_b, dev_partial_c ); // copy the array 'c' back from the GPU to the CPU HANDLE_ERROR( cudaMemcpy( partial_c, dev_partial_c, blocksPerGrid*sizeof(float), cudaMemcpyDeviceToHost ) ); // finish up on the CPU side c = 0; for (int i=0; i<blocksPerGrid; i++) { c += partial_c[i]; } #define sum_squares(x) (x*(x+1)*(2*x+1)/6) printf( "Does GPU value %.6g = %.6g?\n", c, 2 * sum_squares( (float)(N - 1) ) ); // free memory on the gpu side HANDLE_ERROR( cudaFree( dev_a ) ); HANDLE_ERROR( cudaFree( dev_b ) ); HANDLE_ERROR( cudaFree( dev_partial_c ) ); // free memory on the cpu side free( a ); free( b ); free( partial_c );}
主函数中的CUDA 内存申请,拷贝,释放我们就不具体介绍了,我们看核函数中的__shared__ float cache[threadsPerBlock] 这里的__shared__同样是CUDA中的关键字,和__global__一样,意思是申请一个共享内存,这里申请的大小为每个线程块的线程数目,这样每个线程都可以访问一个存储体,不会造成存储体冲突。并且会为每一个线程块分配相同大小的共享内存。这样我们就可以计算每个线程块的向量点乘之和,但是这里有个潜在的危险,如下图所示,如果一部分线程已经计算好了,另一部分还没,那么相加的结果肯定是错的。所以我们需要一个同步机制来保证每个线程块都计算好了,这里就用到了__syncthreads();
总结:本次主要介绍了块内线程如何通过共享内存进行协作,除了线程级通信之外,我们之后还将学校线程块之间的通信。
参考: [1]CUDA并行程序设计 —— GPU编程指南
[2] Addison.Wesley.CUDA.By.Example.Jul.2010.ISBN.0131387685
[3]https://segmentfault.com/a/1190000007533157
- CUDA入门学习(三):共享内存与线程同步
- cuda《学习笔记三》——共享内存和同步
- CUDA编程(七)共享内存与Thread的同步
- CUDA编程(七)共享内存与Thread的同步
- CUDA入门(三) 初探线程与块
- CUDA 学习(十一)、共享内存
- CUDA学习笔记(6) 共享内存与全局内存
- CUDA学习--内存处理之共享内存(3)
- 共享内存与同步线程的使用例子
- CUDA学习三:线程协作
- 多线程学习(三)线程的同步与锁
- 信号量学习 & 共享内存同步
- CUDA入门(4):CUDA内存模型
- CUDA学习(三)
- MFC多线程与线程同步 (三)
- Java线程(三):同步与锁
- Linux入门:线程同步与互斥(三)——信号量
- [CUDA]共享内存
- hive学习笔记
- 【Java】Class.forName() vs Class.loadClass()
- H5基础知识第十课时(JS函数调用和内部属性)
- hdoj 2186
- i-09
- CUDA入门学习(三):共享内存与线程同步
- hdoj 2187
- java创建简单注册登录网页
- 【leetcode】longest-consecutive-sequence
- Gi分支-分支的查询相关命令
- n-14
- hdoj 2188(选拔志愿者)
- hdu Problem-4277(dfs+set)
- codeforcces Karen and Game 思维