CUDA 性能优化 | 共享内存机制 / 向量化访存策略

注：本文为“CUDA 性能优化”相关文章合辑。

图片清晰度受引文原图所限。
重传部分 CSDN 转储失败图片。
略作重排，未整理去重。
如有内容异常，请看原文。

Shared Memory 上的广播机制和 Bank Conflict 到底是怎么回事？

发表于 2023 年 12 月 26 日

  • 🔥2024.03.24 添加了关于对 128 Bit 下的 Bank Conflict 的讨论
  • 🔥2024.04.01 修正了 128 Bit 下第五个例子的错误代码，感谢知乎用户 @Alan小分享 . 指证～

NVIDIA GPU 上的内存结构从慢到快分为 Global Memory、L2 缓存、L1TEX 缓存/Shared Memory 和寄存器，从 Volta/Turning 开始其中 L1TEX 缓存和 Shared Memory 在物理上放在了同一块芯片上，拥有相似的延时和带宽[1]，因此如何掌握 Shared Memory 对性能而言就变得尤为重要。可惜的是，NVIDIA 官方的 CUDA 编程手册中介绍 Shared Memory 的教程其实只介绍了在每个线程访问一个 4 字节（32 位宽）的元素时 Bank Conflict 和广播机制，但对使用常用的向量化访存指令如 LDS.64 或 LDS.128 这种一次能访问 8 个字节（64 位宽）或 16 个字节（128 位宽）元素的情况却鲜有资料讨论。这篇文章大概就是想结合网络上的一些讨论以及通过 Microbenchmark 对这些细节来一探究竟。需要注意的是这篇文章的结论仅在 Turing 架构的 GPU 上验证过，其他架构的 GPU 可能会产生变化（欢迎评论区交流）。

Shared Memory 模型

我们可以把 Shared Memory 它看作是一个长度为 NN 的数组，每个数组元素的大小是 4 字节（比如可以是一个 int 或 float），这个数组对于同一个 Thread Block 中的所有线程都是可见的，但不同 Thread Block 之间的 Shared Memory 不能互相访问。其中 Shared Memory 最值得注意的点机制是其本身被划分称了 32 个 Bank，其中数组中第 ii 个元素存储在了第 i \bmod 32imod32 个 Bank 上。Shared Memory 访存机制可以总结为如下两条：

  • 如果一个 Warp 中的两个不同的线程同时访问了同一个元素，那么就会触发广播机制，即可以合并成对 Shared Memory 的一次访问；
  • 如果一个 Warp 中两个不同的线程同时访问了同一个 Bank 中的不同元素，那么就会产生Bank Conflict，可以理解成一个 Bank 同一时间只能产生吞吐一个元素，因此这两个线程的访存请求会被串行，因而会影响性能；

在 Nsight Compute 上，我们可以通过 Shared Memory 上的 Wavefront 数目来理解 Shared Memory 访存性能，Wavefront 越多说明访存需要的时间越长。

下面几张图片举了几个例子方便理解：

在这里插入图片描述
只会触发广播机制，没有 Bank Conflict，需要 1 个 Wavefront

在这里插入图片描述
不会触发广播机制，没有 Bank Conflict，需要 1 个 Wavefront

在这里插入图片描述
既会触发广播机制，也有 Bank Conflict，需要 4 个 Wavefront（注意第 18 个 Bank）

向量化访存指令

前面在讨论 Shared Memory 上的访存时，我们的 Shared Memory 模型只讨论了一个 Warp 内每个线程所访问的元素。在涉及到向量化访存时这样的模型就不起效果了，因为通过一个 LDS.64 或 LDS.128 指令就可以让一个线程一次性访问 8 个或 16 个字节（相当于 2 个或 4 个元素）。

正确的做法应该是就每个 Wrap 内所产生的每个 Memory Transaction而非每个 Warp 或每条指令来讨论（参考这里）。那么一个在 Shared Memory 上的向量化指令 LDS.64 或 LDS.128 指令到底对应多少个 Memory Transaction？我并没有找到 NVIDIA 给出的官方答案，通过一些网络上的讨论和我自己的 Microbenchmark，我对结合了向量化指令的 Shared Memory 的访存机制的推测如下。

首先一个原则是一个 Warp 中所有线程在同时执行一条 Shared Memory 访存指令时会对应到 1 个或多个 Memory Transaction，一个 Memory Transaction 最长是 128 字节。如果一个 Warp 内在同一时刻所需要的访存超过了 128 字节，那么会则被拆成多个 Transaction 进行。因为一个 Warp 同一时刻执行的访存指令的位宽应该是一样的（即例如不存在线程 0 执行 LDS.32 而线程 1 执行 LDS.128），因此我们只需要对 64 位宽和 128 位宽的访存指令分别讨论即可。

64 位宽的访存指令

对于 64 位宽的访存指令而言，除非触发广播机制，否则一个 Warp 中有多少个活跃的 Half-Warp 就需要多少个 Memory Transaction，一个 Half-Warp 活跃的定义是这个 Half-Warp 内有任意一个线程活跃。触发广播机制只需满足以下条件中的至少一个：

  • 对于 Warp 内所有活跃的第 ii 号线程，第 i\mathrm{xor}1i xor 1 号线程不活跃或者访存地址和其一致；
  • 对于 Warp 内所有活跃的第 ii 号线程，第 i\mathrm{xor}2i xor 2 号线程不活跃或者访存地址和其一致；

如果触发了广播机制，那么两个 Half-Warp 内的 Memory Transaction 可以合并成一个。

我们看几个例子：

在这里插入图片描述
Case 1: 活跃线程全部在第 1 个 Half-Warp 内，需要 1 个 Memory Transaction，没有 Bank Conflict，需要 1 个 Wavefront

在这里插入图片描述
Case 2: 活跃线程分散在了 2 个 Half-Warp 内，需要 2 个 Memory Transaction，没有 Bank Conflict，需要 2 个 Wavefront（注意第 15 号和第 16 号线程）

