Select a result to preview

CUDA2_GPU的内存体系及其优化指南

CUDA（二）：GPU的内存体系及其优化指南

原文作者：紫气东来
原文链接：https://zhuanlan.zhihu.com/p/654027980

概述

在冯·诺依曼架构的硬件中实现高性能计算，最重要的两点是：访存和计算。这两点分别对应着 IO bound 和 compute bound，硬件系统的内存体系深刻影响着这两点。本文深入探讨GPU的内存体系，并通过规约（Reduction）操作的实践案例，展示如何进行CUDA编程优化。

一、GPU的内存体系

1.1 各级内存及其特点

CUDA 内存模型结合了主机和设备的内存系统，具有完整的层次结构，可以显式地进行控制和优化。

1.1.1 全局内存（Global Memory）

特点：

主要用途：

使用方式：

// 数据传输
cudaMemcpy(d_x, h_x, M, cudaMemcpyHostToDevice);

// 静态声明
__device__ T x;        // 单个变量
__device__ T y[N];     // 固定长度的数组

1.1.2 常量内存（Constant Memory）

特点：

使用方式：

// 在核函数外定义
__constant__ T constant_data;

// 从主机复制数据到常量内存
cudaMemcpyToSymbol(constant_data, host_data, size);

1.1.3 纹理内存（Texture Memory）和表面内存（Surface Memory）

特点：

1.1.4 寄存器（Register）

特点：

使用方式：

// 核函数中不加任何限定符的变量通常存放在寄存器中
const int n = blockDim.x * blockIdx.x + threadIdx.x;
c[n] = a[n] + b[n];  // n 就是寄存器变量

可见性和生命周期：

1.1.5 局部内存（Local Memory）

特点：

注意： 使用过多局部内存会降低程序性能。

1.1.6 共享内存（Shared Memory）

特点：

使用方式：

__shared__ real s_y[128];  // 静态共享内存

可见性：

主要作用：

1.1.7 L1和L2缓存

结构：

用途：

物理结构：

1.1.8 内存特点对比表

内存类型 位置 访问速度 容量 作用域 生命周期
寄存器 片上 最快 很小 线程 线程
局部内存 片外（全局内存的一部分） 慢 较大（512KB/线程） 线程 线程
共享内存 片上 很快 有限 线程块 线程块
全局内存 片外 慢 很大 全局 应用程序
常量内存 片外（有缓存） 较快 64KB 全局 应用程序
纹理内存 片外（有缓存） 较快 较大 全局 应用程序

1.2 SM构成及典型GPU对比

1.2.1 SM（Streaming Multiprocessor）的组成

一个 GPU 由多个 SM 构成。一个 SM 包含以下资源：

存储资源：

计算资源：

1.2.2 主流GPU产品对比

GPU V100 A100 H100 L40S
架构 Volta Ampere Hopper Ada Lovelace
内存接口 4096-bit HBM2 5120-bit HBM2 5120-bit HBM3 GDDR6
内存大小 32GB/16GB 40GB 80GB 48GB
内存带宽 900 GB/s 1555 GB/s 3000 GB/s 864 GB/s
SM数量 80 108 132 142
纹理单元 320 432 528 576
L2缓存大小 6144 KB 40 MB 50 MB 96 MB
共享内存/SM 最高96KB 最高164KB 最高228KB 最高128KB
寄存器文件/SM 256 KB 256 KB 256 KB 256 KB
峰值FP16算力 31.4 TFLOPS 78 TFLOPS 120 TFLOPS 90.52 TFLOPS

关键观察：

1.3 GPU之外：近存计算与存算一体

为了进一步降低访存成本，获得更高的性能，近存计算与存算一体逐渐成为热门方向。

1.3.1 近存计算：Graphcore IPU

架构特点：

计算模式：

1.3.2 存算一体：后摩智能 H30

核心思想：
通过对存储器单元本身进行算法嵌入，使得计算可以在存储器单元内完成。

特点：

二、通过规约（Reduction）操作理解GPU内存体系

2.0 规约算法概述

Reduce 算法可以描述为：

x = x_0 ⊗ x_1 ⊗ x_2 ⊗ x_3 ... ⊗ x_{n-1} ⊗ x_n

其中 ⊗ 可表示为求 sum、min、max、avg 等操作，最后获得的输出相比于输入一般维度上会递减。

