共享内存启用块中线程之间的协作。当一个块中的多个线程使用全局内存中的相同数据时，共享内存只能用于从全局内存访问一次数据。共享内存还可以用来避免未合并的内存访问，方法是从全局内存中加载和存储合并模式的数据，然后将其重新排列在共享内存中。除了存储体冲突之外，共享内存中的变形对非连续或未对齐访问不会造成任何损失。
对于具有维度Mxw的A，具有维度wxN的维度B以及维度MxN的维度C的情况，通过矩阵乘法C = AB的简单示例来说明共享存储器的使用。为了简化内核，M和N是32的倍数，对于计算能力2.0或更高的设备，w为32。
问题的自然分解是使用wxw线程的块和瓦片大小。因此，就wxw瓦片而言，A是列矩阵，B是行矩阵，C是它们的外积;参见图9.以M / w块为单位的N / W网格被启动，其中每个线程块从A的单个块和B的单个块计算C中不同块的元素
为此，simpleMultiply内核（未优化矩阵乘法）计算矩阵C的一个瓦片的输出元素。
未经优化的矩阵乘法：
__global__ void simpleMultiply(float *a, float* b, float *c,
int N)
{
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
float sum = 0.0f;
for (int i = 0; i < TILE_DIM; i++) {
sum += a[row*TILE_DIM + i] * b[i*N + col];
}
c[row*N + col] = sum;
}在未优化的矩阵乘法中，a，b和c分别是矩阵A，B和C的全局存储器的指针; blockDim.x，blockDim.y和TILE_DIM都等于w。 wxw-thread块中的每个线程计算一个图块中的一个元素
C.行和列是C中元素的行和列由特定的线程计算。 for循环结束我将一行A乘以B的列，然后写入C.
该内核的有效带宽在NVIDIA Tesla K20X上仅为6.6GB / s（ECC关闭）。 为了分析性能，有必要考虑warp在for循环中如何访问全局内存。 每个线程的变形计算C的一个图块的一行，这取决于A的单个行和B的整个图块，如图10所示。
对于for循环的每个迭代i，warp中的线程读取B tile的一行，这是所有计算功能的顺序和合并访问。
但是，对于每次迭代i而言，warp中的所有线程都从矩阵A的全局内存中读取相同的值，因为索引行* TILE_DIM + i在warp内是不变的。 尽管在计算能力2.0的设备上这种访问只需要1次事务 或更高，因为事务中存在浪费的带宽，因为高速缓存行中的32个字中只有一个4字节的字被使用。 我们可以在循环的后续迭代中重用这个缓存行，我们最终将使用全部32个字; 然而，当许多warps同时在同一个多处理器上执行时，就像通常情况一样，高速缓存行可能容易从迭代i和i + 1之间的高速缓存中逐出。
