二. CUDA编程入门-共享内存以及Bank Conflict

本文主要是介绍二. CUDA编程入门-共享内存以及Bank Conflict，希望对大家解决编程问题提供一定的参考价值，需要的开发者们随着小编来一起学习吧！

前言

自动驾驶之心推出的 《CUDA与TensorRT部署实战课程》，链接。记录下个人学习笔记，仅供自己参考

Note：关于共享内存杜老师之前也讲过，感兴趣的可以看看 3.6.cuda运行时API-共享内存的学习

本次课程我们来学习课程第二章—CUDA 编程入门，一起来学习共享内存以及 Bank Conflict

课程大纲可以看下面的思维导图

0. 简述

这节课我们来讲第二章第 3 小节，共享内存以及Bank Conflict，这个部分我们主要分为以下两个部分讲解：

• shared memory
• bank conflict

1. shared memory

本小节目标：理解如何使用 shared memory，为什么使用 shared memory 会有加速效果以及在 shared memory 中使用动态/静态变量的注意事项

这个部分我们来看 shared memory 共享内存，共享内存是 CUDA 编程中比较重要的内容，大家在写核函数的时候如果能够把 shared memory 利用好，它其实给你程序带来的速度提升还是比较明显的，所以通过这个小节希望大家能够去理解如何使用 shared memory 以及 shared memory 加速主要体现在哪里，还有如何在 shared memory 中使用动态/静态变量，使用动态/静态变量的一些注意事项都有哪些

1.1 执行一下我们的第七个CUDA程序

源代码获取地址：https://github.com/kalfazed/tensorrt_starter

我们先来看这节课的案例代码 2.7-matmul-shared-memory，如下图所示：

这节课的案例代码相比于之前的主要多了 matmul_gpu_shared.cu 这个文件，它是一个利用共享内存实现矩阵乘法的 CUDA 程序，我们先来执行下这个案例，执行结果如下：

可以看到我们得到了四个结果，第一个是 gpu warmup 的耗时，第二个是 general 不使用共享内存进行矩阵乘法的耗时大概是 xxx ms，第三个是使用共享内存并且是使用静态变量进行矩阵乘法的耗时大概是 72.74 ms，我们可以看到它的一个加速比大概是比 general 快了 1/3 左右，最后一个是使用共享内存但是使用动态变量进行矩阵乘法的耗时大概是 93.31 ms，这个加速效果就没有那么明显了

我们进入到代码中看下它是怎么做的，首先是 main 函数其代码如下：

#include <stdio.h>
#include <cuda_runtime.h>#include "utils.hpp"
#include "timer.hpp"
#include "matmul.hpp"int seed;
int main(){Timer timer;int width     = 1<<12; // 4,096int low       = 0;int high      = 1;int size      = width * width;int blockSize = 16;bool statMem  = true;char str[100];float* h_matM = (float*)malloc(size * sizeof(float));float* h_matN = (float*)malloc(size * sizeof(float));float* h_matP = (float*)malloc(size * sizeof(float));float* d_matP = (float*)malloc(size * sizeof(float));// seed = (unsigned)time(NULL);seed = 1;initMatrix(h_matM, size, low, high, seed);seed += 1;initMatrix(h_matN, size, low, high, seed);LOG("Input size is %d x %d", width, width);/* GPU warmup */timer.start_gpu();MatmulOnDevice(h_matM, h_matN, h_matP, width, blockSize);timer.stop_gpu();timer.duration_gpu("matmul in gpu(warmup)");/* GPU general implementation <<<256, 16>>>*/timer.start_gpu();MatmulOnDevice(h_matM, h_matN, d_matP, width, blockSize);timer.stop_gpu();std::sprintf(str, "matmul in gpu(without shared memory)<<<%d, %d>>>", width / blockSize, blockSize);timer.duration_gpu(str);compareMat(h_matP, d_matP, size);// /* GPU general implementation <<<256, 16>>>*/timer.start_gpu();MatmulSharedOnDevice(h_matM, h_matN, d_matP, width, blockSize, statMem);timer.stop_gpu();std::sprintf(str, "matmul in gpu(with shared memory(static))<<<%d, %d>>>", width / blockSize, blockSize);timer.duration_gpu(str);compareMat(h_matP, d_matP, size);/* GPU general implementation <<<256, 16>>>*/statMem = false;timer.start_gpu();MatmulSharedOnDevice(h_matM, h_matN, d_matP, width, blockSize, statMem);timer.stop_gpu();std::sprintf(str, "matmul in gpu(with shared memory(dynamic))<<<%d, %d>>>", width / blockSize, blockSize);timer.duration_gpu(str);compareMat(h_matP, d_matP, size);return 0;
}

main 函数中修改的东西不多，为了让测速尽量公平我们让所有核函数的 grid size 和 block size 保持一致，主要对比不使用 shared memory、使用 shared memory 和静态变量以及使用 shared memory 和动态变量三种情况下的运行时间。此外为了实现更加精准的 kernel 核函数测速，这里我们采用了 event 进行标记，因为 CPU 上的 std::chrono 库测试会存在一些延迟

event 中文名字叫做事件，可以用来标记 cuda 的 stream 中某一个执行点，一般用在多个 stream 同步或者监听某个 stream 的执行，有关 stream 和 event 我们在之后的案例中再细讲，这里大家了解下即可