在这里插入图片描述
Case 3: 活跃线程分散在了 2 个 Half-Warp 内，但因为触发了广播机制中的第一条，因此仍然只需要 1 个 Memory Transaction，没有 Bank Conflict，需要 1 个 Wavefront

在这里插入图片描述
Case 4: 活跃线程分散在了 2 个 Half-Warp 内，看似好像触发了广播机制，但其实并没有，因为第一个 Half-Warp 触发的是第一条，第二个 Half-Warp 触发的是第二条，因此仍然需要 2 个 Memory Transaction，没有 Bank Conflict，需要 2 个 Wavefront

在这里插入图片描述
Case 5: 活跃线程分散在了 2 个 Half-Warp 内，没有触发广播机制，需要 2 个 Memory Transaction，没有 Bank Conflict，需要 2 个 Wavefront

可以通过 Nsight Compute 跑一跑下面的代码并观察和 Shared Memory 的相关 Metric 来验证上面这五个例子：
smem_64bit.cu

#include <cstdint>__global__ void smem_1(uint32_t *a) {__shared__ uint32_t smem[128];uint32_t tid = threadIdx.x;for (int i = 0; i < 4; i++) {smem[i * 32 + tid] = tid;}__syncthreads();if (tid < 16) {reinterpret_cast<uint2 *>(a)[tid] =reinterpret_cast<const uint2 *>(smem)[tid];}
}__global__ void smem_2(uint32_t *a) {__shared__ uint32_t smem[128];uint32_t tid = threadIdx.x;for (int i = 0; i < 4; i++) {smem[i * 32 + tid] = tid;}__syncthreads();if (tid < 15 || tid == 16) {reinterpret_cast<uint2 *>(a)[tid] =reinterpret_cast<const uint2 *>(smem)[tid == 16 ? 15 : tid];}
}__global__ void smem_3(uint32_t *a) {__shared__ uint32_t smem[128];uint32_t tid = threadIdx.x;for (int i = 0; i < 4; i++) {smem[i * 32 + tid] = tid;}__syncthreads();reinterpret_cast<uint2 *>(a)[tid] =reinterpret_cast<const uint2 *>(smem)[tid / 2];
}__global__ void smem_4(uint32_t *a) {__shared__ uint32_t smem[128];uint32_t tid = threadIdx.x;for (int i = 0; i < 4; i++) {smem[i * 32 + tid] = tid;}__syncthreads();uint32_t addr;if (tid < 16) {addr = tid / 2;} else {addr = (tid / 4) * 4 + (tid % 4) % 2;}reinterpret_cast<uint2 *>(a)[tid] =reinterpret_cast<const uint2 *>(smem)[addr];
}__global__ void smem_5(uint32_t *a) {__shared__ uint32_t smem[128];uint32_t tid = threadIdx.x;for (int i = 0; i < 4; i++) {smem[i * 32 + tid] = tid;}__syncthreads();reinterpret_cast<uint2 *>(a)[tid] =reinterpret_cast<const uint2 *>(smem)[tid % 16];
}int main() {uint32_t *d_a;cudaMalloc(&d_a, sizeof(uint32_t) * 128);smem_1<<<1, 32>>>(d_a);smem_2<<<1, 32>>>(d_a);smem_3<<<1, 32>>>(d_a);smem_4<<<1, 32>>>(d_a);smem_5<<<1, 32>>>(d_a);cudaFree(d_a);cudaDeviceSynchronize();return 0;
}

上面的例子中并没有列举出有 Bank Conflict 的情况，那么 Bank Conflict 在这种情况下应该如何考虑呢？正如前面提到的那样，我们只需要计算每个 Memory Transaction 中的 Bank Conflict 数目然后加起来就好了（因为 Memory Transaction 是串行的）。

128 位宽的访存指令

128 位宽的访存指令和 64 位宽的访存指令是类似的，不同的是需要以 Half-Warp 为单位来计算，对于每个 Half-Warp 而言，除非触发广播机制，这个 Half-Warp 中有多少个活跃的 Quarter-Warp 就需要多少个 Memory Transaction，一个 Quarter-Warp 活跃的定义是这个 Quarter-Warp 内有任意一个线程活跃。类似地，如果触发广播机制那么两个 Quarter-Warp 中的 Transaction 就可以被合并成一个。 触发广播机制的条件和 64 位宽的访存指令是一样的（注意广播机制是以整个 Warp 为单位考虑）。这也就意味着假设一个 Warp 中 32 个线程都活跃，即使它们的访存地址都一样，也需要 2 个 Memory Transaction。

同样来看几个例子：

在这里插入图片描述

Case 1: 活跃线程分散在了 2 个 Half-Warp 和 2 个 Quarter-Warp 内，每个 Half-Warp 需要 1 个 Memory Transaction，总共需要 2 个 Memory Transaction，没有 Bank Conflict，需要 2 个 Wavefront

在这里插入图片描述

Case 2: 活跃线程分散在了 1 个 Half-Warp 和 2 个 Quarter-Warp 内，需要 1 个 Memory Transaction，没有 Bank Conflict，需要 1 个 Wavefront

在这里插入图片描述
Case 3: 活跃线程分散在了 2 个 Half-Warp 和 4 个 Quarter-Warp 内，但触发了广播机制（第一条），每个 Half-Warp 需要 1 个 Memory Transaction，总共需要 2 个 Memory Transaction，没有 Bank Conflict，需要 2 个 Wavefront

在这里插入图片描述
Case 4: 活跃线程分散在了 2 个 Half-Warp 和 4 个 Quarter-Warp 内，没有触发广播机制，每个 Half-Warp 需要 2 个 Memory Transaction，总共需要 4 个 Memory Transaction，没有 Bank Conflict，需要 4 个 Wavefront

在这里插入图片描述
Case 5: 活跃线程分散在了 2 个 Half-Warp 和 4 个 Quarter-Warp 内，但触发了广播机制（第一条和第二条），每个 Half-Warp 需要 1 个 Memory Transaction，总共需要 2 个 Memory Transaction，但因为每个 Memory Transaction 中有 1 个 Bank Conflict，因此会拆分成 4 个 Memory Transaction，对应需要 4 个 Wavefront

