为什么要使用共享内存呢,因为共享内存的访问速度快。这是首先要明确的,下面详细研究。
cuda程序中的内存使用分为主机内存(host memory) 和 设备内存(device memory),我们在这里关注的是设备内存。设备内存都位于gpu之上,前面我们看到在计算开始之前,每次我们都要在device上申请内存空间,然后把host上的数据传入device内存。cudaMalloc()申请的内存,还有在核函数中用正常方法申请的变量的内存。这些内存叫做全局内存,那么还有没有别的内存种类呢?常用的还有共享内存,常量内存,纹理内存,他们都用一些不正常的方法申请。
他们的申请方法如下:
共享内存:__shared__ 变量类型 变量名;
常量内存:__constant__ 变量类型 变量名;
纹理内存:texture<变量类型> 变量名;
存储类型 | 寄存器 | 共享内存 | 纹理内存 | 常量内存 | 全局内存 |
带宽 | ~8TB/s | ~1.5TB/s | ~200MB/s | ~200MB/s | ~200MB/s |
延迟 | 1个周期 | 1~32周期 | 400~600周期 | 400~600周期 | 400~600周期 |
他们在不同的情况下有各自的作用,他们最大的区别就是带宽不同,通俗说就是访问速度不同。后面三个看起来没什么不同,但是他们在物理结构方面有差别,适用于不同的情况。
共享内存实际上是可受用户控制的一级缓存。申请共享内存后,其内容在每一个用到的block被复制一遍,使得在每个block内,每一个thread都可以访问和操作这块内存,而无法访问其他block内的共享内存。这种机制就使得一个block之内的所有线程可以互相交流和合作。下面的例子中就显示了线程之间的交流和合作。
这个例子计算的是两个向量的点积。
1 /* 2 * Copyright 1993-2010 NVIDIA Corporation. All rights reserved. 3 * 4 * NVIDIA Corporation and its licensors retain all intellectual property and 5 * proprietary rights in and to this software and related documentation. 6 * Any use, reproduction, disclosure, or distribution of this software 7 * and related documentation without an express license agreement from 8 * NVIDIA Corporation is strictly prohibited. 9 * 10 * Please refer to the applicable NVIDIA end user license agreement (EULA) 11 * associated with this source code for terms and conditions that govern 12 * your use of this NVIDIA software. 13 * 14 */ 15 16 17 #include "../common/book.h" 18 19 #define imin(a,b) (a<b?a:b) 20 21 const int N = 33 * 1024; 22 const int threadsPerBlock = 256; 23 const int blocksPerGrid = 24 imin( 32, (N+threadsPerBlock-1) / threadsPerBlock ); 25 26 27 __global__ void dot( float *a, float *b, float *c ) { 28 __shared__ float cache[threadsPerBlock]; 29 int tid = threadIdx.x + blockIdx.x * blockDim.x; 30 int cacheIndex = threadIdx.x; 31 32 float temp = 0; 33 while (tid < N) { 34 temp += a[tid] * b[tid]; 35 tid += blockDim.x * gridDim.x; 36 } 37 38 // set the cache values 39 cache[cacheIndex] = temp; 40 41 // synchronize threads in this block 42 __syncthreads(); 43 44 // for reductions, threadsPerBlock must be a power of 2 45 // because of the following code 46 int i = blockDim.x/2; 47 while (i != 0) { 48 if (cacheIndex < i) 49 cache[cacheIndex] += cache[cacheIndex + i]; 50 __syncthreads(); 51 i /= 2; 52 } 53 54 if (cacheIndex == 0) 55 c[blockIdx.x] = cache[0]; 56 } 57 58 59 int main( void ) { 60 float *a, *b, c, *partial_c; 61 float *dev_a, *dev_b, *dev_partial_c; 62 63 // allocate memory on the cpu side 64 a = (float*)malloc( N*sizeof(float) ); 65 b = (float*)malloc( N*sizeof(float) ); 66 partial_c = (float*)malloc( blocksPerGrid*sizeof(float) ); 67 68 // allocate the memory on the GPU 69 HANDLE_ERROR( cudaMalloc( (void**)&dev_a, 70 N*sizeof(float) ) ); 71 HANDLE_ERROR( cudaMalloc( (void**)&dev_b, 72 N*sizeof(float) ) ); 73 HANDLE_ERROR( cudaMalloc( (void**)&dev_partial_c, 74 blocksPerGrid*sizeof(float) ) ); 75 76 // fill in the host memory with data 77 for (int i=0; i<N; i++) { 78 a[i] = i; 79 b[i] = i*2; 80 } 81 82 // copy the arrays 'a' and 'b' to the GPU 83 HANDLE_ERROR( cudaMemcpy( dev_a, a, N*sizeof(float), 84 cudaMemcpyHostToDevice ) ); 85 HANDLE_ERROR( cudaMemcpy( dev_b, b, N*sizeof(float), 86 cudaMemcpyHostToDevice ) ); 87 88 dot<<<blocksPerGrid,threadsPerBlock>>>( dev_a, dev_b, 89 dev_partial_c ); 90 91 // copy the array 'c' back from the GPU to the CPU 92 HANDLE_ERROR( cudaMemcpy( partial_c, dev_partial_c, 93 blocksPerGrid*sizeof(float), 94 cudaMemcpyDeviceToHost ) ); 95 96 // finish up on the CPU side 97 c = 0; 98 for (int i=0; i<blocksPerGrid; i++) { 99 c += partial_c[i]; 100 } 101 102 #define sum_squares(x) (x*(x+1)*(2*x+1)/6) 103 printf( "Does GPU value %.6g = %.6g?\n", c, 104 2 * sum_squares( (float)(N - 1) ) ); 105 106 // free memory on the gpu side 107 HANDLE_ERROR( cudaFree( dev_a ) ); 108 HANDLE_ERROR( cudaFree( dev_b ) ); 109 HANDLE_ERROR( cudaFree( dev_partial_c ) ); 110 111 // free memory on the cpu side 112 free( a ); 113 free( b ); 114 free( partial_c ); 115 }
我们首先关注核函数dot。__shared__ float cache[threadsPerBlock];就是这节重点,申请cache数组时,由于使用了共享内存,则每一个block里面都有一份cache,使得block内的thread都可以访问和操作其各自的cache数组。
1 while (tid < N) { 2 temp += a[tid] * b[tid]; 3 tid += blockDim.x * gridDim.x; 4 }
这一段我们相当熟悉,每个线程计算若干对a,b的乘积,然后相加。然后这样cache[cacheIndex] = temp;将结果存入cache中。这时,每一个线程的结果都被存在了cache数组中,我们知道接下来要对数组求和,然而这里有潜在的危险,那就是我们不知道所有线程是否已经将数据写入了cache,也就是说,是否每一个线程都已经执行完了第39行。这里就需要等待,等待所有线程执行到同一位置,这就是 __syncthreads();的作用。这个函数称为同步函数,即在所有线程全部执行到__syncthreads()为止,谁也不许动,其后任何代码都无法执行。
因此,我们可以很清楚的明白所有线程全部执行完了第39行,然后同步解除,大家再一起往前走。做加法。
1 int i = blockDim.x/2; 2 while (i != 0) { 3 if (cacheIndex < i) 4 cache[cacheIndex] += cache[cacheIndex + i]; 5 __syncthreads(); 6 i /= 2; 7 } 8 9 if (cacheIndex == 0) 10 c[blockIdx.x] = cache[0];
这段就不难理解了,逐对相加,最后cache【0】位置的数就是结果。将其值存入c数组,准备导出。
剩下的main函数部分是如下几步操作(和前面学习的差不多):
1.为输入输出数组分配内存
2.将a,b数组付初值,然后复制给device中,cudaMemcpy()
3.调用核函数执行并行计算。
4.device值返回后数组c求和。
很明显,由于我们使用了共享内存存储cache数组,使得在操作cache数组时的速度有了大幅提高(相比于全局内存)。共享内存的意义也就在此。
现在,请观察下面的两组代码:
while (i != 0) { if (cacheIndex < i) cache[cacheIndex] += cache[cacheIndex + i]; __syncthreads(); i /= 2; }
while (i != 0) { if (cacheIndex < i) { cache[cacheIndex] += cache[cacheIndex + i]; __syncthreads(); } i /= 2; }
下面的代码中由于if的存在,只有部分线程包含同步操作。代码似乎得到了优化。但是真的如此吗
当然不是的,上面的红字“所有线程全部执行到__syncthreads()为止”,所有很重要,<<<>>>中launch了多少个threadperblock,那么就必须要等待所有的线程,一个都不能少。由于if的存在,上例中部分线程永远都不可能执行到cache[cacheIndex] += cache[cacheIndex + i];这一步,因此就要永远等待下去,因而程序无法执行。
总结:在能用共享内存的时候尽量用,进而提高block内的执行效率,但是在同步问题上一定要慎重。。。