我们来看下 event 事件如何使用，首先我们在 timer.hpp 头文件中定义了两个 cuda event 成员变量 _gStart 和 _gStop，代码如下所示：

class Timer {
public:using s  = std::ratio<1, 1>;using ms = std::ratio<1, 1000>;using us = std::ratio<1, 1000000>;using ns = std::ratio<1, 1000000000>;public:Timer();~Timer();public:void start_cpu();void start_gpu();void stop_cpu();void stop_gpu();template <typename span>void duration_cpu(std::string msg);void duration_gpu(std::string msg);private:std::chrono::time_point<std::chrono::high_resolution_clock> _cStart;std::chrono::time_point<std::chrono::high_resolution_clock> _cStop;cudaEvent_t _gStart;cudaEvent_t _gStop;float _timeElasped;
};

之后我们在 timer.cpp 中利用 cudaEventRecord() 定义了两个函数，代码如下所示：

void Timer::start_gpu() {cudaEventRecord(_gStart, 0);
}void Timer::stop_gpu() {cudaEventRecord(_gStop, 0);
}void Timer::duration_gpu(std::string msg){CUDA_CHECK(cudaEventSynchronize(_gStart));CUDA_CHECK(cudaEventSynchronize(_gStop));cudaEventElapsedTime(&_timeElasped, _gStart, _gStop);// cudaDeviceSynchronize();// LAST_KERNEL_CHECK();LOG("%-60s uses %.6lf ms", msg.c_str(), _timeElasped);
}

当我们要测速的时候，我们计算这两个事件之间的事件差即可，通过调用 cudaEventElapsedTime 这个 API 就能完成，关于使用静态共享变量和动态共享变量的差异我们在后续代码分析中再跟大家去讲解

1.2 CUDA Core的矩阵乘法计算

OK，在我们进入到共享内存使用之前，我们先回顾下上节课讲的利用 CUDA Core 进行矩阵乘法计算，如下图所示：

A 为 4x8 的矩阵，B 为 8x4 的矩阵，C = A * B，我们说过 CUDA 中的每一个线程负责 C 中某一个元素的计算，比如上图中的 c(0,0)，C 矩阵中共有 16 个元素，那么我们可以同时启动 16 个线程一起去计算，这样就只需要 8 个时钟周期，这是我们上节课讲过的

现在我们考虑这么一个问题，比如说我们现在有 16 个线程，第一个线程它计算的是 c(0,0) 访问的是 A 矩阵中绿色的第一行和 B 矩阵中紫色的第一列，我们继续看我们下一个线程计算的是 c(0,1)，它访问的是什么呢？

可以看到它访问的依旧是 A 矩阵中绿色的第一行，数据没有变，地址没有变，而 B 矩阵则访问的是紫色的第二列，稍微变了一下，我们继续看 c(0,2) 的计算，如下图所示：

c(0,2) 访问 A 矩阵的数据依旧没有改变，B 矩阵的数据则是变成了第三列

同理 c(0,3) 也是矩阵 A 的数据没有改变，矩阵 B 的数据变成了另外一列，所以我们就会发现 c(0,0) 到 c(0,3) 线程的计算中它们去矩阵 A 取数据的时候，其实都是从 global memory 中反复的取同一个数据，这个的话其实效率不高，那既然每一次都是去同一个地方取相同的数据，我们就想只取一次然后反复利用不就可以提高效率吗

看完了 C 矩阵的一行元素在计算时的共性特点后，我们再来看 C 矩阵的一列元素在计算时是不是也有类似的特点呢

我们可以发现 c(1,0) 计算时 B 矩阵的那一列的数据是没有改变的，只是 A 的那一行数据进行了改变，同理 c(2,0)、c(3,0) 也是一样的，如下图所示：

所以我们在利用 CUDA Core 进行矩阵乘法计算的过程中会发现有很多冗余的操作，比如同样的地址空间，我们要从 global memory 中一次次的访问

那么从上面的图我们也能看到 global memory 它其实距离我们的计算单元还是比较远的，这意味着我们在访问它里面的一些数据时并不是非常高效，因为我们都知道你的内存距离计算单元越远，它访问的速度就越慢。所以我们自然而然的就会去想，既然是多次访问同样的数据，那我们是不是可以利用距离计算单元更近的 shared memory 共享内存呢

上面是 NVIDIA Ampere 架构中 SM 的框架图，可以看到 SM 中有 L1 Instruction Cache、L1 Data Cache 以及 shared memory，我们从图中可以看到 NVIDIA 的 SM 框架设计中将 L1 Data Cache 和 shared memory 给放到一起了。再深入看，图中共有四个 Warp Scheduler，每个 Warp Scheduler 调度 32 个线程，这 32 个线程可以调用我们刚才说的 L1 Instruction Cache、L1 Data Cache/Shared Memroy 这些资源

我们的内存从广义上可以分为以下两种：

• on-chip memory：片上内存
• off-chip memory：片外内存（DRAM）

L1/L2 cache 和 shared memory 都是属于 on-chip memory，memory load/store 的 overhead 会比较小，是可以高速访问的 memory，L1 cache 相比 L2 cache 更快，它的延迟大概会低 20~30 倍左右，带宽大概是 10 倍左右，而 global memory 的延迟是最高的，它属于 off-chip memory 也叫 DRAM，我们一般在 cudaMalloc 时都是在 global memory 上进行访问的

