1. CUDA 内存模型解决什么问题CUDA 程序通常包含两类处理器HostCPU 及其内存DeviceGPU 及其内存CPU 和 GPU 可能同时执行任务并且 GPU 上有大量线程并行访问数据。因此我们需要回答几个问题数据应该放在哪种内存中哪些线程可以访问这份数据一个线程写入的数据其他线程什么时候能看到多个线程同时访问同一数据时如何避免错误如何让内存访问尽可能快CUDA 内存模型就是用来描述这些规则的。2. CUDA 内存层次CUDA 提供了多种内存空间。它们在容量、速度、生命周期和访问范围上各不相同。GPU | ------------------------ | | 寄存器 片上缓存 每线程私有 L1 Cache / L2 Cache | | ------------------------ | 全局内存 容量大但相对较慢 | ------------------------ | | 共享内存 常量/纹理内存 Block 内线程共享 适合特定访问模式2.1 寄存器Register每个线程私有。速度最快。容量很小。用于保存局部变量、中间计算结果等。线程结束后寄存器中的数据也就失效。__global__voidadd(int*c){intvalue10;// 通常优先放在寄存器中c[0]value;}寄存器使用过多会导致线程可并发数量下降。如果寄存器不够编译器可能把部分变量放到局部内存中。2.2 局部内存Local Memory局部内存的逻辑作用域是线程私有但它通常位于设备全局内存中并不是真正的片上高速内存。可能使用局部内存的情况包括线程的局部数组太大。寄存器空间不足。某些复杂的局部变量无法放入寄存器。因此局部内存虽然“私有”但访问速度可能较慢。2.3 共享内存Shared Memory同一个 Thread Block 内的线程可以访问。生命周期通常与 Block 相同。位于 GPU 片上速度较快。容量比全局内存小。常用于线程之间共享中间结果。__global__voidexample(float*data){__shared__floattile[256];inttidthreadIdx.x;tile[tid]data[tid];__syncthreads();// 确保 Block 内线程都完成写入// 读取其他线程写入的共享数据data[tid]tile[(tid1)%256];}共享内存只能在同一个 Block 内直接协作。不同 Block 之间不能直接依靠共享内存通信。2.4 全局内存Global MemoryGPU 上容量最大的内存。Kernel 中的所有线程通常都可以访问。Kernel 执行结束后数据仍然可以保留。访问延迟较高访问模式对性能影响很大。典型用途输入数据。输出数据。多个 Kernel 之间需要传递的数据。2.5 常量内存Constant Memory设备端只读。适合保存较小的常量数据。当一个 Warp 中的线程读取相同地址时通常效率较高。适合查找表、模型参数和固定配置等数据。2.6 纹理内存Texture Memory纹理内存适合具有特定访问模式的数据读取例如图像数据。空间局部性较强的数据。需要纹理缓存或特殊寻址方式的场景。现代 CUDA 程序中很多场景也可以通过只读缓存和普通全局内存实现因此是否使用纹理内存要根据具体访问模式决定。3. 线程与内存访问范围理解 CUDA 内存时可以同时关注两个维度数据存在哪里哪些线程可以访问内存空间访问范围生命周期典型特点寄存器单个线程线程执行期间最快但容量小局部内存单个线程线程执行期间逻辑私有物理上通常在全局内存共享内存同一个 BlockBlock 执行期间快适合 Block 内协作全局内存设备上的多个线程较长可跨 Kernel容量大但访问较慢常量内存Kernel 中的线程只读较长适合只读常量主机内存CPU 访问部分情况可映射由主机管理CPU 与 GPU 数据交换可以用下面的规则快速判断只给当前线程使用优先考虑寄存器。Block 内线程反复共享考虑共享内存。多个 Kernel 或大量线程都需要访问使用全局内存。所有线程读取少量固定数据考虑常量内存。4. CUDA 内存一致性与可见性4.1 什么是可见性假设线程 A 修改了共享内存中的数据线程 Ashared_value 42 线程 B读取 shared_value线程 B 是否一定能读到42取决于几个条件线程 A 是否已经完成写入。写入结果是否已经对线程 B 可见。线程 B 是否在正确的时间读取。是否存在其他线程同时修改该数据。因此“写入完成”和“其他线程能够安全读取”不是同一件事。4.2 数据竞争当多个线程访问同一个内存位置并且至少有一个线程执行写操作同时没有正确的同步机制就可能产生数据竞争。例如__global__voidwrong(int*result){result[0]1;}如果多个线程同时执行result[0] 1这个操作实际上可以拆成读取 result[0] 加 1 写回 result[0]多个线程可能同时读取相同的旧值最后写回的结果会覆盖其他线程的更新。4.3 Kernel 结束通常提供一个执行边界在默认执行方式下Kernel 内部的线程可能并行执行。一个 Kernel 完成后再启动的下一个 Kernel 通常可以看到前一个 Kernel 写入的全局内存结果。因此常见做法是把需要全局协作的计算拆成多个 KernelKernel 1生成中间结果 ↓ Kernel 1 完成 ↓ Kernel 2读取中间结果并继续计算不同 Block 之间通常不能在同一个 Kernel 中随意使用普通读写完成全局同步。5. 同步机制与内存栅栏5.1__syncthreads()__syncthreads()是 Block 内的线程屏障。它通常用于保证Block 内所有线程都到达这个位置。之前对共享内存的写入对 Block 内线程可见。示例__global__voidcopy_tile(constfloat*input,float*output){__shared__floattile[128];inttidthreadIdx.x;tile[tid]input[tid];__syncthreads();output[tid]tile[127-tid];}如果没有__syncthreads()某个线程可能在其他线程完成写入前就开始读取共享内存。5.2__syncthreads()的边界__syncthreads()只能同步同一个 Block 内的线程能同步 Block 内线程。不能同步不同 Block。不能让尚未到达屏障的线程“自动完成工作”。需要特别注意如果同一个 Block 中部分线程进入屏障而其他线程因为分支没有进入可能导致程序无法正常继续。5.3 内存栅栏内存栅栏主要关注内存操作的顺序和可见性而不是让线程停下来等待。可以粗略理解为屏障大家先停下来等人到齐。栅栏保证我之前的内存操作不会被重新排到后面。栅栏不能替代线程屏障。如果既需要等待其他线程又需要保证内存可见性通常需要组合使用同步机制。5.4 原子操作原子操作可以保证一个特定的读-改-写过程不可被其他线程打断。__global__voidcount(int*counter){atomicAdd(counter,1);}这里每个线程都对counter加 1atomicAdd能避免多个线程同时更新造成覆盖。但原子操作也有代价大量线程竞争同一个地址时性能可能下降。原子操作只保护特定的更新操作。它不是通用的线程屏障。6. 内存访问模式与性能CUDA 的性能不仅取决于“使用哪种内存”还取决于线程如何访问内存。6.1 合并访问GPU 通常以 Warp 为单位执行线程。一个 Warp 中的线程如果访问连续的全局内存地址硬件更容易合并这些访问从而减少内存事务。较好的访问模式线程 0 → data[0] 线程 1 → data[1] 线程 2 → data[2] 线程 3 → data[3]较差的访问模式可能是线程 0 → data[0] 线程 1 → data[1024] 线程 2 → data[2048] 线程 3 → data[3072]实际效果还与数据类型、对齐方式、GPU 架构等因素有关但基本原则是让相邻线程尽量访问相邻数据。6.2 共享内存银行冲突共享内存被划分为多个 Bank。若同一个 Warp 中的多个线程同时访问同一个 Bank 的不同地址就可能产生银行冲突导致访问被分成多个步骤。常见优化方向调整数据布局。修改数组维度。使用填充元素改变地址分布。让同一个 Warp 中线程的访问尽量分散到不同 Bank。6.3 共享内存不是总比全局内存快共享内存适合数据会被多个线程重复使用。数据可以提前加载到片上。线程之间需要频繁交换中间结果。但如果数据只访问一次或者加载和同步成本很高使用共享内存未必有收益。优化时应结合实际测量而不是机械地把所有数据搬到共享内存。7. 主机与设备之间的数据传输7.1 传统显式传输最基本的方式是CPU 主机内存 | | cudaMemcpy v GPU 全局内存 | | Kernel 访问 v 计算结果 | | cudaMemcpy v CPU 主机内存典型流程在主机端准备输入数据。在设备端分配内存。将输入数据复制到设备。启动 Kernel。将结果复制回主机。释放设备内存。这种方式的数据位置清晰但需要开发者管理拷贝过程。7.2 统一内存Unified Memory统一内存为 CPU 和 GPU 提供统一的地址空间。开发者可以使用一个指针访问数据运行时负责在主机和设备之间迁移数据。优点编程简单。减少显式拷贝代码。适合原型开发和数据结构复杂的场景。缺点数据迁移可能产生额外开销。页面在 CPU 和 GPU 之间频繁迁移时性能可能下降。高性能程序通常仍需要关注数据预取和访问位置。7.3 异步传输通过 CUDA Stream可以让数据传输和 Kernel 执行在满足条件时重叠时间 → Stream 1拷贝数据 A → 执行 Kernel A Stream 2拷贝数据 B → 执行 Kernel B这有助于隐藏部分数据传输延迟但通常需要合适的 Stream 使用方式。可异步传输的主机内存例如页锁定内存。设备和传输路径支持并发。8. 一条完整的因果链可以用下面的思路分析 CUDA 内存问题选择数据位置 ↓ 确定哪些线程能够访问 ↓ 确定线程之间是否共享数据 ↓ 如果共享判断是否存在数据竞争 ↓ 选择屏障、栅栏或原子操作 ↓ 检查访问是否合并、是否存在银行冲突 ↓ 根据实际测量进行优化例如一个矩阵乘法 Kernel 可能这样设计矩阵数据放在全局内存中。每个 Block 把一小块矩阵加载到共享内存。Block 内线程通过共享内存重复使用数据。每次加载后使用__syncthreads()。每个线程使用寄存器保存自己的累加结果。全局内存访问尽量按照连续地址进行。9. 常见误解误解一局部内存一定很快“局部”描述的是访问范围不代表物理位置。局部内存通常位于设备全局内存可能比寄存器慢很多。误解二共享内存可以跨 Block 使用共享内存只在同一个 Block 内共享。不同 Block 之间需要通过全局内存、多个 Kernel 或其他专门机制协作。误解三__syncthreads()可以同步整个 Grid不能。它只能同步当前 Block 中的线程。误解四原子操作可以替代所有同步原子操作主要保护一个具体的更新操作不能自动保证一组复杂操作的整体顺序也不能让整个 Grid 的线程停下来等待。误解五统一内存一定比显式拷贝快统一内存更方便但运行时自动迁移数据可能产生开销。性能取决于访问模式、数据迁移次数和硬件环境。误解六把数据放进共享内存就一定更快如果数据不重复使用或者同步和加载开销较大共享内存可能反而降低性能。10. 总结CUDA 内存模型可以从四个核心问题理解数据在哪里寄存器、共享内存、全局内存、常量内存等。谁能访问单线程、同一 Block、整个设备或主机与设备。什么时候能看到需要考虑内存可见性、执行顺序和 Kernel 边界。如何访问更快合并访问、减少银行冲突、合理使用共享内存和寄存器。最重要的判断原则是先确定数据的访问范围和生命周期再选择内存类型如果多个线程共享数据再设计同步最后根据访问模式进行性能优化。
