[CUDA]共享内存

博客分类

    于GPU上启动的每个线程块上的共享内存，CUDA C编译器都会创建该变量的一个副本。同一线程块的每个线程都共享这块内存，但是线程无法看到也不能修改其他线程块中的共享内存。这样做可以使得一个线程块中的多个线程能够在计算上通信和协作。

    共享内存缓冲区驻留在物理GUP上，因此访问共享内存的延迟远远低于访问普通缓冲区的延迟。

    共享内存的声明方式是在前面加上  __shared__

伦理片 http://www.dotdy.com/

    为了保持进程同步，可以使用cuda的函数__syncthreads();。这个函数的作用是为了确保线程块的每个线程都执行完__syncthreads();之前的语句后，才会执行下面的语句。

 

    出于易于理解，写了一个简单的程序，大致功能就是对于一列数，每四个数字进行逆转位置

 1 2 3 4 5 6 7 8  ----》 4 3 2 1 8 7 6 5

Cpp代码  

1. #include<cuda_runtime.h>  
2. #include<windows.h>  
3. #include<iostream>  
4. using namespace std;  
5. const int nMax = 50;  
6. __global__ void exchangeKernel(float *aaa)  
7. {  
8.     int offset = threadIdx.x + blockDim.x * blockIdx.x;  
9.     int x = threadIdx.x;  
10.     __shared__ float tmp[4];  
11.     int a = offset / 4;  
12.     a = (a + 1) * 4 - (offset - a * 4) - 1; ///a为同一个block对应位置的offset  
13.     tmp[x] = aaa[a];  
14.   
15.     __syncthreads();  
16.     aaa[offset] = tmp[x];  
17. }  
18.   
19. int main(){  
20.     float a[nMax];  
21.     float *devA;  
22.     for (int i = 0; i < nMax; i++){  
23.         a[i] = i;  
24.     }  
25.     cudaMalloc((void**)&devA, nMax*sizeof(float));  
26.   
27.     cudaMemcpy(devA, a, nMax*sizeof(float), cudaMemcpyHostToDevice);  
28.   
29.     exchangeKernel << <10, 4 >> >(devA );  
30.   
31.     cudaMemcpy(a, devA, nMax*sizeof(float), cudaMemcpyDeviceToHost);  
32.   
33.   
34.     for (int i = 0; i < 40; i++){  
35.         cout << a[i] << " ";  
36.     }cout << endl;  
37.     cudaFree(devA);  
38.     cin >> a[0];  
39.     return 0;  
40. }  

 影音先锋电影 http://www.iskdy.com/