这里再稍微扩展下，shared memory 的大小一般是 48KB 左右，但这个不是一个固定值我们是可以进行调整的，我们可以利用 cuda runtime API 把你的 shared memory 大小和 L1 Cache 大小进行一个调整，比如你的程序它是优先使用 shared memory 的话，我们可以把 shared 大小从 48KB 变成 64KB，相应的我们 L1 Cache 就可能会小一点。

此外 shared memory 叫做共享内存还是有原因的，每一个 block 线程块里面都有一块 shared memory，一个 block 内部所有的 thread 线程它其实都是共享这同一块 shared memory 的

不同架构的不同memory的bandwidth的不同

上图是从 Is Data Placement Optimization Still Relevant On Newer GPUs 这篇论文中截取的，它这里对比了 Kepler、Maxwell、Pascal、Volta 四个 GPU 架构，纵轴代表的是带宽，横轴是不同的内存比如 L1、shared memory、L2 等等，越往右走我们内存的带宽越小，我们可以发现 Global memory 和前面的 L1、shared memory 相比带宽差异非常明显，所以我们如果能少用就尽量少用 global memory 而更多的使用 L1、shared memory 这些东西

最后我们来简单对比下上面提到的这些内容：

• register：线程独享，访问最快，大小最小，on-chip
• L1 cache：SM 内共享，on-chip
• shared memory：SM 内共享，on-chip
• L2 cache：SM 间共享，on-chip
• local memory：线程独享，寄存器不足的时候使用，off-chip
• global memory：设备内所有线程共享，off-chip
• constant memory：只读内存，用于避免线程读数据冲突，off-chip
• texture memory：只读内存，可以实现硬件插值，off-chip

1.3 代码分析

在之前的分析中我们知道矩阵乘法的计算过程中我们经常要去同一块内存中访问同一批数据，既然如此我们是不是可以把这些数据提前给放到 shared memory 中做一个缓存呢

如上图所示，我们一开始可以把 A 矩阵中绿色的 16 个元素和 B 矩阵中紫色的 16 个元素先给放到 shared memory 访存中，让它们先做一次计算，全部计算做完不需要访问它们的时候，我们再做下一块的计算，如下图所示：

依次类推，将计算完的结果进行累加放到 C 中就行了，这个就是使用 shared memory 来做矩阵乘法的一个方法，下面我们来详细分析下具体的实现代码

我们先从 main.cpp 开始，代码如下所示：

#include <stdio.h>
#include <cuda_runtime.h>#include "utils.hpp"
#include "timer.hpp"
#include "matmul.hpp"int seed;
int main(){Timer timer;int width     = 1<<12; // 4,096int low       = 0;int high      = 1;int size      = width * width;int blockSize = 16;bool statMem  = true;char str[100];float* h_matM = (float*)malloc(size * sizeof(float));float* h_matN = (float*)malloc(size * sizeof(float));float* h_matP = (float*)malloc(size * sizeof(float));float* d_matP = (float*)malloc(size * sizeof(float));// seed = (unsigned)time(NULL);seed = 1;initMatrix(h_matM, size, low, high, seed);seed += 1;initMatrix(h_matN, size, low, high, seed);LOG("Input size is %d x %d", width, width);/* GPU warmup */timer.start_gpu();MatmulOnDevice(h_matM, h_matN, h_matP, width, blockSize);timer.stop_gpu();timer.duration_gpu("matmul in gpu(warmup)");/* GPU general implementation <<<256, 16>>>*/timer.start_gpu();MatmulOnDevice(h_matM, h_matN, d_matP, width, blockSize);timer.stop_gpu();std::sprintf(str, "matmul in gpu(without shared memory)<<<%d, %d>>>", width / blockSize, blockSize);timer.duration_gpu(str);compareMat(h_matP, d_matP, size);// /* GPU general implementation <<<256, 16>>>*/timer.start_gpu();MatmulSharedOnDevice(h_matM, h_matN, d_matP, width, blockSize, statMem);timer.stop_gpu();std::sprintf(str, "matmul in gpu(with shared memory(static))<<<%d, %d>>>", width / blockSize, blockSize);timer.duration_gpu(str);compareMat(h_matP, d_matP, size);/* GPU general implementation <<<256, 16>>>*/statMem = false;timer.start_gpu();MatmulSharedOnDevice(h_matM, h_matN, d_matP, width, blockSize, statMem);timer.stop_gpu();std::sprintf(str, "matmul in gpu(with shared memory(dynamic))<<<%d, %d>>>", width / blockSize, blockSize);timer.duration_gpu(str);compareMat(h_matP, d_matP, size);return 0;
}

main.cpp 中的内容和之前矩阵乘法的内容差不多，包括矩阵初始化，warmup 等等，main 中主要是比较了使用共享内存与不使用共享内存在性能上的差异，通过调用三个核函数来进行对比，分别是：

• 不使用共享内存的矩阵乘法的核函数
• 使用静态共享内存的矩阵乘法的核函数
• 使用动态共享内存的矩阵乘法的核函数

下面我们来看具体的实现代码 matmul_gpu_shared.cu，如下所示：

