GPU共享内存小结---pycuda

先来点题外话，由于导师的要求使用GPU加速但是又没有系统的学习就开始了使用pycuda进行GPU的编程，变用边学。

在优化程序的时候shared memory是必用的，所以理解共享内存是个什么东西才可以做到自己对数据是如何跑的有数。
先看看这张GPU的存储结构图（偷的图，哈哈^_^）：

在图中我们关心的重点是，每一个block都对应的一个shared memory，线程块之可以相互交流通信就是通过共享内存进行传递的。
我们要记住的一点就是每一个线程块都有自己对用的一块共享内存。这么说，是感受不到的，下面我就用一个简单的程序来形象的看看这个效果。

import numpy as np
import pycuda.autoinit
import pycuda.driver as cuda
from pycuda.compiler import SourceModule
mod = SourceModule('''
__global__ void Text_GPU(float *A , float *B){

    int bid = blockIdx.x;  
    int tid = threadIdx.x;

    __shared__ float s_data[128];

    s_data[tid] = A[bid*128 + tid];
    __syncthreads();

    for(int i = 64;i>0;i/=2){
        if(tid < i)
            s_data[tid] = s_data[tid] + s_data[tid +i];
        __syncthreads();
    }
    if(tid == 0)
        B[bid] = s_data[0];

}
''')
Text_GPU = mod.get_function("Text_GPU")
A = np.ones((32,128),dtype=np.float32)
B = np.ones((32,),dtype=np.float32)
Text_GPU(cuda.In(A),cuda.InOut(B),grid=(32,1),block=(128,1,1))
print(B)

这个程序非常的简单，我们可以看到创建了A（32,128）的一个矩阵再复制给共享内存进行一个规约之后将和存入数组B。
其中值得体会的是程序中的这句话：

s_data[tid] = A[bid*128 + tid];

s_data只有128 而 A的大小是128*32，要是我们不知道GPU的存储结构那我们就一头雾水。
下面讲解下这个知识点：
一个线程块开了128的共享内存，而同时调用的有32个线程块，每一个块都有128的共享内存，所以就满足了和A的大小相等了。
看看输出的结果：

对A矩阵的每一列进行求和，结果放入B中。结果有32个说明有32共享内存一起运行。
好了，这个知识点记录完了，88。0.0