GPU中的树形计算方式：
在GPU中，reduce采用了一种树形的计算方式。由于GPU没有针对global数据的同步操作，只能针对block的数据进行同步，因此一般将reduce分为两个阶段：

  1. 第一阶段：开启m个block计算出m个小份的reduce值
  2. 第二阶段：使用一个block将m个小份再次进行reduce，得到最终结果

实验代码开源地址： https://github.com/ifromeast/cuda_learning/tree/main/02_reduce

2.1 仅使用全局内存实现规约

2.1.1 折半归约法

对于数组归约的并行计算问题，要从一个数组出发，最终得到一个数。假如数组元素个数是2的整数次方，可以将数组后半部分的各个元素与前半部分对应的数组元素相加。如果重复此过程，最后得到的第一个数组元素就是最初的数组中各个元素的和。这就是所谓的折半归约（binary reduction）法。

2.1.2 实现代码

void __global__ reduce_global(real *d_x, real *d_y)
{
    const int tid = threadIdx.x;
    real *x = d_x + blockDim.x * blockIdx.x;

    for (int offset = blockDim.x >> 1; offset > 0; offset >>= 1)
    {
        if (tid < offset)
        {
            x[tid] += x[tid + offset];
        }
        __syncthreads();
    }

    if (tid == 0)
    {
        d_y[blockIdx.x] = x[0];
    }
}

2.1.3 关键点说明

  1. 同步函数 __syncthreads()：保证一个线程块中的所有线程在执行该语句后面的语句之前都完全执行了该语句前面的语句

  2. 指针操作：real *x = d_x + blockDim.x * blockIdx.x; 将数组 d_x 中第 blockDim.x * blockIdx.x 个元素的地址赋给指针 x

  3. 独立归约：for 循环内在各个线程块内对其中的数据独立地进行归约，不同线程块之间不需要同步

  4. 位操作优化：在 offset 的计算过程中使用了位操作，对于2的幂而言更加高效

  5. 部分归约：该核函数仅将一个长度为 10^8 的数组 d_x 归约到一个长度为 10^8/128 的数组 d_y

  6. 内存限制：因为 __global__ 修饰符的限制，d_y 和 x 均是全局内存的变量

2.2 使用共享内存实现规约

2.2.1 优化动机

由于全局内存访问速度最低，因此性能较低。本节使用对整个线程块可见的共享内存来实现同样的规约操作。

2.2.2 实现方式

使用 __shared__ 修饰共享内存变量 s_y，长度等于线程块大小。然后将全局内存中的数据复制到共享内存中，每个线程块都有一个共享内存变量的副本。

const int tid = threadIdx.x;
const int idx = blockIdx.x * blockDim.x + threadIdx.x;
__shared__ real s_y[128];
s_y[tid] = (idx < N) ? d_x[idx] : 0.0;
__syncthreads();

2.2.3 资源分配说明

GPU需要分配两种资源：

  1. 存储资源：共享内存数组
  2. 计算资源：根据thread数量确定
    • 一个block中分配128个thread线程
    • 32个线程为一组（即一个 warp）
    • 绑定在一个SIMD单元
    • 128个线程 = 4组SIMD单元

2.2.4 使用共享内存的优势

  1. 不再要求全局内存数组的长度 N 是线程块大小的整数倍
  2. 在规约的过程中不会改变全局内存数组中的数据（在仅使用全局内存时，数组 d_x 中的部分元素被改变）
  3. 访问速度更快：在核函数中对共享内存访问的次数越多，加速效果越明显

2.3 使用动态共享内存实现规约

2.3.1 动态共享内存的必要性

在使用共享内存数组时，如果指定了一个固定的长度（如128），这种静态的方式可能会导致错误的发生，因此有必要使用动态操作。

2.3.2 修改方法

修改1：执行配置中添加第三个参数

<<<grid_size, block_size, sizeof(real) * block_size>>>

前两个参数分别是网格大小和线程块大小，第三个参数是核函数中每个线程块需要定义的动态共享内存的字节数，其默认值为零。

修改2：改变共享内存变量的声明方式

extern __shared__ real s_y[];

与静态共享内存的声明方式有两点不同：

  1. 必须加上限定词 extern
  2. 不能指定数组大小

2.3.3 性能对比

编译运行代码：

nvcc reduce_gpu.cu -o reduce