#include "cuda_runtime_api.h"
#include "utils.hpp"#define BLOCKSIZE 16/* 使用shared memory把计算一个tile所需要的数据分块存储到访问速度快的memory中
*/
__global__ void MatmulSharedStaticKernel(float *M_device, float *N_device, float *P_device, int width){__shared__ float M_deviceShared[BLOCKSIZE][BLOCKSIZE];__shared__ float N_deviceShared[BLOCKSIZE][BLOCKSIZE];/* 对于x和y, 根据blockID, tile大小和threadID进行索引*/int x = blockIdx.x * blockDim.x + threadIdx.x;int y = blockIdx.y * blockDim.y + threadIdx.y;float P_element = 0.0;int ty = threadIdx.y;int tx = threadIdx.x;/* 对于每一个P的元素，我们只需要循环遍历width / tile_width 次就okay了，这里有点绕，画图理解一下*/for (int m = 0; m < width / BLOCKSIZE; m ++) {M_deviceShared[ty][tx] = M_device[y * width + (m * BLOCKSIZE + tx)];N_deviceShared[ty][tx] = N_device[(m * BLOCKSIZE + ty)* width + x];__syncthreads();for (int k = 0; k < BLOCKSIZE; k ++) {P_element += M_deviceShared[ty][k] * N_deviceShared[k][tx];}__syncthreads();}P_device[y * width + x] = P_element;
}__global__ void MatmulSharedDynamicKernel(float *M_device, float *N_device, float *P_device, int width, int blockSize){/* 声明动态共享变量的时候需要加extern，同时需要是一维的 注意这里有个坑, 不能够像这样定义： __shared__ float M_deviceShared[];__shared__ float N_deviceShared[];因为在cuda中定义动态共享变量的话，无论定义多少个他们的地址都是一样的。所以如果想要像上面这样使用的话，需要用两个指针分别指向shared memory的不同位置才行*/extern __shared__ float deviceShared[];int stride = blockSize * blockSize;/* 对于x和y, 根据blockID, tile大小和threadID进行索引*/int x = blockIdx.x * blockSize + threadIdx.x;int y = blockIdx.y * blockSize + threadIdx.y;float P_element = 0.0;int ty = threadIdx.y;int tx = threadIdx.x;/* 对于每一个P的元素，我们只需要循环遍历width / tile_width 次就okay了 */for (int m = 0; m < width / blockSize; m ++) {deviceShared[ty * blockSize + tx] = M_device[y * width + (m * blockSize + tx)];deviceShared[stride + (ty * blockSize + tx)] = N_device[(m * blockSize + ty)* width + x];__syncthreads();for (int k = 0; k < blockSize; k ++) {P_element += deviceShared[ty * blockSize + k] * deviceShared[stride + (k * blockSize + tx)];}__syncthreads();}if (y < width && x < width) {P_device[y * width + x] = P_element;}
}/*使用Tiling技术一个tile处理的就是block, 将一个矩阵分为多个小的tile，这些tile之间的执行独立，并且可以并行
*/
void MatmulSharedOnDevice(float *M_host, float *N_host, float* P_host, int width, int blockSize, bool staticMem){/* 设置矩阵大小 */int size = width * width * sizeof(float);long int sMemSize = blockSize * blockSize * sizeof(float) * 2;/* 分配M, N在GPU上的空间*/float *M_device;float *N_device;CUDA_CHECK(cudaMalloc((void**)&M_device, size));CUDA_CHECK(cudaMalloc((void**)&N_device, size));/* 分配M, N拷贝到GPU上*/CUDA_CHECK(cudaMemcpy(M_device, M_host, size, cudaMemcpyHostToDevice));CUDA_CHECK(cudaMemcpy(N_device, N_host, size, cudaMemcpyHostToDevice));/* 分配P在GPU上的空间*/float *P_device;CUDA_CHECK(cudaMalloc((void**)&P_device, size));;/* 调用kernel来进行matmul计算, 在这个例子中我们用的方案是：使用一个grid，一个grid里有width*width个线程 */dim3 dimBlock(blockSize, blockSize);dim3 dimGrid(width / blockSize, width / blockSize);if (staticMem) {MatmulSharedStaticKernel <<<dimGrid, dimBlock>>> (M_device, N_device, P_device, width);} else {MatmulSharedDynamicKernel <<<dimGrid, dimBlock, sMemSize, nullptr>>> (M_device, N_device, P_device, width, blockSize);}/* 将结果从device拷贝回host*/CUDA_CHECK(cudaMemcpy(P_host, P_device, size, cudaMemcpyDeviceToHost));CUDA_CHECK(cudaDeviceSynchronize());/* 注意要在synchronization结束之后排查kernel的错误 */LAST_KERNEL_CHECK(); /* Free */cudaFree(P_device);cudaFree(N_device);cudaFree(M_device);
}

这段代码是 CUDA 程序中用于加速矩阵乘法的一个高级示例，通过利用 shared memory（共享内存）来提高内存访问效率。在 CUDA 编程中，shared memory 是一种位于每个 block 内部的内存类型，它允许同一 block 内的线程快速共享数据。与访问全局内存相比，访问 shared memory 的延迟更低，因此，合理使用 shared memory 可以显著提高程序的性能。这个示例展示了两种方法来实现矩阵乘法的加速：一种使用静态共享内存，另一种使用动态共享内存。(from chatGPT)

静态共享内存核函数 MatmulSharedStaticKernel

这个核函数使用静态分配的共享内存数组 M_deviceShared 和 N_deviceShared 来存储每个 block 需要的 M 和 N 矩阵的一部分（称为 tile）。这种方法称为 tiling 技术，目的是减少对全局内存的访问次数，因为每个元素被多次用于计算。