在这里插入图片描述

Case 6: 活跃线程分散在了 2 个 Half-Warp 和 4 个 Quarter-Warp 内，没有触发广播机制，每个 Half-Warp 需要 2 个 Memory Transaction，总共需要 4 个 Memory Transaction，没有 Bank Conflict，需要 4 个 Wavefront

同样可以通过 Nsight Compute 跑一跑下面的代码并观察和 Shared Memory 的相关 Metric 来验证上面这几个例子：

smem_128bit.cu

#include <cstdint>__global__ void smem_1(uint32_t *a) {__shared__ uint32_t smem[128];uint32_t tid = threadIdx.x;for (int i = 0; i < 4; i++) {smem[i * 32 + tid] = tid;}__syncthreads();if (tid == 15 || tid == 16) {reinterpret_cast<uint4 *>(a)[tid] =reinterpret_cast<const uint4 *>(smem)[4];}
}__global__ void smem_2(uint32_t *a) {__shared__ uint32_t smem[128];uint32_t tid = threadIdx.x;for (int i = 0; i < 4; i++) {smem[i * 32 + tid] = tid;}__syncthreads();if (tid == 0 || tid == 15) {reinterpret_cast<uint4 *>(a)[tid] =reinterpret_cast<const uint4 *>(smem)[4];}
}__global__ void smem_3(uint32_t *a) {__shared__ uint32_t smem[128];uint32_t tid = threadIdx.x;for (int i = 0; i < 4; i++) {smem[i * 32 + tid] = tid;}__syncthreads();reinterpret_cast<uint4 *>(a)[tid] = reinterpret_cast<const uint4 *>(smem)[(tid / 8) * 2 + ((tid % 8) / 2) % 2];
}__global__ void smem_4(uint32_t *a) {__shared__ uint32_t smem[128];uint32_t tid = threadIdx.x;for (int i = 0; i < 4; i++) {smem[i * 32 + tid] = tid;}__syncthreads();uint32_t addr;if (tid < 16) {addr = (tid / 8) * 2 + ((tid % 8) / 2) % 2;} else {addr = (tid / 8) * 2 + ((tid % 8) % 2);}reinterpret_cast<uint4 *>(a)[tid] =reinterpret_cast<const uint4 *>(smem)[addr];
}__global__ void smem_5(uint32_t *a) {__shared__ uint32_t smem[128];uint32_t tid = threadIdx.x;for (int i = 0; i < 4; i++) {smem[i * 32 + tid] = tid;}__syncthreads();reinterpret_cast<uint4 *>(a)[tid] =reinterpret_cast<const uint4 *>(smem)[(tid / 16) * 4 + (tid % 16) / 8 + (tid % 8) / 4 * 8];
}__global__ void smem_6(uint32_t *a) {__shared__ uint32_t smem[128];uint32_t tid = threadIdx.x;for (int i = 0; i < 4; i++) {smem[i * 32 + tid] = tid;}__syncthreads();uint32_t addr = (tid / 16) * 4 + (tid % 16 / 8) * 8;if (tid < 16) {addr += (tid % 4 / 2) * 2;} else {addr += (tid % 4 % 2) * 2;}reinterpret_cast<uint4 *>(a)[tid] =reinterpret_cast<const uint4 *>(smem)[addr];
}int main() {uint32_t *d_a;cudaMalloc(&d_a, sizeof(uint32_t) * 128);smem_1<<<1, 32>>>(d_a);smem_2<<<1, 32>>>(d_a);smem_3<<<1, 32>>>(d_a);smem_4<<<1, 32>>>(d_a);smem_5<<<1, 32>>>(d_a);smem_6<<<1, 32>>>(d_a);cudaFree(d_a);cudaDeviceSynchronize();return 0;
}

总结

可以看到实际上 32 位宽的访存指令和 64/128 位宽的访存指令的广播机制以及 Bank Conflict 的计算都有很大的不同，但是官方文档中没有出现相关的描述（或者我没看到😵‍💫）。这篇文章通过 Microbenchmark 以及前人的一些讨论总结了几套规则，但需要注意的是这些规则有一定的局限性，其一是我只评测了以上图片中的例子，因此是不清楚更加复杂的访存情况是不是仍然符合这些规则的，其二是这些规则不是官方记录的，因此很有可能在将来被新发布的 GPU 架构所改写。

参考

  • How to understand the bank conflict of shared_mem

  • Unexpected shared memory bank conflict

  • Volta GPU 白皮书

  • Turing GPU 白皮书

  • volta-architecture-whitepaper.pdf

  • NVIDIA-Turing-Architecture-Whitepaper.pdf

本文允许以 CC BY-NC-SA 4.0 的方式授权转载，背景图片来源于 KARL EGGER

最后更新于 2024 年 4 月 1 日 22 时 58 分

CUDA 程序优化

1. 基础介绍

简介

本合集主要介绍我在开发分布式异构训练框架时的 CUDA 编程实践和性能优化的相关内容。主要包含以下几个部分:

  1. 介绍 CUDA 的基本概念和架构，帮助读者建立对 CUDA 的初步认识，包括硬件架构 / CUDA 基础等内容
  2. 介绍一些性能优化技巧和工具，帮助读者优化 CUDA 程序的执行效率
  3. 结合具体的代码示例来说明一个 cuda 程序的优化思路和结果，帮助读者更好地理解和掌握 CUDA 编程和性能优化的实践方法

希望通过本文档，能够帮助大家写出更高效的 CUDA 程序。下面我们就开始吧～

硬件架构

要说清楚为什么 GPU 比 CPU 更适合大规模并行计算，要从硬件层面开始说起