通过 nvprof 命令查看GPU各部分的耗时：

nvprof ./reduce

单精度结果（由于该例子中不存在对内存的频繁读写，因此几种方式的性能差别不大）：

Type  Time(%)      Time     Calls       Avg       Min       Max  Name
GPU activities:   
  98.25%  30.8438s       300  102.81ms  80.686ms  259.15ms  [CUDA memcpy HtoD]
   0.44%  137.86ms       100  1.3786ms  1.3756ms  1.3822ms  reduce_global(float*, float*)
   0.44%  137.66ms       300  458.85us  343.93us  863.09us  [CUDA memcpy DtoH]
   0.43%  136.51ms       100  1.3651ms  1.3647ms  1.3666ms  reduce_shared(float*, float*)
   0.43%  136.50ms       100  1.3650ms  1.3647ms  1.3664ms  reduce_dynamic(float*, float*)

双精度结果（放大访存速度的差别）：

编译方式：

nvcc reduce_gpu.cu -DUSE_DP -o reduce_dp

此时可以看到全局内存的性能出现了明显下降：

Type  Time(%)      Time     Calls       Avg       Min       Max  Name
GPU activities:   
  98.65%  60.8587s       300  202.86ms  189.54ms  304.60ms  [CUDA memcpy HtoD]
   0.65%  398.54ms       300  1.3285ms  1.2460ms  2.9135ms  [CUDA memcpy DtoH]
   0.26%  157.32ms       100  1.5732ms  1.5677ms  1.5803ms  reduce_global(double*, double*)
   0.22%  137.32ms       100  1.3732ms  1.3716ms  1.3746ms  reduce_shared(double*, double*)
   0.22%  137.31ms       100  1.3731ms  1.3714ms  1.3750ms  reduce_dynamic(double*, double*)

2.4 规约的其他优化方法

2.4.1 使用原子函数

优化动机：

在前边几个版本的数组归约函数中，核函数并没有做全部的计算，而只是将一个长一些的数组 d_x 变成了一个短一些的数组 d_y。所有操作所用时间约为 7.5 ms（单精度），而在 GPU 中实际计算的时间仅为 1.4ms 左右。

如果能在 GPU 中计算出最终结果，则有望显著地减少整体的计算时间。有两种方法：

  1. 用另一个核函数将较短的数组进一步归约
  2. 在先前的核函数的末尾利用原子函数进行归约，直接得到最终结果

问题分析：

如果简单地将代码改为：

if (tid == 0) {
    d_y[0] += s_y[0];
}

这会导致错误，因为：

原子函数解决方案：

if (tid == 0)
{
    atomicAdd(d_y, s_y[0]);
}

原子函数特点：

性能提升：

使用原子函数后总时间变为 2.8 ms，相比于之前的方式，性能提升接近3倍：

Type  Time(%)      Time     Calls       Avg       Min       Max  Name
GPU activities:   
  66.95%  194.36ms       100  1.9436ms  1.8185ms  2.0081ms  reduce(float const *, float*, int)
  33.00%  95.797ms       101  948.48us  1.4400us  95.642ms  [CUDA memcpy HtoD]
   0.05%  133.09us       100  1.3300us  1.2160us  2.6560us  [CUDA memcpy DtoH]

分析：

2.4.2 使用线程束函数与协作组

线程束（Warp）基础：

优化1：使用线程束同步函数

当所涉及的线程都在一个线程束内时，可以将线程块同步函数 __syncthreads 换成更加廉价的线程束同步函数 __syncwarp：

for (int offset = blockDim.x >> 1; offset >= 32; offset >>= 1)
{
    if (tid < offset)
    {
        s_y[tid] += s_y[tid + offset];
    }
    __syncthreads();
}

for (int offset = 16; offset > 0; offset >>= 1)
{
    if (tid < offset)
    {
        s_y[tid] += s_y[tid + offset];
    }
    __syncwarp();
}

优化2：使用线程束洗牌函数

函数 __shfl_down_sync 的作用是将高线程号的数据平移到低线程号中去，这正是归约问题中需要的操作：

for (int offset = 16; offset > 0; offset >>= 1)
{
    y += __shfl_down_sync(FULL_MASK, y, offset);
}

相比之前的版本，有两处不同：

  1. 在进行线程束内的循环之前，将共享内存中的数据复制到了寄存器。因为寄存器一般来说比共享内存更高效
  2. 去掉了束同步函数，因为洗牌函数能够自动处理同步与读-写竞争问题