• 索引计算：使用 blockIdx、blockDim 和 threadIdx 计算出每个线程在全局和 tile 中的位置。
• 数据载入：每个线程加载 M 和 N 矩阵中相应的元素到共享内存中，然后等待所有线程完成加载（通过 __syncthreads() 同步）。
• 计算：在共享内存中完成 tile 内的矩阵乘法，并累加结果到 P_element 中。完成每个 tile 的计算后，再次同步线程以确保所有线程都完成了计算。
• 写回结果：将计算结果写回到全局内存中的输出矩阵 P。

动态共享内存核函数 MatmulSharedDynamicKernel

与静态共享内存版本不同，这个核函数使用动态共享内存来存储 tile 数据。动态共享内存的大小在 kernel 调用时指定，允许更灵活地控制内存使用。

• 动态共享内存声明：通过 extern __shared__ 声明一个未指定大小的共享内存数组 deviceShared。实际大小在 kernel 调用时通过第三个参数确定。
• 索引计算和数据载入：与静态版本相似，但使用动态计算的偏移来访问共享内存中的数据。
• 计算和写回结果：计算过程与静态共享内存版本类似，但注意到共享内存的布局需要根据动态指定的 blockSize 进行调整。

主机函数 MatmulSharedOnDevice

这个函数封装了 GPU 上的矩阵乘法计算过程，包括内存分配、数据传输和核函数调用。它根据传入的参数 staticMem 决定是调用静态共享内存版本还是动态共享内存版本的核函数。

• 内存分配和数据传输：为输入矩阵 M 和 N，以及输出矩阵 P 在 GPU 上分配内存，并将数据从主机复制到 GPU。
• 核函数调用：根据 staticMem 的值选择合适的核函数进行调用。对于动态共享内存版本，需要额外指定共享内存的大小。
• 结果回传：将计算结果从 GPU 内存复制回主机内存。

总结

通过使用 shared memory 来缓存每个 block 所需的矩阵部分，这个示例程序大幅减少了对全局内存的访问次数，从而提高了矩阵乘法的计算效率。动态和静态共享内存的使用展示了 CUDA 编程中灵活处理内存的能力，对于优化高性能计算应用至关重要。

OK，以上就是对 2.7-matmul-shared-memory 案例代码的分析，最后我们执行下看下输出结果，如下图所示：

1.4 static vs. dynamic

这个部分为补充内容，是博主询问 chatGPT 有关静态共享内存和动态共享内存之间的区别以及 CUDA 编程中使用推荐(from chatGPT)

在 CUDA 编程中，共享内存是一种高速但容量有限的内存，位于每个线程块（Block）中，允许该块内的所有线程访问。共享内存的使用可以显著提高内存访问效率，尤其是在需要多个线程读取相同数据时。共享内存分为两种类型：静态共享内存和动态共享内存。

静态共享内存

静态共享内存在编译时就确定了大小，通过在 CUDA 核函数内部使用 __shared__ 关键字和静态数组声明来分配。静态共享内存的大小必须在编译时已知，这意味着它的大小不能基于运行时的数据计算得出。

__global__ void kernelExample() {__shared__ float staticSharedMem[256];// 使用 staticSharedMem...
}

动态共享内存

与静态共享内存不同，动态共享内存的大小在运行时指定，允许根据不同的需求为每个线程块分配不同大小的共享内存。动态共享内存通过在核函数的执行配置中使用第三个参数来指定。

extern __shared__ float dynamicSharedMem[];

动态共享内存在核函数调用时通过额外参数指定大小：

kernelExample<<<numBlocks, blockSize, sharedMemSize>>>(...);

静态共享内存与动态共享内存的区别

• 声明方式：静态共享内存在编译时声明，大小固定；动态共享内存在运行时指定大小。
• 灵活性：动态共享内存的大小可以根据需要在运行时调整，提供了更大的灵活性。
• 使用简便性：静态共享内存由于在编译时就已确定，使用起来更简单直观。
• 性能差异：理论上，静态共享内存和动态共享内存的性能应该是相近的，因为它们访问的是相同的物理内存。然而，动态共享内存需要在核函数调用时动态分配，这可能会导致微小的性能开销，特别是在频繁启动核函数的情况下。

为什么静态共享内存表现更好？

静态共享内存的性能通常更稳定且略优于动态共享内存，原因可能包括：

1. 开销较低：静态共享内存在编译时就已分配，无需运行时处理额外的分配过程。
2. 编译器优化：编译器可能能够更好地优化静态共享内存的使用，因为其大小在编译时已知。

动态共享内存的不稳定性

动态共享内存性能不稳定可能的原因包括：

1. 运行时分配：动态共享内存的分配在运行时进行，可能受到当前 GPU 状态的影响，如其他并发执行的任务。
2. 大小变化：如果动态共享内存的大小在不同的核函数调用之间变化，这可能影响性能，因为每次调整大小都可能引入额外的开销。

推荐使用哪个？

在 CUDA 编程中，推荐根据具体需求选择共享内存的类型：

• 如果共享内存的大小在编译时已知且不变，推荐使用静态共享内存，因为它简单且效率高。
• 如果需要根据不同的输入动态调整共享内存的大小，或者同一程序的不同部分需要不

同大小的共享内存，那么动态共享内存是更灵活的选择。

总的来说，静态共享内存由于其简单性和潜在的性能优势，通常是首选。但在需要额外灵活性的场合，动态共享内存则提供了必要的调整能力。

2. bank conflict

