关于CUDA等从硬件到软件的记录

1. 前言

砚上三五笔，落墨鹧鸪啼

本文主要记录：最近要在CUDA的运算分配上做一个比较详细的分析，看看运算效率是怎么计算的。
如有不对，欢迎评论区指正！

2. 正文

2.1 CUDA

2.1.1 计算层面

统一计算设备架构CUDA（Compute Unified Device Architecture）, NVIDIA的并行架构，可以理解为一整套GPU软件开发的操作系统。

线程（thread）：是基本执行单元，有硬件支持，每个线程执行相同代码
线程块（block）：是若干线程的分组，Block内一个块至多512个线程、或1024个线程（根据不同的GPU规格），线程块可以是一维、二维或者三维的
线程网络（grid）：是若干线程块的网格，Grid是一维和二维的。

我们会把一组thread归为一个block，而block又会被组织成一个grid。

GPU上有很多计算核心也就是Streaming Multiprocessor (SM)，在具体的硬件执行中，一个SM会同时执行一组线程，在CUDA里叫warp，直接可以理解这组硬件线程会在这个SM上同时执行一部分指令，这一组的数量一般为32或者64个线程。假如一个SM同时能执行64个线程，但一个block有1024个线程，那这1024个线程是分1024/64=16次执行。

Thread Block：程序里定义的一组线程。
Warp：每 32 个线程会被组成一个 warp（硬件调度的基本单位）。
Warp Scheduler：把 warp 分配给 Streaming Multiprocessor (SM)。
SM：包含很多 CUDA cores，warp 的线程会被分配到这些 cores 上执行。

2.1.2 硬件层面

GPU主要由显存和计算单元组成：
显存（Global Memory）：显存是在GPU板卡上的DRAM，类似于CPU的内存，就是那堆DDR啊，GDDR5啊之类的。特点是容量大（可达16GB），速度慢，CPU和GPU都可以访问。
计算单元（Streaming Multiprocessor）：执行计算的。每一个SM都有自己的控制单元（Control Unit），寄存器（Register），缓存（Cache），指令流水线（execution pipelines）。

对于CPU可能会说有几个核，比如8核CPU，16核CPU。但GPU 处理器由数百或数千个并行运行以执行复杂图形操作的小核心或单元组成。英伟达将这些核心或处理器称为 Cuda cores，而 AMD/ATI 将其称为流处理器（Stream Processor）。
下图为GPU的显存颗粒：

DDR3 RAM。这是一种用于计算机和笔记本电脑的标准类型 RAM。它比旧主板上长期使用的旧型 DDR2 RAM 速度快得多。与 GDDR RAM 相比，它们也便宜得多。与 DDR2 相比，DDR3 还能在更低的电压下运行，这为您节省了一些功耗。

比如NVIDIA的Pascal架构的芯片GP100，包含一组 GPC（图形处理簇，Graphics Processing Clusters）、TPC（纹理处理簇，Texture Processing Clusters）、SM（流多处理器，Stream Multiprocessors）以及内存控制器。：

一块完成的芯片（仅上图所示的，方便说明），包含60个流多处理器（SM），每个SM包含64个CUDA 核心，那么整块显卡一共有3840和CUDA 核心，但其实实际上并不会有效使用60个SM。

其中每一个SM可以单独来看：

图中为一个 SM 的架构。其中绿色的“Core” 为单精度 CUDA 核心，共有 64 个，同时支持 32 位单精度浮点计算和 16 位半精度浮点计算，其中 16 位计算吞吐是 32 位计算吞吐的两倍。图中橘黄色的 “DP Unit” 为双精度计算单元，支持 64 位双精度浮点计算，数量为 32 个。每个 GP100 SM 双精度计算吞吐为单精度的一半。下面这张处理性能表证实了这一点。

2.1.3 API使用

比如向量加法：c=a+b
在C++中，这是一个比较直观的运算操作流程。

#include <stdio.h>
#include <cuda_runtime.h>
 
// CUDA 核kernel函数：每个线程负责加一个元素
__global__ void vectorAdd(int *a, int *b, int *c, int n) {
    int i = threadIdx.x;  // 获取当前线程在 block 中的编号（0 ~ N-1）
    if (i < n)
        c[i] = a[i] + b[i];
}
 
int main() {
    const int N = 10;
    int a[N], b[N], c[N];               // 在 CPU（host）上的数组
    int *d_a, *d_b, *d_c;               // GPU（device）上的指针
 
    // 初始化数据
    for (int i = 0; i < N; i++) {
        a[i] = i;
        b[i] = i * 2;
    }
 
    // 1. 分配 GPU 内存
    cudaMalloc((void**)&d_a, N * sizeof(int));
    cudaMalloc((void**)&d_b, N * sizeof(int));
    cudaMalloc((void**)&d_c, N * sizeof(int));
 
    // 2. 拷贝数据到 GPU
    cudaMemcpy(d_a, a, N * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, b, N * sizeof(int), cudaMemcpyHostToDevice);
 
    // 3. 启动 Kernel    <<<blocks, threads>>> 启动 kernel   	设置并行维度、执行函数
    vectorAdd<<<1, N>>>(d_a, d_b, d_c, N);
 
    // 4. 把结果拷回 CPU
    cudaMemcpy(c, d_c, N * sizeof(int), cudaMemcpyDeviceToHost);
 
    // 5. 打印结果
    printf("Result:\n");
    for (int i = 0; i < N; i++)
        printf("%d + %d = %d\n", a[i], b[i], c[i]);
 
    // 6. 释放 GPU 内存
    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
 
    return 0;
}

