课程内容（CUDA Shared Memory）
此次讲座是OLCF CUDA培训系列的第二部分，重点讲解了如何在CUDA编程中使用共享内存（Shared Memory）来优化性能。主讲人Tom Papatheodore详细说明了共享内存的特点、使用方法及其与全局内存的区别，并通过具体示例展示如何实现数据缓存和线程同步。最后，课程通过模板操作（Stencil operation）展示了在CUDA编程中共享内存的使用方式。
共享内存（Shared Memory）
在上一章节Cuda编程学习笔记（一）中讲到，我们需要为程序中的数据分配设备内存，在准备阶段我们需要在CPU和GPU上初始化我们需要的内存空间大小，然后把CPU的数据复制到GPU内存上，这里提到的GPU内存是全局内存（Global Memory）。但是，在本章节中我们所提到的内存则是共享内存（Shared Memory）。
共享内存有以下特性区别于全局内存：
- 全局内存一般指的是硬件上的内存，换句话讲就是他的实现是在GPU的DRAM，而不是在GPU本身（on-die），它是有一定数量的与GPU链接高速内存所组成。相反，共享内存指的是真正在GPU内部实现的内存也叫做on-chip memory。
- 共享内存有用更快的速度以及更高的带宽，粗略估算比全局内存大概快5倍。
- 共享内存是由用户自行管理的。
- 共享内存使用
__shared__关键字区分。 - 共享内存是单个区块（block）的逻辑资源，这意味着每个区块都可以管理一个只属于他自己的共享内存，这种资源是可以被区块内的线程所共享的，但却无法被其他区块以及其他区块的内存访问。
实现共享内存
视频中Tom通过一个模板操作（Stencil operation）的例子来详细解释了共享内存的用法，模板操作有点类似于神经网络的卷积操作，主要是通过特定的模式对一系列的数值数据进行处理，详细可见维基百科解释。图2，Tom在视频中解释道，这个stencil kernel就是一个滑动窗口，对于每一个元素计算包括左右三个元素在内共七个元素之和。
使用CUDA并行计算可以将数组分成多个block并行地计算，在这里使用共享内存就可以让一个block内的线程更快地读取到所需要的数据。我们可以分配一个blockDim+6大小的共享内存空间用来存储滑动窗口所需要的数据。详细代码会在课后练习中展示。
线程同步操作（__syncthreads）
在上面的例子中往共享内存中填入数据这一步操作是并行的，这会造成一个问题，当block内第i个线程需要做stencil操作的时候他需要读取共享内存中[i-3, i+3]共7个元素，但是我们无法保证共享内存中其他元素已经被填入，因此我们在这里要使用一个同步操作，保证共享内存的数组填入操作已经被完成。
这里我们可以使用__syncthreads()函数，这个函数会同步块内的所有线程，当且仅当所有线程执行到这个屏障（barrier），线程才会被允许执行后续的操作。
协作组（Cooperative Groups）
Tom在这节课中也讲到实际CUDA 9中也添加了不少的特性，例如协作组，这可以让我们更加灵活地去同步线程或者建立线程间的通信，值得注意的是这里的线程不再是一个线程块（block）内的线程而是可以跨线程块，同时，也可以是线程块内的某些线程而并非所有线程。这样的设计利用不同颗粒度的线程组给程序员提供了更加灵活、动态的线程分组。
但是Tom表示，协作组作为更高阶的内容会在稍晚的课程中讲解，这里仅仅一笔带过。我Google了一下相关资料，官方给出的技术博客讲的十分详细，有兴趣可以看一下。
共享内存与缓存（Shared Memory vs Cache）
共享内存（Shared Memory）是由用户管理的高速缓存，而缓存（Cache）则是隐式地管理数据。缓存遵循着某些模式去管理未来可能会用到的数据，与CPU缓存类似，GPU缓存也有自己的调度算法。在早些年的时候，GPU是没有缓存的，用户只能显式地管理自己的数据缓存，因此程序的效率会因为程序员对算法的理解而产生巨大的差异。
随着时代的发展，NVIDIA的GPU添加了大小可观的L1、L2缓存，例如Volta架构的GPU。在如此充沛的缓存空间的支持下，来自于精妙管理的共享内存所带了的优势就会被极大的缩小。图5展示了Volta架构与Pascal架构下共享内存所带来的性能提升对比，相比Pascal，Volta拥有更大缓存，因此共享内存所带来的优势显著减少。
正如官方论坛的一篇讨论帖子中，一位叫做njuffa的老哥说的：
I think it is fair to say that the importance of shared memory in CUDA programming has decreased with the advent of L1/L2 caches of competitive size in GPUs
我认为可以很公平地讲，共享内存在CUDA编程中的地位随着更具有竞争力的大小的L1/L2缓存的来临，已经在降低。
课程练习
hw2 – stencil1d
完善stencil1d的代码，代码相对简单直白，和课程中的内容高度一致，所以并没有额外的收获。
__global__ void stencil_1d(int *in, int *out) {
__shared__ int temp[BLOCK_SIZE + 2*RADIUS];
int gindex = threadIdx.x + blockIdx.x * blockDim.x;
int lindex = threadIdx.x + RADIUS;
// Read input elements into shared memory
temp[lindex] = in[gindex];
if (threadIdx.x < RADIUS) {
temp[lindex - RADIUS] = in[gindex - RADIUS];
temp[lindex + BLOCK_SIZE] = in[gindex + BLOCK_SIZE];
}
// Synchronize (ensure all the data is available)
__syncthreads();
// Apply the stencil
int result = 0;
for (int offset = -RADIUS; offset <= RADIUS; offset++)
result += temp[lindex + offset];
// Store the result
out[gindex] = result;
}
hw2 – matrix_mul_shared
这个其实与第一章节中hw1 – matrix_mul的代码差距不大，关键在于这次练习使用了共享内存，将内积需要的行列拆成若干个线程块，然后每个线程块记录一份共享内存。开始写的时候多少有点不习惯，因为对于输出矩阵的每个元素来说，都要管理As、Bs两块共享内存，但是这两块贡献内存中只有其中的第threadIdx.x行与第threadIdx.y列会被使用，就让我觉得这个共享内存存储了很多无用的数据。
__global__ void mmul(const float *A, const float *B, float *C, int ds) {
// declare cache in shared memory
__shared__ float As[block_size][block_size];
__shared__ float Bs[block_size][block_size];
int idx = threadIdx.x+blockDim.x*blockIdx.x; // create thread x index
int idy = threadIdx.y+blockDim.y*blockIdx.y; // create thread y index
if ((idx < ds) && (idy < ds)){
float temp = 0;
for (int i = 0; i < ds/block_size; i++) {
// Load data into shared memory
As[threadIdx.y][threadIdx.x] = A[idy * ds + (i * block_size + threadIdx.x)];
Bs[threadIdx.y][threadIdx.x] = B[(i * block_size + threadIdx.y) * ds + idx];
// Synchronize
__syncthreads();
// Keep track of the running sum
for (int k = 0; k < block_size; k++)
temp += As[threadIdx.y][k] * Bs[k][threadIdx.x]; // dot product of row and column
__syncthreads();
}
// Write to global memory
C[idy*ds+idx] = temp;
}
}最后我们来对比下使用共享内存与全局内存的代码执行效率的区别：
全局内存
(base) qibin:~/projects/cuda-training-series/exercises/hw1$ ./matrix_mul
Init took 0.161087 seconds. Begin compute
Done. Compute took 4.861984 seconds
Success!
共享内存
(base) qibin:~/projects/cuda-training-series/exercises/hw2$ ./matrix_mul_shared
Init took 0.157803 seconds. Begin compute
Done. Compute took 1.450261 seconds
Success!从结果上来看有足足2.4倍速度的提升，差不多70%的相对提升，相当的不错。