本小节目标：理解在 shared memory 中的 bank 是什么，什么时候会发生 bank conflict，以及如何减缓 bank conflict

这部分给大家讲 shared memory 中另一个重要的概念叫做 bank conflict（存储体冲突），bank（存储体）是 shared memory 中的一个概念，我们如果说想要去利用好 shared memory 的话，其实我们也需要去理解我们的数据在 shared memory 中是以什么样的方式去存储的，所以我们就得理解 bank 在 shared memory 中是什么样子的，以及什么时候会产生 bank conflict，如果产生了 bank conflict 我们应该怎么去缓解它，这是我们这个小节需要解决的问题。

2.1 执行一下我们的第八个CUDA程序

这个小节对应的案例是 2.8-bank-conflict

相比于之前的案例多了 matmul_gpu_bank_conflict.cu 和 matmul_gpu_bank_conflict_pad.cu 两个文件，make 执行之后的输出结果如下图所示：

输出结果显示了使用 shared memory 中没有 bank conflict 和有 bank conflict 时的速度差异，我们发现无论是静态共享内存还是动态共享内存，只要发生了 bank conflict 矩阵乘法计算的时间就会变得特别长，比普通的正常使用 global memory 进行矩阵乘法的时间还要长，这个是一个非常不好的现象

那么好，如果我们发生了存储体冲突 bank conflict 我们怎么办呢？我们就要用 pad 这个方法去解决它这个我们后面再讲，使用 pad 解决之后可以看到对于静态共享内存的 bank conflict 而言时间从 153ms 降低到了 114ms，速度还是很慢，动态共享内存的 bank conflict 而言时间从 215ms 降低到了 114ms，也就是我们使用了 pad 之后能够将 bank conflict 给减轻

我们这个程序其实跟之前的矩阵乘法几乎差不多，只不过我们把它的一个索引的方式给改了，在 matmul 的过程中，将矩阵的遍历方式从行优先转为列优先，从而人为的让 matmul 的过程发生 bank conflict。然后在 shared memory 申请内存时人为的进行 padding，从而防止 bank conflict 发生

matmul_gpu_bank_conflict.cu(左:静态,右:动态)

matmul_gpu_bank_conflict_pad.cu(左:静态,右:动态)

2.2 shared memory中存放数据的特殊方式

我们先来回顾一下 shared memory 中存放数据的方式是什么样子的，我们知道在 CUDA 编程中，32 个 threads 组成一个 warp，一般程序在执行的时候是以 warp 为单位去执行的，也就是说每 32 个 threads 一起执行统一指令，比如同时读/取数据，如下图所示：

NVIDIA 硬件设计者为了能够比较高效的访问我们的 shared memory 把它也给分成了 32 个不同的部分，我们称之为 *bank，分别对应 warp 中 32 个线程，之后让每一个线程去访问它们其中的一个部分，如下图所示：

*空间有限，上图以16个bank为例

我们来看 bank 中一个比较常见的图：

一个 warp 是 32 个线程，对应着 shared memory 中的 32 个 bank，也就是上图中 Bank0-Bank31，值得注意的是 bank 的宽度代表的是一个 bank 所存储的数据的大小宽度

• 可以是 4 个字节（32bit，单精度浮点数）
• 也可以是 8 个字节（64bit，双精度浮点数）

值得注意的是我们 bank 有一个特点即每 31 个 bank，就会进行一次 stride，比如说 bank 的宽度是 4 字节，我们在 shared memory 中申请了 float A[256] 大小的空间，那么 A[0], A[1], …, A[31] 分别在 bank0, bank1, …, bank31 中，A[32], A[33], …, A[63] 也分在了 bank0, bank1, …, bank31 中，所以 A[0], A[32] 是共享同一个 bank 的

2.3 bank conflict

一个很理想的情况就是 32 个 thread 分别访问 shared memory 中的 32 个不同的 bank，没有 bank conflict，一个 memory 周期完成所有的 memory read/write（row major/行优先矩阵访问）

但是我们想如果说我们不用行优先去访问，而是采用列优先访问一个矩阵会发生什么情况呢？如下图所示：

可以看到我们所有的 thread 它们都在访问同一个 bank 即 bank0，我们只不过在访问 bank0 中不同的地址空间，一个访问 A[0]，一个访问 A[32]，一个访问 A[64] 等等，那这样的访问效率就会很慢。

上图就是一个最不理想的情况，32 个 thread 访问 shared memory 中的同一个 bank，bank conflict 最大化，需要 32 个 memory 周期才能完成所有的 memory read/write（column major/列优先矩阵访问）

不知道大家有没有发现这个情况其实还是比较常见的，一般出现在哪里呢，大家可以想一想，这个不就是矩阵的一个转置吗，矩阵转置的时候其实就经常会发生，要不然就是列优先访问，要不就是行优先访问，总之它很容易产生这个 bank conflict

2.4 使用padding缓解bank conflict

如果产生 bank conflict 的话我们怎么去解决它呢，我们以下面这张图为例来讲解：

为了方便解释，这里使用 8 个 bank 一次 stride 进行举例，在实际 CUDA 设计中，依然是 32 个 bank 一次 stride。上图中我们所有 thread 都一起集中去访问 Bank2 导致 bank conflict，我们要去解决它，我们怎么做呢，我们在申请 shared memory 的时候多添加一列，那再看下此时会发生什么