Kernel 是 CUDA 程序中由 GPU 执行的函数，它是并行计算的核心。

__global__ void myKernel(...) {
    // 每个线程执行这段代码
}

CPU（Host）和 GPU（Device）之间数据传输。

2.1.4 kernel和core

从上面也不难看出来，CUDA core指的就是SM里面的计算单元。而CUDA kernel其实是软件层面的概念。在 CUDA 编程里，你写的 __global__ void myKernel(...) { ... } 就是一个 kernel 函数。当你调用 kernel 时，比如：
myKernel<<<gridDim, blockDim>>>(...);
就会创建成千上万个线程，这些线程去执行 kernel 里的代码。作用：定义所有线程要干的事。

你写了一个 kernel（任务说明书）。启动时，会生成很多线程。GPU 硬件的调度器把这些线程分配给 CUDA cores 执行。

2.1.5 原子加法

在现代计算机中，很多操作都是多线程或并行执行的，例如 CPU 多核或者 GPU 上成千上万的线程同时运行。假设多个线程同时对同一个内存地址执行加法操作：

变量 x = 5
线程 A: x = x + 2
线程 B: x = x + 3

如果操作不是原子的，可能会出现竞态条件（race condition）：

线程 A 读 x（5）
线程 B 读 x（5）
线程 A 写 x = 7
线程 B 写 x = 8

最终结果可能是 8，而不是我们期望的 10。这是因为加法操作实际上包含多个步骤：读、加、写。如果被其他线程打断，就会出错。
原子加法（Atomic Addition）就是一种操作，它保证加法过程对其他线程不可中断，整个操作要么完全完成，要么完全不做，永远不会被其他线程“打断”。
用伪代码表示：

atomic_add(x, value):
    x = x + value   // 这个操作是原子的

无论多少线程同时执行：

atomic_add(x, 2)
atomic_add(x, 3)
atomic_add(x, 1)

最终 x 的值一定是正确累加后的结果（比如原来是 5，最终一定是 11）。

原子加法操作可以确保每个核心的计算结果都能最终反映在总和中。但是，它并不能保证这些结果的累加顺序。累加顺序完全取决于哪个核心先完成计算，这是一种不确定性行为。

2.2 存储

2.2.1 DRAM 和 SRAM

静态随机存取存储器（Static Random-Access Memory，SRAM是随机存取存储器的一种。所谓的“静态”，是指这种存储器只要保持通电，里面储存的数据就可以恒常保持。
每个 bit 用电容 + 晶体管存储，电容会漏电。

动态随机存取存储器（Dynamic Random Access Memory，DRAM里面所储存的数据就需要周期性地更新。 SRAM有它的缺点，即它的集成度较低，功耗较DRAM大，相同容量的DRAM内存可以设计为较小的体积，但是SRAM却需要很大的体积。同样面积的硅片可以做出更大容量的DRAM，因此SRAM显得更贵。
每个 bit 用触发器 (6个晶体管) 存储，数据只要有电就能保持，不需要刷新。通常是 1T1C的结果，也就是一个晶体管一个电容的组合。

2.2.2 SDRAM

同步动态随机存储器（Synchronous Dynamic RAM），采用3.3V工作电压，内存数据位宽64位，SDRAM与CPU通过一个相同的时钟频率锁在一起，使两者以相同的速度同步工作。它和 CPU 时钟同步工作。现在常说的 DDR、DDR2、DDR3、DDR4、DDR5 内存都是 SDRAM 的家族成员。

2.2.3 DDR DDR2 DDR3 DDR4

双倍速率同步动态随机存储器（Double data Rate SDRAM，DDR SDRAM，简称DDR），采用2.5V工作电压，内存数据位宽64位，一个时钟脉冲传输两次数据，分别在时钟脉冲的上升沿和下降沿各传输一次数据，因此称为双倍速率的SDRAM。
拥有两倍于上一代DDR内存的预读取能力，从2，3，4，5开始编号。

2.2.3 FLASH

FLASH闪存是属于内存器件的一种。闪存则是一种非易失性（ Non-Volatile ）内存，在没有电流供应的条件下也能够长久地保持数据，其存储特性相当于硬盘，这项特性正是闪存得以成为各类便携型数字设备的存储介质的基础。Flash 属于存储器，更接近硬盘/ROM 的角色，是长期数据保存。SSD（固态硬盘）里面的核心存储介质就是 Flash，一般是 NAND Flash（因为容量大、成本低、擦写速度合适）

2.2.3 GDDR5

显卡用的是GDDR5，G就是graphic也就是图形的意思显存是为图形处理器特殊定制和打造的，另外显存与图形处理器之间使用的是固化连接，自然可以提供比插接更高的频率，另外还可以根据不同的GPU来定制合适的显存另外GDDR5已经是DDR3时代的产物，随着内存进入DDR4时代相信下一代显存也快要来了。

引用

引用自：https://zhuanlan.zhihu.com/p/604571345
引用自：https://blogs.novita.ai/zh-CN/what-are-cuda-cores-a-deep-dive-into-gpu-parallel-processing/
引用自：https://blog.csdn.net/asasasaababab/article/details/80447254
引用自：https://www.akshatblog.com/graphics-card-components-explained-in-detail/
引用自：https://blog.csdn.net/kkk584520/article/details/53814067
引用自：https://zhuanlan.zhihu.com/p/375053359

3. 后记

To be continued.......

posted @ 2025-09-18 16:47 泪水下的笑靥阅读(27) 评论(0) 收藏举报

刷新页面返回顶部

myleaf