[image](https://img2024.cnblogs.com/blog/1439743/202406/1439743-20240606194543448-637778562.png =700x)

以当前较主流的硬件 i9-14900k 和 A100 为例:

i9-14900k: 24 核心，32 线程 (只能在 16 个能效核上进行超线程), L2: 32MB, L3: 36MB, 内存通信带宽 89.6GB/s

A100: 108 SM, 6912 CUDA core, 192KB L1, 60MB L2, 40GB DRAM.

我个人的理解，GPU 的运算核心之所以远多于 CPU, 是因为远少于 CPU 的控制逻辑. GPU 每个 core 内不需要考虑线程调度的情况，不需要保证严格一致的运算顺序，另外每个 sm 都有自己独立的寄存器和 L1, 对线程的切换重入非常友好，所以更适合大规模数据的并行运算。而这种设计方式也会对程序员提出更高的要求，纯 CPU 程序可能写的最好的代码和最差的情况有个 2/3 倍的性能差距就很大了，而 CUDA kernel 可能会相差几十倍甚至几百倍.

[image](https://img2024.cnblogs.com/blog/1439743/202406/1439743-20240606194556694-1290930234.png =700x)

HBM（High-Bandwidth Memory） ：HBM 是高带宽内存，也就是常说的显存，这张图里的 DRAM。 带宽: 1.5TB/s

L2 Cache：L2 Cache 是 GPU 中更大容量的高速缓存层，可以被多个 SM 访问。L2 Cache 还可以用于协调 SM 之间的数据共享和通信。 带宽: 4TB/s

SM (Streaming Multiprocessor) ：GPU 的主要计算单元，负责执行并行计算任务。每个 SM 都包含多个 CUDA core，也就是 CUDA 里 Block 执行的地方，关于 block_size 如何设置可以参考 block_size 设置, 跟随硬件不同而改变，通常为 128/256

L1 Cache/SMEM：, 也叫 shared_memory, 每个 SM 独享一个 L1 Cache，CUDA 里常用于单个 Block 内部的临时计算结果的存储，比如 cub 里的 Block 系列方法就经常使用，带宽: 19TB/s

SMP (SM partition): A100 中有 4 个。每个有自己的 warp 调度器，寄存器等.

CUDA Core: 图里绿色的 FP32/FP64/INT32 等就是，是 thread 执行的基本单位

Tensor Core: Volta 架构之后新增的单元，主要用于矩阵运算的加速

WARP (Wavefront Parallelism) ：WARP 指的是一组同时执行的 Thread，固定 32 个，不够 32 时也会按 32 分配. warp 一个线程对内存操作后，其他 warp 内的线程是可见的.

Dispatch Unit： 从指令队列中获取和解码指令，协调指令的执行和调度

Register File: 寄存器用于存储临时数据、计算中间结果和变量。GPU 的寄存器比 CPU 要多很多

2. cuda 基础

cuda 基础语法上和 c/c++ 是一致的。引入了 host/device 定义，host 指的是 cpu 端，device 指的是 gpu 端

个人感觉最难的部分在于并行的编程思想和 cpu 编程的思想差异比较大。我们以一个向量相加的 demo 程序举例:

__global__ void add_kernel(int *a, int *b, int *c, int n) {int index = threadIdx.x + blockIdx.x * blockDim.x;if (index < n) {c[index] = a[index] + b[index];}
}int main() {int *a, *b, *c;int *d_a, *d_b, *d_c;int n = 10000;int size = n * sizeof(int);cudaMalloc((void**)&d_a, size);cudaMalloc((void**)&d_b, size);cudaMalloc((void**)&d_c, size);a = (int*)malloc(size);random_ints(a, n);b = (int*)malloc(size);random_ints(b, n);c = (int*)malloc(size);cudaMemcpy(d_a, a, size, cudaMemcpyHostToDevice);cudaMemcpy(d_b, b, size, cudaMemcpyHostToDevice);//cuda kerneladd_kernel<<<(n + threads_per_block - 1)/threads_per_block, threads_per_block>>>(d_a, d_b, d_c, n);cudaMemcpy(c, d_c, size, cudaMemcpyDeviceToHost);cudaFree(d_a);cudaFree(d_b);cudaFree(d_c);return 0;
}

描述符

cuda 新增了三个描述符:

__global__: 在 device 上运行，可以从 host/device 上调用，返回值必须是 void, 异步执行.

__device__: 在 device 上运行和调用

__host__: 只能在 host 上执行和调用

CUDA Kernel

cuda_kernel 是由 <<<>>> 围起来的，里面主要有 4 个参数用来配置这个 kernel <<<grid_size, block_size, shared_mem_size, stream>>>

grid_size: 以一维 block 为例，grid_size 计算以 (thread_num + block_size - 1) /block_size 计算大小

block_size: 见上面 SM 部分介绍

shared_mem_size: 如果按 __shared__ int a [] 方法声明共享内存，需要在这里填需要分配的共享内存大小。注意不能超过硬件限制，比如 A100 192KB

stream: 异步多流执行时的 cuda 操作队列，在这个流上的所有 kernel 是串行执行的，多个流之间是异步执行的。后续会在异步章节里详细介绍

整个过程如下图，先通过 cudaMemcpy 把输入数据 copy 到显存 ->cpu 提交 kernel->gpu kernel_launch-> 结果写回线程 ->DeviceToHost copy 回内存.

[image](https://img2024.cnblogs.com/blog/1439743/202406/1439743-20240606194617352-1175795907.png =700x)

add_kernel 相当于我们将 for 循环拆分为了每个线程只处理一个元素的相加的并行执行。通过 nvcc 编译后就完成了第一个 kernel 的编写。下一篇会以一个具体的例子来讲如何进行 kernel 的性能分析和调优.

常用库

thrust: cuda 中类似于 c++ STL 的定位，一些类似于 STL 的常见算法可以在这里找到现成的实现，比如 sort/reduce/unique/random 等。文档: https://nvidia.github.io/cccl/thrust/api/namespace_thrust.html

cudnn: 神经网络加速的常用库。包含卷积 /pooling/softmax/normalization 等常见 op 的优化实现.

cuBlas: 线性代数相关的库。进行矩阵运算时可以考虑使用，比如非常经典的矩阵乘法实现 cublasSgemm

Cub: warp/block/device 级的编程组件，非常常用。文档: https://nvidia.github.io/cccl/cub/

nccl: 集合通信库。用于卡间通信 / 多机通信

相关资料

cuda 编程指导手册: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#programming-model

性能分析工具 Nsight-System & Nsight-Compute: https://docs.nvidia.com/nsight-systems/index.html

2. 访存优化

简介

在 CUDA 程序中，访存优化个人认为是最重要的优化项。往往 kernel 会卡在数据传输而不是计算上，为了最大限度利用 GPU 的计算能力，我们需要根据 GPU 硬件架构对 kernel 访存进行合理的编写.

这章主要以计算一个 tensor 的模为例，来看具体如何优化访存从而提升并行效率。以下代码都只列举了 kernel 部分，省略了 host 提交 kernel 的部分. Block_size 均为 256, 代码均在 A100 上评估性能.
∥ v ⃗ ∥ = v 1 2 + v 2 2 + ⋯ + v n 2 \| \vec{v} \| = \sqrt{v_1^2 + v_2^2 + \cdots + v_n^2} ∥v ∥=v12​+v22​+⋯+vn2​ ​

CPU 代码

void cpu_tensor_norm(float* tensor, float* result, size_t len) {double sum = 0.0;for (auto i = 0; i < len; ++i) {sum += tensor[i] * tensor[i];}*result = sqrt(sum);
}

是一个非常简单的函数，当然这里针对 CPU 程序也有优化方法，当前实现并不是最优解。此处忽略不表

GPU 实现 1 - 基于 CPU 并行思想

在传统的多线程 CPU 任务中，如果想处理一个超大数组的求和，很自然的会想到起 N 个线程，每个线程算 1/N 的和，最后再把这 N 个和加到一块求 sqrt. 根据这个思路实现了 kernel1, 使用 1 个 block, block 中每个线程计算 1/N 连续存储的数据.

__global__ void norm_kernel1(float* tensor, float* result, size_t len) {auto tid = threadIdx.x;auto size = len / blockDim.x;double sum = 0;for (auto i = tid * size; i < (tid + 1) * size; i++) {sum += tensor[i] * tensor[i];}result[tid] = float(sum);
}
//norm_kernel1<<<1, block_size>>>(d_tensor, d_result1, len);

[image](https://img2024.cnblogs.com/blog/1439743/202406/1439743-20240617152246087-1253896473.png =700x)

接下来我们使用 nsight-compute 来分析下第一个 kernel 实现哪些地方不合理。具体使用方法可以从官方文档获取，这里举个命令例子: nv-nsight-cu-cli -f --target-processes all -o profile --set full --devices 0 ./output/norm_kernel_bench
[image](https://img2024.cnblogs.com/blog/1439743/202406/1439743-20240617152307166-1642995465.png =700x)

ncu 分析给出了 2 个主要问题:

grid 太小

因为我们只使用了 256 个线程。而 A100 光 CUDA core 就 7000 个。在 GPU 中不像 CPU 中线程上下文切换具有很高的成本，如果想充分利用算力，就要用尽量多的线程来提升并发度。同时线程数也不能无限制增加，因为如果每个线程使用了 32 个寄存器，而 SM 中最多 16384 的寄存器空间的话，多于 16384/32=512 个线程后，这些多出来的线程就需要把数据存到显存里，反而会降低运行效率.

读显存瓶颈

下面打开 detail 后也给出了问题日志: Uncoalesced global access, expected 262144 transactions, got 2097152 (8.00x) at PC

coalesced 指的是显存读取需要是连续的，这里也许你会有疑问，在 kernel1 里就是按照连续的显存读的呀。这里涉及到 GPU 的实际执行方式，当一个 thread 在等读显存数据完成的时候，GPU 会切换到下一个 thread, 也就是说是需要让 thread1 读显存的数据和 thread2 的数据是连续的才会提升显存的读取效率，在 kernel1 中明显不连续。同时参考 ProgrammingGuide 中的描述，每个 warp 对显存的访问是需要对齐 32/64/128 bytes (一次数据传输的数据量在默认情况下是 32 字节), 如果所有数据传输处理的数据都是该 warp 所需要的，那么合并度为 100%，即合并访问

GPU 实现 2 - 优化 coalesced

__global__ void norm_kernel2(float* tensor, float* result, size_t len) {auto tid = threadIdx.x;double sum = 0.0;while (tid < len) {sum += tensor[tid] * tensor[tid];tid += blockDim.x;}result[threadIdx.x] = float(sum);
}
//norm_kernel2<<<1, block_size>>>(d_tensor, d_result1, len);

依然还是 1 个 block256 个线程，kernel2 将每个 thread 读取方式改成了每隔 256 个 float 读一个，这样 uncoalesced 的报错就不见了。但是！一跑 bench 却发现为啥反而还变慢了呢？
[image](https://img2024.cnblogs.com/blog/1439743/202406/1439743-20240617152401330-978761037.png =700x)

此时可以点开 memory WorkLoad Analysis

[image](https://img2024.cnblogs.com/blog/1439743/202406/1439743-20240617152422066-825943602.png =700x)

从这里可以看到 L2 cache 命中率降低了 50% 左右，原因是因为按照 kernel1 的访问方式，第一次访问了 32bytes 长度，但是只用了一部分后，剩下的会缓存在 L2 中，而 Kernel2 虽然访问显存连续了，但每次的 cache 命中率会随着读入数据利用效率的变高而降低，根本原因是因为线程和 block 太少导致的。另外这张图上还有个明显的叹号，我们没有合理的用到 shared_memory. 接下来的 kernel3 重点优化这两部分

GPU 实现 3 - 增大并发 & 利用 shared_memory

__global__ void norm_kernel3(float* tensor, float* result, size_t len) {auto tid = threadIdx.x + blockIdx.x * blockDim.x;extern __shared__ double sum[];auto loop_stride = gridDim.x * blockDim.x;sum[threadIdx.x] = 0;while (tid < len) {sum[threadIdx.x] += tensor[tid] * tensor[tid];tid += loop_stride;}__syncthreads();if (threadIdx.x == 0) {for (auto i = 1; i < blockDim.x; ++i) {sum[0] += sum[i];}result[blockIdx.x] = float(sum[0]);}
}
//grid_size3=64
//norm_kernel3<<<grid_size3, block_size, block_size * sizeof(double)>>>(d_tensor, d_result3, len);
//for (auto i = 1; i < grid_size3; ++i) {
//    h_result3[0] += h_result3[i];
//}
//sqrt(h_result3[0])

[image](https://img2024.cnblogs.com/blog/1439743/202406/1439743-20240617152437547-1649475252.png =700x)

短短几行改动，让程序快了 27 倍，这是什么魔法 (黑人问号)? kernel3 做了如下几个优化:

  1. 使用 shared_memory 用来存储临时加和，最后在每个 block 的第一个 thread 里把这些加和再加到一块，最后再写回 HBM. shared_memory 访问速度 19T/s, HBM 速度才 1.5TB/s, 所以我们如果有需要临时存储的变量，要尽可能的把 shared_mem 利用起来.
  2. 这次使用了 64 个 block, GPU 的其他 SM 终于不用看戏了。但其实还是可以增加的，A100 有 108 个 SM 呢，让我们把他用满再看下性能 auto grid_size3 = (len + block_size - 1) /block_size; 可以看到我们终于让计算吞吐打满了～

[image](https://img2024.cnblogs.com/blog/1439743/202406/1439743-20240617152524570-318241188.png =700x)

到这里还有 2 个问题需要解决: 1. 我们只是在 GPU 里求了局部加和，全局和还得挪到 CPU 算好挫. 2. 每个 block 在 syncthread 之后只有第一个线程在计算，能不能加快计算的同时减少一下计算量

GPU 实现 4 - 优化加和

template <int64_t BLOCK_SIZE>
__global__ void norm_kernel4(float* tensor, double* result, size_t len) {using BlockReduce = cub::BlockReduce<double, BLOCK_SIZE>;__shared__ typename BlockReduce::TempStorage temp_storage;int tid = threadIdx.x + blockIdx.x * blockDim.x;double sum = 0.0;if (tid < len) {sum += tensor[tid] * tensor[tid];}double block_sum = BlockReduce(temp_storage).Sum(sum);if (threadIdx.x == 0) {atomicAdd(result, block_sum);}
}
//norm_kernel4<block_size><<<grid_size, block_size>>>(d_tensor, d_result4, len);
//sqrt(h_result4)

[image](https://img2024.cnblogs.com/blog/1439743/202406/1439743-20240617152540197-967725268.png =700x)

对与第一个线程连续加和的问题，我偷懒使用了 cub::BlockReduce 方法，BlockReduce 的原理是经典的树形规约算法。利用分治的思想，把原来的 32 轮加和可以简化为 5 轮加和，这样就能极大减少长尾线程的计算量.

[image](https://img2024.cnblogs.com/blog/1439743/202406/1439743-20240617152547698-1157082575.png =700x)

对于显存上的全局求和问题，由于 block 之间是没有任何关联的，我们必须使用原子操作来解决对全局显存的操作。这里为了减少原子写冲突，只在 block 的第一个线程上进行原子加。另外我们可以不用 cub 的 BlockReduce 优化到和他差不多的性能吗？

GPU 实现 5 - 自己实现 BlockReduce

template <int64_t BLOCK_SIZE>
__global__ void norm_kernel5(float* tensor, double* result, size_t len) {using WarpReduce = cub::WarpReduce<double>;const int64_t warp_size = BLOCK_SIZE / 32;__shared__ typename WarpReduce::TempStorage temp_storage[warp_size];__shared__ float reduce_sum[BLOCK_SIZE];int tid = threadIdx.x + blockIdx.x * blockDim.x;int warp_id = threadIdx.x / 32;float sum = 0.0;if (tid < len) {sum += tensor[tid] * tensor[tid];}auto warp_sum = WarpReduce(temp_storage[warp_id]).Sum(sum);reduce_sum[threadIdx.x] = warp_sum; //这里尽量避免wrap内的分支__syncthreads();//树形规约int offset = warp_size >> 1;if (threadIdx.x % 32 == 0) {while (offset > 0) {if (warp_id < offset) {reduce_sum[warp_id * 32] += reduce_sum[(warp_id + offset) * 32];}offset >>= 1;__syncthreads();}}if (threadIdx.x == 0) {atomicAdd(result, reduce_sum[0]);}
}
//norm_kernel5<block_size><<<grid_size5, block_size>>>(d_tensor, d_result5, len);
//sqrt(h_result5)

[image](https://img2024.cnblogs.com/blog/1439743/202406/1439743-20240617152605137-331016775.png =700x)

kernel5 主要分为两个部分，第一步进行了 warp 内的规约，第二步手动实现了树形规约的方法。耗时基本与 BlockReduce 一致，由于 warp 内的 32 个线程会共享寄存器和 shared_mem 读写，在 warp 内先做一些规约可以适当减少后续的 sync_threads 轮数.
其实这个 kernel 还是有很大优化空间，篇幅受限原因更深层次的优化技巧在后续说明

计算分析

[image](https://img2024.cnblogs.com/blog/1439743/202406/1439743-20240617152622218-867073834.png =700x)

在汇编指令统计这块可以看到 LDS (从 shared_mem 加载指令), LOP3 (logic 操作), STS (往 mem 中写入操作), BAR (barrier), BSYNC (线程同步时的 barrier, 对应 atomic 操作), WARPSYNC (warp 同步) 这些相较于 kernel4 多了很多. WARPSYNC/LOP3/BAR 这些变多是正常的，是 kernel5 里新增的逻辑. LDS/STS 增加应该是因为我们在 warp_reduce 时的频率比 block_reduce 对共享内存的访问频率更高导致.

访存注意要点 - bank conflicts

[image](https://img2024.cnblogs.com/blog/1439743/202406/1439743-20240617152637810-1790818641.png =700x)

内存分析里出现了一个新名词 bank conflicts , 先根据 官方文档 解释下这个名词.

SharedMemory 结构：放在 shared memory 中的数据是以 4bytes 作为 1 个 word，依次放在 32 个 banks 中。第 i 个 word 存放在第 (i % 32) 个 bank 上。每个 bank 在每个 cycle 的 bandwidth 为 4bytes。所以 shared memory 在每个 cycle 的 bandwidth 为 128bytes。这也意味着每次内存访问只会访问 128bytes 数据

[image](https://img2024.cnblogs.com/blog/1439743/202406/1439743-20240617152647400-1139929431.png =700x)

如果同一个 warp 内的多个 threads 同时访问同一个 bank 内的同一个 word, 会触发 broadcast 机制，会同时发给多个 thread. 不会产生冲突

冲突主要产生在多个 threads 访问同一个 bank 内的不同 word, 如上图的第二列。这样就会导致本来的一次 memory transaction 被强制拆分成了 2 次，而且需要 ** 串行 ** 访问 memory

解决方法:

通过错位的方式访问数组，避免访问步长和 32 产生交集，每个线程根据线程编号 tid 与访问步长 s 的乘积来访问数组的 32-bits 字 (word)：

extern __shared__ float shared[];
float data = shared[baseIndex + s * tid];

如果按照上面的方式，那么当 s*n 是 bank 的数量 (即 32) 的整数倍时或者说 n 是 32/d 的整数倍 (d 是 32 和 s 的最大公约数) 时，线程 tid 和线程 tid+n 会访问相同的 bank。我们不难知道如果 tid 与 tid+n 位于同一个 warp 时，就会发生 bank 冲突，相反则不会。

仔细思考你会发现，只有 warp 的大小 (即 32) 小于等于 32/d 时，才不会有 bank 冲突，而只有当 d 等于 1 时才能满足这个条件。要想让 32 和 s 的最大公约数 d 为 1，s 必须为奇数。于是，这里有一个显而易见的结论：当访问 ** 步长 s 为奇数 ** 时，就不会发生 bank 冲突。

via：

  • CUDA Shared Memory 在向量化指令下的访存机制 孤独代码 发表于 2023 年 12 月 26 日
    https://code.hitori.moe/post/cuda-shared-memory-access-mechanism-with-vectorized-instructions/

  • CUDA 程序优化 - 1. 基础介绍 - SunStriKE - 博客园 posted @ 2024-06-06 19:49
    https://www.cnblogs.com/sunstrikes/p/18235920

  • cuda 程序优化 - 2. 访存优化 - SunStriKE - 博客园 posted @ 2024-06-17 15:31
    https://www.cnblogs.com/sunstrikes/p/18252517

本文来自互联网用户投稿，该文观点仅代表作者本人，不代表本站立场。本站仅提供信息存储空间服务，不拥有所有权，不承担相关法律责任。如若转载，请注明出处：http://www.pswp.cn/web/81097.shtml

如若内容造成侵权/违法违规/事实不符，请联系多彩编程网进行投诉反馈email:809451989@qq.com，一经查实，立即删除！

相关文章

NVMe高速传输之摆脱XDMA设计1

NVMe高速传输之摆脱XDMA设计1

NVMe IP放弃XDMA原因 选用XDMA做NVMe IP的关键传输模块&#xff0c;可以加速IP的设计&#xff0c;但是XDMA对于开发者来说&#xff0c;还是不方便&#xff0c;原因是它就象一个黑匣子&#xff0c;调试也非一番周折&#xff0c;尤其是后面PCIe4.0升级。 因此决定直接采用PCIe设…
阅读更多...
企业级单元测试流程

企业级单元测试流程

企业级的单元测试流程不仅是简单编写测试用例&#xff0c;而是一整套系统化、自动化、可维护、可度量的工程实践&#xff0c;贯穿从代码编写到上线部署的全生命周期。下面是一个尽可能完善的 企业级单元测试流程设计方案&#xff0c;适用于 Java 生态&#xff08;JUnit Mockit…
阅读更多...
关于vector、queue、list哪边是front、哪边是back，增加、删除元素操作

关于vector、queue、list哪边是front、哪边是back，增加、删除元素操作

容器的 front、back 及操作方向 1.1vector&#xff08;动态数组&#xff09; 结构&#xff1a;连续内存块&#xff0c;支持快速随机访问。 操作方向&#xff1a; front&#xff1a;第一个元素&#xff08;索引 0&#xff09;。 back&#xff1a;最后一个元素&#xff08;索引…
阅读更多...
嵌入式之汇编程序示例

嵌入式之汇编程序示例

目录 经典例子:求阶乘 一:数组求和 二:数据压栈退栈 三:函数嵌套调用 经典例子:求阶乘 知识点: BGT 用于判断 r2 > r0&#xff0c;确保循环执行 恰好 r0 次。BNE 用于判断 r2 ≠ r0&#xff0c;会导致循环多执行一次&#xff0c;得到错误结果。 这就是阶乘代码中必须…
阅读更多...
【MySQL】第九弹——索引(下)

【MySQL】第九弹——索引(下)

文章目录 &#x1f30f;索引(上)回顾&#x1f30f;使用索引&#x1fa90;自动创建索引&#x1fa90;手动创建索引&#x1f680;主键索引&#x1f680;普通索引&#x1f680;唯一索引&#x1f680;复合索引 &#x1fa90;查看索引&#x1fa90;删除索引&#x1f680;删除主键索引…
阅读更多...
毕业论文格式（Word）

毕业论文格式（Word）

目录 Word目录怎么自动生成&#xff1f;快速生成试试这3个方法&#xff01; - 知乎https://zhuanlan.zhihu.com/p/692056836目录生成需要先设置标题样式&#xff0c;这个不仅是目录生成需要&#xff0c;和后续的图表也有关系。 最好不要自己创建新的样式&#xff0c;而是在现有…
阅读更多...
PostGIS实现栅格数据转二进制应用实践【ST_AsBinary】

PostGIS实现栅格数据转二进制应用实践【ST_AsBinary】

ST_AsBinary解析与应用实践&#xff08;同ST_AsWKB&#xff09; 一、函数概述二、核心参数解析三、典型用法示例四、Out-DB 波段处理机制五、二进制格式与其他格式的转换六、性能与存储优化七、应用场景八、注意事项九、扩展应用&#xff1a;基于Python Web的栅格二进制数据的…
阅读更多...
线性回归原理推导与应用（七）：逻辑回归原理与公式推导

线性回归原理推导与应用（七）：逻辑回归原理与公式推导

逻辑回归是一种分类算法&#xff0c;常用于二分类&#xff0c;也就是得出的结果为是和不是&#xff0c;例如通过各种因素判断一个人是否生病&#xff0c;信用卡是否违约等。逻辑回归在社会和自然科学中应用非常广泛&#xff0c; 前置知识 线性回归 逻辑回归的底层方法就是线…
阅读更多...
Fastrace：Rust 中分布式追踪的现代化方案

Fastrace：Rust 中分布式追踪的现代化方案

原文链接&#xff1a;Fastrace: A Modern Approach to Distributed Tracing in Rust | FastLabs / Blog 摘要 在微服务架构中&#xff0c;分布式追踪对于理解应用程序的行为至关重要。虽然 tokio-rs/tracing 在 Rust 中被广泛使用&#xff0c;但它存在一些显著的挑战&#xf…
阅读更多...
水果系列数据集- 葡萄grapes＞＞ DataBall

水果系列数据集- 葡萄grapes＞＞ DataBall

该数据集可以用于目标检测&#xff0c;水果分类 &#xff0c;文生图相关项目。 以下是图片样例&#xff1a;
阅读更多...
HTTP协议接口三种测试方法之-postman

HTTP协议接口三种测试方法之-postman

HTTP协议作为现代Web开发的基石&#xff0c;其接口测试是开发过程中不可或缺的环节。Postman作为最流行的API测试工具之一&#xff0c;能够极大提升我们的测试效率。本文将详细介绍如何使用Postman进行HTTP接口测试。 一、HTTP协议基础回顾 在开始使用Postman之前&#xff0c…
阅读更多...
佰力博科技与您探讨半导体电阻测试常用的一些方法

佰力博科技与您探讨半导体电阻测试常用的一些方法

一、两探针法​ 两探针法是一种较为基础的测试方法。该方法将两根探针与半导体样品表面紧密接触&#xff0c;通过电源在两根探针之间施加电压&#xff0c;同时使用电流表测量通过样品的电流&#xff0c;再根据欧姆定律计算电阻。​这种方法的优点在于操作简单、设备要求较低&a…
阅读更多...
机器学习的一些基本概念

机器学习的一些基本概念

看了b站一个清华博士的视频做的笔记&#xff0c;对于人工智能的底层原理&#xff0c;训练方式&#xff0c;以及生成式文本输出&#xff0c;图片生成的底层原理有了一个了解&#xff0c;算是一个还不错的科普文。之前一直想要了解一下机器学习的入门原理&#xff0c;神经网络相关…
阅读更多...
Python爬虫实战：研究Grab 框架相关技术

Python爬虫实战：研究Grab 框架相关技术

1. 引言 1.1 研究背景与意义 随着互联网的快速发展，网络上的数据量呈爆炸式增长。如何高效地获取和利用这些数据成为了当前的研究热点。网络爬虫作为一种自动获取网页内容的技术，能够按照一定的规则，自动地抓取万维网信息，在搜索引擎、数据挖掘、信息整合等领域有着广泛的…
阅读更多...
uniapp 嵌入鸿蒙原生组件 具体步骤

uniapp 嵌入鸿蒙原生组件 具体步骤

关于怎么使用uniapp 嵌入鸿蒙原生组件 HBuilder X 版本 4.64 app-harmony文件下新建 index.uts button.ets button.ets里面复制uniapp 官方提供的 示例代码 https://uniapp.dcloud.net.cn/tutorial/harmony/native-component.html button.ets import { NativeEmbedBuilderO…
阅读更多...
阿里云 OS Copilot 使用指南

阿里云 OS Copilot 使用指南

安装&#xff1a; AlibabaCloudLinux: sudo yum install -y os-copilotUbuntu&#xff1a; curl -#S https://mirrors.aliyun.com/os-copilot/os-copilot-all-in-one-latest.sh | bash添加RAM用户 打开 https://ram.console.aliyun.com/users 复制AccessKey&#xff0c;Ac…
阅读更多...
枚举类扩充处理

枚举类扩充处理

问题背景 由于 Java 不允许枚举继承另一个枚举&#xff08;enum cannot extend enum&#xff09;&#xff0c;但可以通过 组合方式 或 工具类 来实现类似功能。 ✅ 解决方案一&#xff1a;组合方式引入原始枚举值 示例代码&#xff1a; public enum CustomErrorCodeEnum imp…
阅读更多...
Spring Security探索与应用

Spring Security探索与应用

Spring Security核心概念 框架定位与核心能力 Spring Security是Spring生态中实现应用级安全的核心框架，其官方定义为"强大且高度可定制的认证与访问控制框架"。作为Spring应用程序安全防护的事实标准解决方案，它通过模块化设计提供以下核心能力： 认证(Authenti…
阅读更多...
蓝桥杯国14 不完整的算式

蓝桥杯国14 不完整的算式

&#xff01;&#xff01;&#xff01;&#xff01;&#xff01;&#xff01;&#xff01;&#xff01;&#xff01;&#xff01;&#xff01;&#xff01;&#xff01;&#xff01;&#xff01;理清思路 然后一步步写 问题描述 小蓝在黑板上写了一个形如 AopBC 的算式&#x…
阅读更多...
扫描电镜：打开微观世界的“超维相机“

扫描电镜：打开微观世界的“超维相机“

当你用手机拍摄一朵花的微距照片时&#xff0c;放大100倍已足够惊艳。但如果告诉你&#xff0c;科学家手中的"相机"能将物体放大百万倍&#xff0c;连病毒表面的蛋白突触都清晰可见&#xff0c;你是否会好奇这背后的黑科技&#xff1f;这把打开微观宇宙的钥匙&#x…
阅读更多...
最新文章

Copyright @ 2022~2023