我们上面说了在这个例子中 shared memory 中 8 个 bank 进行一个 stride（实际是 32），那我们多申请了一列导致内部的 bank 的布局发生了改变，如上图所示，之前是 8x8 的矩阵排列现在布局改变后变成了 9x8 的矩阵排列

并且我们发现一个非常神奇的现象，之前我们说所有的 thread 都在访问 Bank2 导致 bank conflict 冲突了，现在布局改变后我们发现每一个 thread 它在访问 Bank2 的时候其实不会产生任何冲突的，同样的如果所有的 thread 都访问 Bank3 也会冲突。

那我们之前行访问 Bank0-Bank7 没有冲突，布局改变后会不会产生冲突呢？可以从上图中看到如果按照行访问的方式也是没有冲突的，这个就是很神奇的地方，所以大家记住如果想要解决 bank conflict 的话在申请 shared memory 的时候多添加一列就好了。

下面我们去代码里面看看

2.5 代码分析

2.8-bank-conflict 案例代码和之前 2.7 小节差不多，其中 matmul_gpu_bank_conflict.cu 是手动按列优先访问使其产生 bank conflict，而 matmul_gpu_bank_conflict_pad.cu 则是利用 pad 解决出现的 bank conflict，因此我们重点来分析它就行

代码如下所示：

#include "cuda_runtime_api.h"
#include "utils.hpp"#define BLOCKSIZE 16/* 使用shared memory把计算一个tile所需要的数据分块存储到访问速度快的memory中
*/
__global__ void MatmulSharedStaticConflictPadKernel(float *M_device, float *N_device, float *P_device, int width){/* 添加一个padding，可以防止bank conflict发生，结合图理解一下*/__shared__ float M_deviceShared[BLOCKSIZE][BLOCKSIZE + 1];__shared__ float N_deviceShared[BLOCKSIZE][BLOCKSIZE + 1];/* 对于x和y, 根据blockID, tile大小和threadID进行索引*/int x = blockIdx.x * BLOCKSIZE + threadIdx.x;int y = blockIdx.y * BLOCKSIZE + threadIdx.y;float P_element = 0.0;int ty = threadIdx.y;int tx = threadIdx.x;/* 对于每一个P的元素，我们只需要循环遍历width / tile_width 次就okay了，这里有点绕，画图理解一下*/for (int m = 0; m < width / BLOCKSIZE; m ++) {/* 这里为了实现bank conflict, 把tx与tx的顺序颠倒，同时索引也改变了*/M_deviceShared[tx][ty] = M_device[x * width + (m * BLOCKSIZE + ty)];N_deviceShared[tx][ty] = M_device[(m * BLOCKSIZE + tx)* width + y];__syncthreads();for (int k = 0; k < BLOCKSIZE; k ++) {P_element += M_deviceShared[tx][k] * N_deviceShared[k][ty];}__syncthreads();}/* 列优先 */P_device[x * width + y] = P_element;
}__global__ void MatmulSharedDynamicConflictPadKernel(float *M_device, float *N_device, float *P_device, int width, int blockSize){/* 声明动态共享变量的时候需要加extern，同时需要是一维的 注意这里有个坑, 不能够像这样定义： __shared__ float M_deviceShared[];__shared__ float N_deviceShared[];因为在cuda中定义动态共享变量的话，无论定义多少个他们的地址都是一样的。所以如果想要像上面这样使用的话，需要用两个指针分别指向shared memory的不同位置才行*/extern __shared__ float deviceShared[];int stride = (blockSize + 1) * blockSize;/* 对于x和y, 根据blockID, tile大小和threadID进行索引*/int x = blockIdx.x * blockSize + threadIdx.x;int y = blockIdx.y * blockSize + threadIdx.y;float P_element = 0.0;int ty = threadIdx.y;int tx = threadIdx.x;/* 对于每一个P的元素，我们只需要循环遍历width / tile_width 次就okay了 */for (int m = 0; m < width / blockSize; m ++) {/* 这里为了实现bank conflict, 把tx与tx的顺序颠倒，同时索引也改变了*/deviceShared[tx * (blockSize + 1) + ty]             = M_device[x * width + (m * blockSize + ty)];deviceShared[stride + (tx * (blockSize + 1) + ty)]  = N_device[(m * blockSize + tx) * width + y];__syncthreads();for (int k = 0; k < blockSize; k ++) {P_element += deviceShared[tx * (blockSize + 1) + k] * deviceShared[stride + (k * (blockSize + 1 ) + ty)];}__syncthreads();}/* 列优先 */P_device[x * width + y] = P_element;
}/*使用Tiling技术一个tile处理的就是block, 将一个矩阵分为多个小的tile，这些tile之间的执行独立，并且可以并行
*/
void MatmulSharedConflictPadOnDevice(float *M_host, float *N_host, float* P_host, int width, int blockSize, bool staticMem){/* 设置矩阵大小 */int size = width * width * sizeof(float);long int sMemSize = (blockSize + 1) * blockSize * sizeof(float) * 2;/* 分配M, N在GPU上的空间*/float *M_device;float *N_device;CUDA_CHECK(cudaMalloc((void**)&M_device, size));CUDA_CHECK(cudaMalloc((void**)&N_device, size));/* 分配M, N拷贝到GPU上*/CUDA_CHECK(cudaMemcpy(M_device, M_host, size, cudaMemcpyHostToDevice));CUDA_CHECK(cudaMemcpy(N_device, N_host, size, cudaMemcpyHostToDevice));/* 分配P在GPU上的空间*/float *P_device;CUDA_CHECK(cudaMalloc((void**)&P_device, size));;/* 调用kernel来进行matmul计算, 在这个例子中我们用的方案是：使用一个grid，一个grid里有width*width个线程 */dim3 dimBlock(blockSize, blockSize);dim3 dimGrid(width / blockSize, width / blockSize);if (staticMem) {MatmulSharedStaticConflictPadKernel <<<dimGrid, dimBlock>>> (M_device, N_device, P_device, width);} else {MatmulSharedDynamicConflictPadKernel <<<dimGrid, dimBlock, sMemSize, nullptr>>> (M_device, N_device, P_device, width, blockSize);}/* 将结果从device拷贝回host*/CUDA_CHECK(cudaMemcpy(P_host, P_device, size, cudaMemcpyDeviceToHost));CUDA_CHECK(cudaDeviceSynchronize());/* 注意要在synchronization结束之后排查kernel的错误 */LAST_KERNEL_CHECK(); /* Free */cudaFree(P_device);cudaFree(N_device);cudaFree(M_device);
}