优化3：使用协作组（Cooperative Groups）

协作组可以看作是线程块和线程束同步机制的推广，提供了更为灵活的线程协作方式。

使用协作组需要包含头文件：

#include <cooperative_groups.h>
using namespace cooperative_groups;

可以用函数 tiled_partition 将一个线程块划分为若干片（tile），每一片构成一个新的线程组。目前仅可以将片的大小设置为 2 的正整数次方且不大于 32（即 2、4、8、16 和 32）。

real y = s_y[tid];

thread_block_tile<32> g = tiled_partition<32>(this_thread_block());
for (int i = g.size() >> 1; i > 0; i >>= 1)
{
    y += g.shfl_down(y, i);
}

性能对比：

Type  Time(%)      Time     Calls       Avg       Min       Max  Name
GPU activities:   
  29.49%  190.45ms       100  1.9045ms  1.8181ms  2.0072ms  reduce_syncwarp(float const *, float*, int)
  27.84%  179.82ms       100  1.7982ms  1.7960ms  1.8183ms  reduce_shfl(float const *, float*, int)
  27.82%  179.65ms       100  1.7965ms  1.7957ms  1.7976ms  reduce_cp(float const *, float*, int)
  14.80%  95.571ms       301  317.51us  1.4390us  95.122ms  [CUDA memcpy HtoD]
   0.06%  384.69us       300  1.2820us  1.2150us  1.7920us  [CUDA memcpy DtoH]

这3种方式的性能相比之前的方式都有所提高。

2.4.3 进一步分析和优化

线程利用率问题：

在前边的例子中，使用大小为 128 的线程块：

优化策略：在归约之前累加多个数据

用一个寄存器变量 y，在循环体中对读取的全局内存数据进行累加：

real y = 0.0;
const int stride = blockDim.x * gridDim.x;
for (int n = bid * blockDim.x + tid; n < N; n += stride)
{
    y += d_x[n];
}
s_y[tid] = y;
__syncthreads();

包装函数实现：

real reduce(const real *d_x)
{
    const int ymem = sizeof(real) * GRID_SIZE;
    const int smem = sizeof(real) * BLOCK_SIZE;

    real h_y[1] = {0};
    real *d_y;
    CHECK(cudaMalloc(&d_y, ymem));

    reduce_cp<<<GRID_SIZE, BLOCK_SIZE, smem>>>(d_x, d_y, N);
    reduce_cp<<<1, 1024, sizeof(real) * 1024>>>(d_y, d_y, GRID_SIZE);

    CHECK(cudaMemcpy(h_y, d_y, sizeof(real), cudaMemcpyDeviceToHost));
    CHECK(cudaFree(d_y));

    return h_y[0];
}

这里，将GRID_SIZE取为10240，将BLOCK_SIZE取为128。当数据量为 N = 100000000 时，在归约前每个线程将先累加几十个数据。

性能结果：

完整计算时间只需要0.85ms，GPU上核函数的计算时间极大缩减：

Type  Time(%)      Time     Calls       Avg       Min       Max  Name
GPU activities:   
  62.32%  95.388ms         1  95.388ms  95.388ms  95.388ms  [CUDA memcpy HtoD]
  37.59%  57.529ms       200  287.64us  6.4000us  572.79us  reduce_cp(float const *, float*, int)
   0.09%  135.46us       100  1.3540us  1.3110us  2.5600us  [CUDA memcpy DtoH]

进一步优化：使用静态全局内存

在上面的包装函数中，需要为数组 d_y 分配与释放设备内存。实际上，设备内存的分配与释放是比较耗时的。一种优化方案是使用静态全局内存代替动态全局内存，因为静态内存是编译期间就会分配好的。

__device__ real static_y[GRID_SIZE];

real reduce(const real *d_x)
{
    real *d_y;
    CHECK(cudaGetSymbolAddress((void**)&d_y, static_y));

    const int smem = sizeof(real) * BLOCK_SIZE;

    reduce_cp<<<GRID_SIZE, BLOCK_SIZE, smem>>>(d_x, d_y, N);
    reduce_cp<<<1, 1024, sizeof(real) * 1024>>>(d_y, d_y, GRID_SIZE);

    real h_y[1] = {0};
    CHECK(cudaMemcpy(h_y, d_y, sizeof(real), cudaMemcpyDeviceToHost));

    return h_y[0];
}

