CUDA的存储单元蕴含以下类型：

如下表所示：名称地位用处应用办法限度备注Register寄存器GPU的SM上存储局部变量每个SM上有成千上万个一个线程最大数量为256个须要省着用线程公有，最快线程退出则生效Shared memoryGPU芯片上实现Block内的线程通信，目前最快的多Thread沟通的中央__shared__修饰符须要__syncThreads()同步分为32个banks须要省着用，会影响流动warp数量可被1个block所有thread拜访，次快高带宽，低提早Local memory寄存单线程的大型数组和变量（Register不够时用它）没有特定的存储单元线程公有，速度较慢，速度与Global memory靠近Constant memory常量内存驻留在device memory中用于同一warp的所有thread同时拜访同样的常量数据，比方光线追踪__constant__修饰符必须在host端应用 cudaMemcpyToSymbol初始化没有特定的存储单元，然而有独自的缓存只读，全局Global memory等同于GPU显存驻留在device memory中输出数据，写入后果全局，速度较慢Texture memory纹理内存用于减速局部性拜访，比方热传导模型只读，全局，速度次于Shared Memory（提早比Shared Memory高，带宽比hared Memory小）Host memory：可分页内存主机端内存应用malloc拜访应用free开释不能够应用DMA拜访内存页能够置换到磁盘中另一种Host memory：又称：Page-locked Memory,Zero-Copy Memory主机端内存应用cudaMallocHost拜访应用cudaFreeHost开释属于另一种Global memory如何应用Shared Memory优化CUDA利用呢？

Shared Memory的特点是快的时候特地快，慢的时候特地慢。什么时候快？同一warp中所有线程拜访不同的banks或者 同一warp中所有线程读取同一地址（通过播送）什么时候慢？同一warp中多个线程拜访同一个bank的不同地址（此时将产生 bank conflict）串行拜访请留神：bank conflict产生的起因就是 warp的调配和bank的调配重叠了：

如何防止bank conflict，简略的办法是Padding法（如同叫做补边）：

通过减少一个空列，让bank强行错位，使得每段间断的数据被调配到不同的bank中。具体做法很简略：

就是在设置Shared Memory的时候，不设置成 方阵BLOCK_SIZE X BLOCK_SIZE,而设置成 BLOCK_SIZE X (BLOCK_SIZE+1).最初，咱们能够应用Shared Memory优化mXn, nXk的矩阵乘 的代码，进步访存的效率。具体方法如下：申请两块 Shared Memory，都是BLOCK_SIZE X BLOCK_SIZE 大小。一个沿着矩阵mXn滑动，一个沿着矩阵 nXk滑动。将 子集的后果累加到 目标矩阵中：

具体的代码如下：__global__ void gpu_matrix_mult_shared(int d_a, int d_b, int *d_result, int m, int n, int k)
{

__shared__ int tile_a[BLOCK_SIZE][BLOCK_SIZE];__shared__ int tile_b[BLOCK_SIZE][BLOCK_SIZE];int row = blockIdx.y * BLOCK_SIZE + threadIdx.y;int col = blockIdx.x * BLOCK_SIZE + threadIdx.x;int tmp = 0;int idx;for (int sub = 0; sub < gridDim.x; ++sub) {    idx = row * n + sub * BLOCK_SIZE + threadIdx.x;    tile_a[threadIdx.y][threadIdx.x] = row<n && (sub * BLOCK_SIZE + threadIdx.x)<n? d_a[idx]:0;    idx = (sub * BLOCK_SIZE + threadIdx.y) * n + col;    tile_b[threadIdx.y][threadIdx.x] = col<n && (sub * BLOCK_SIZE + threadIdx.y)<n? d_b[idx]:0;        __syncthreads();    for (int k = 0; k < BLOCK_SIZE; ++k)     {        tmp += tile_a[threadIdx.y][k] * tile_b[k][threadIdx.x];    }    __syncthreads();}if(row < n && col < n){    d_result[row * n + col] = tmp;}

}并将后面 代码中调用矩阵乘的中央：gpu_matrix_mult<<<dimGrid, dimBlock>>>(d_a, d_b, d_c, m, n, k);  改为 gpu_matrix_mult_shared<<<dimGrid, dimBlock>>>(d_a, d_b, d_c, m, n, k); 其余不变。编译，执行：

批改blocksize，将其别离改为 16，8，4，再进行统计汇总：矩阵MXN（m)矩阵NXK(n)矩阵NXK(k)blocksizestop-start(ms)100100100321.83286100100100161.2736510010010081.2329210010010043.528651001001006（补测）2.199910010010012（补测）1.34755从下面的后果来看，blocksize为8，16，32时如同差别不大，然而blocksize为4的时候速度降得比拟厉害。blocksize为4时，其实并没有产生bank conflict！而只是因为4X4，只有16个线程，而一个warp须要32个线程，所以相当于计算时，有一半算力被节约掉了，进而速度慢了一倍。专家建议，至多应该NXN>32比拟好。将 矩阵从100改为1000试试，然而发现一旦改为1000后，CPU计算可能算不过去了，须要将CPU那局部代码和前面比拟的代码屏蔽掉。

再从新统计：矩阵MXN（m)矩阵NXK(n)矩阵NXK(k)blocksizestop-start(ms)10001000100032265.10610001000100016228.091000100010008202.3821000100010004518.3151000100010006（补测）386.17110001000100012（补测）246.29