这段代码演示了在 CUDA 程序中使用共享内存进行矩阵乘法计算，并采用 padding 技术来避免 bank conflict。代码中包含了两个核函数：一个使用静态共享内存，另一个使用动态共享内存。此外，还有一个主机函数负责管理内存分配、数据传输以及核函数的调用。(from chatGPT)

bank conflict 和 padding 技术

CUDA 的共享内存被分成了多个 bank，每个 bank 可以在一个周期内独立地服务一个内存请求。当两个（或多个）线程访问同一 bank 的不同地址时，如果这些访问不能合并，就会发生 bank conflict，导致访问延迟。为了解决或减少这种冲突，可以采用 padding 技术，即在共享内存数组中添加额外的列，使得实际使用的数据不会落在同一 bank 的冲突地址上。

核函数分析

MatmulSharedStaticConflictPadKernel

这个核函数使用静态共享内存来存储每个 tile 的 M 和 N 矩阵片段，每个片段都添加了一个额外的列作为 padding。通过这种方式，当线程组内的线程并行访问共享内存时，可以避免 bank conflict。

• 共享内存分配：为 M 和 N 的每个 tile 分配了 BLOCKSIZE x (BLOCKSIZE + 1) 的共享内存空间。
• 数据载入和计算：每个线程负责将 M 和 N 的一个元素载入到共享内存中，然后计算 P 的一个元素。为了模拟 bank conflict 并演示 padding 的效果，数据载入时 tx 和 ty 的顺序被颠倒，且在共享内存中使用了 [tx][ty] 的索引方式。

MatmulSharedDynamicConflictPadKernel

与静态共享内存的核函数相比，这个核函数使用动态分配的共享内存来存储 tile 数据，并同样采用了 padding 技术来避免 bank conflict。

• 动态共享内存分配：共享内存的实际大小在 kernel 调用时通过第三个参数确定，以 (blockSize + 1) * blockSize 的形式计算，为 M 和 N 的每部分分配了足够的空间，并预留了 padding。
• 数据载入和计算：类似于静态版本，但是通过动态计算的索引来访问共享内存中的数据。

主机函数 MatmulSharedConflictPadOnDevice

这个函数负责在主机端准备数据，调用相应的核函数，并处理数据的回传。它根据 staticMem 参数的值选择调用静态或动态共享内存的核函数。

• 内存分配和数据传输：为 M、N 和 P 矩阵在 GPU 上分配内存，并将 M 和 N 的数据从主机复制到 GPU。
• 核函数调用：根据 staticMem 的值选择合适的核函数进行计算，对于动态共享内存版本，还需要指定共享内存的大小。
• 结果回传和资源释放：计算完成后，将结果从 GPU 内存复制回主机内存，并释放 GPU 上分配的资源。

总结

通过添加 padding 来避免共享内存的 bank conflict 是优化 CUDA 程序性能的有效方法之一。这段代码展示了如何在实现矩阵乘法时，通过静态和动态共享内存以及 padding 技术来减少 bank conflict，从而提高并行计算的效率。

OK，以上就是对 2.8-bank-conflict 案例代码的分析，最后我们执行下看下输出结果，如下图所示：

总结

本次课程我们学习了 shared memory 共享内存，并利用它来加速矩阵乘法的计算，shared memory 位于 block 内部，相比于 global memory 而言访问速度更快、延迟更低。此外我们还分析了使用 shared memory 进行矩阵乘法的计算时可能出现 bank conflict 的问题，也就是多个线程访问同一 bank 的不同地址时造成的访问延迟，为了解决或减少这种冲突，我们采用 padding 技术在 shared memory 数组中添加额外的列让 bank 的内存布局发生改变从而缓解 bank conflict

OK，以上就是第 3 小节有关共享内存以及 Bank Conflict 的全部内容了，下节我们来学习 CUDA 编程中的 Stream 与 Event，敬请期待😄

参考

• NVIDIA Ampere Architecture In-Depth
• Is Data Placement Optimization Still Relevant On Newer GPUs
• CUDA – Streaming Multiprocessors
• avoiding bank conflicts in shared memory
• Shared Memory Banks and Conflicts
• 3.6.cuda运行时API-共享内存的学习

这篇关于二. CUDA编程入门-共享内存以及Bank Conflict的文章就介绍到这儿，希望我们推荐的文章对编程师们有所帮助！

---