计算时间从 0.85ms 缩短到了 0.6 ms。

2.5 优化方法总结

2.5.1 各种方法的性能对比

计算方法 计算结果 计算时间(ms) 单次加速比 累计加速比
CPU 33554432.0 600 1 1
GPU (全局内存) 123633392.0 7.2 83 83
GPU (静态共享内存) 123633392.0 7.2 1 83
GPU (动态共享内存) 123633392.0 7.2 1 83
GPU (原子函数) 123633392.0 2.2 3.3 274
GPU (束同步函数) 123633392.0 2.2 1 274
GPU (洗牌函数) 123633392.0 2.1 1.05 288
GPU (协作组) 123633392.0 2.1 1 288
GPU (增大线程利用率) 123000064.0 0.85 2.5 719
GPU (静态全局内存) 123000064.0 0.6 1.4 1006

2.5.2 关键优化技术总结

  1. 使用共享内存代替全局内存

    • 共享内存位于片上，访问速度远快于全局内存
    • 适用于需要频繁访问的数据

  2. 使用原子函数避免数据竞争

    • 保证操作的原子性
    • 减少主机与设备之间的数据传输

  3. 使用线程束级别的同步和洗牌函数

    • __syncwarp() 比 __syncthreads() 更轻量
    • 洗牌函数自动处理同步问题
    • 避免线程束分化

  4. 提高线程利用率

    • 在归约前让每个线程累加多个数据
    • 减少归约过程中的线程闲置

  5. 使用静态全局内存

    • 避免运行时的内存分配和释放开销
    • 编译期分配，效率更高

三、GPU内存优化的核心原则

3.1 内存访问模式优化

3.1.1 合并访问（Coalesced Access）

定义：
当一个线程束（warp）中的32个线程访问连续的内存地址时，这些访问可以被合并成少数几个内存事务，大大提高内存带宽利用率。

最佳实践：

示例：

// 好的访问模式 - 合并访问
int tid = threadIdx.x + blockIdx.x * blockDim.x;
float value = array[tid];

// 差的访问模式 - 跨步访问
int tid = threadIdx.x + blockIdx.x * blockDim.x;
float value = array[tid * stride];  // stride > 1 会导致非合并访问

3.1.2 Bank冲突（Bank Conflict）

共享内存的Bank结构：

避免Bank冲突的方法：

  1. 确保同一warp中的线程访问不同的bank
  2. 使用padding技术
  3. 调整数据布局

示例：

// 会产生bank冲突
__shared__ float shared[32][32];
float value = shared[threadIdx.x][threadIdx.y];

// 避免bank冲突 - 添加padding
__shared__ float shared[32][33];  // 多一列
float value = shared[threadIdx.x][threadIdx.y];

3.2 内存层次使用策略

3.2.1 寄存器优先原则

3.2.2 共享内存作为缓存

典型模式：

__shared__ float tile[TILE_SIZE][TILE_SIZE];

// 1. 从全局内存加载到共享内存
tile[ty][tx] = global_data[gid];
__syncthreads();

// 2. 在共享内存中进行计算
float result = 0;
for (int i = 0; i < TILE_SIZE; i++) {
    result += tile[ty][i] * tile[i][tx];
}
__syncthreads();

// 3. 写回全局内存
global_result[gid] = result;

3.2.3 常量内存的使用场景

3.3 内存带宽优化

3.3.1 计算访存比（Compute-to-Memory Ratio）

定义：
计算访存比 = 算术操作数量 / 内存访问数量

优化策略：

3.3.2 隐藏内存延迟

方法：

  1. 增加占用率（Occupancy）：让更多的线程块同时运行
  2. 指令级并行（ILP）：在一个线程中执行多个独立操作
  3. 异步内存操作：使用异步拷贝指令

3.4 数据布局优化

3.4.1 AoS vs SoA

AoS（Array of Structures）：

struct Particle {
    float x, y, z;
    float vx, vy, vz;
};
Particle particles[N];

SoA（Structure of Arrays）：

struct Particles {
    float x[N], y[N], z[N];
    float vx[N], vy[N], vz[N];
};

选择原则：

3.4.2 数据对齐

四、性能分析方法

4.1 使用nvprof进行性能分析

4.1.1 基本使用

# 基本性能分析
nvprof ./program

# 详细的内存分析
nvprof --print-gpu-trace ./program

# 分析特定的指标
nvprof --metrics gld_efficiency,gst_efficiency ./program

4.1.2 关键指标

内存相关指标：

计算相关指标：

4.2 使用Nsight Compute

Nsight Compute 是NVIDIA的新一代性能分析工具，提供更详细的分析。

# 基本分析
ncu ./program

# 详细分析特定kernel
ncu --kernel-name myKernel --launch-skip 0 --launch-count 1 ./program

# 生成报告
ncu --export report ./program

4.3 使用Nsight Systems

用于系统级别的性能分析，可以看到CPU和GPU的时间线。

# 生成性能报告
nsys profile -o report ./program

# 在GUI中查看
nsys-ui report.qdrep

4.4 性能优化流程

  1. 识别瓶颈

    • 使用profiler找出最耗时的kernel
    • 分析是内存受限还是计算受限

  2. 针对性优化

    • 内存受限：优化内存访问模式
    • 计算受限：提高指令级并行度

  3. 验证优化效果

    • 重新运行profiler
    • 对比优化前后的指标

  4. 迭代优化

    • 继续寻找新的瓶颈
    • 重复优化过程

五、实践指导与最佳实践

5.1 内存优化检查清单

全局内存优化：

共享内存优化：

寄存器优化：

5.2 常见陷阱与解决方案

5.2.1 线程束分化

问题：

if (threadIdx.x < 16) {
    // 分支A
} else {
    // 分支B
}

解决方案：

5.2.2 过度使用共享内存

问题：
使用过多共享内存会降低占用率。

解决方案：

5.2.3 忽略内存对齐

问题：
未对齐的内存访问会导致性能下降。

解决方案：

// 确保数据结构对齐
struct __align__(16) MyStruct {
    float4 data;
};

5.3 不同应用场景的优化策略

5.3.1 矩阵运算

5.3.2 规约操作

5.3.3 卷积操作

六、总结与展望

6.1 核心要点回顾

  1. GPU内存层次结构

    • 寄存器：最快，容量最小，线程私有
    • 共享内存：很快，容量有限，线程块共享
    • 全局内存：最慢，容量最大，所有线程可访问
    • L1/L2缓存：自动管理，不可编程

  2. 优化原则

    • 最大化内存带宽利用率
    • 最小化内存访问延迟
    • 提高计算访存比
    • 避免线程束分化

  3. 实践技巧

    • 使用合并访问
    • 避免bank冲突
    • 合理使用共享内存
    • 利用原子操作和warp-level函数

6.2 性能优化的一般流程

  1. Profile First：先分析，找出真正的瓶颈
  2. Optimize Bottleneck：针对瓶颈进行优化
  3. Measure Again：验证优化效果
  4. Iterate：持续迭代优化

6.3 未来发展趋势

  1. 更大的片上内存

    • 共享内存容量持续增长
    • L2缓存大幅增加

  2. 新的内存技术

    • HBM3及更高带宽的内存
    • 近存计算和存算一体

  3. 编程模型的演进

    • 更高级的抽象
    • 自动优化工具的改进

6.4 学习资源

官方文档：

推荐书籍：

在线资源：

参考资料

[1] 《CUDA C编程权威指南》程润伟, Max Grossman, Ty McKercher 著，颜成钢, 殷建, 李亮 译，机械工业出版社，2017-6

[2] brucefan1983/CUDA-Programming: Sample codes for my CUDA programming book (github.com)

[3] https://developer.download.nvidia.cn/assets/cuda/files/reduction.pdf

[4] 深度了解 NVIDIA Ampere 架构

[5] 有了琦琦的棍子：深入浅出GPU优化系列：reduce优化

[6] NVIDIA Hopper Architecture In-Depth | NVIDIA Technical Blog

[7] NVIDIA Ada GPU Architecture Whitepaper

[8] GitHub - Liu-xiandong/How_to_optimize_in_GPU

[9] GitHub - BBuf/how-to-optim-algorithm-in-cuda

[10] ifromeast/cuda_learning: learning how CUDA works (github.com)

注： 本文档基于知乎文章"CUDA（二）：GPU的内存体系及其优化指南"整理而成，包含了原文的核心内容，并补充了实践指导和最佳实践建议。