CUDA Programming

https://github.com/ForceInjection/AI-fundermentals

Getting Started

开发环境与工具链：CUDA GPU、驱动、CUDA Toolkit、C 编译器等。

在windows下需要安装 nvidiasmi, nvcc, visual stuido 2022, 检测：

nvidia-smi
nvcc --version

然后新建一个CUDA template项目，运行一下，可以跑通就可以

Introduction to CUDA C

第一个 CUDA 程序、kernel 调用、传参、查询设备、使用 device properties。

不能用 format 因为 cuda 对 std的支持不好

#include <cstdio>
#include <cuda_runtime.h>
#include <source_location>
#include <iostream>

inline void cuda_check(cudaError_t e, std::source_location loc = std::source_location::current()) {
    if (e != cudaSuccess) {
        //std::cerr << std::format("CUDA error {}:{}: {}\n", loc.file_name(), loc.line(), cudaGetErrorString(e));
        //std::exit(1);
        std::fprintf(stderr, "CUDA error %s:%u: %s\n",
            loc.file_name(), loc.line(), cudaGetErrorString(e));
        std::exit(1);
    }
}

int main() {
    int n = 0;
    cuda_check(cudaGetDeviceCount(&n));
    printf("CUDA devices: %d\n", n);

    for (int i = 0; i < n; i++) {
        cudaDeviceProp p{};
        cuda_check(cudaGetDeviceProperties(&p, i));
        printf("\n== Device %d ==\n", i);
        printf("Name: %s\n", p.name);
        printf("Compute capability: %d.%d\n", p.major, p.minor);
        printf("SM count: %d\n", p.multiProcessorCount);
        printf("Warp size: %d\n", p.warpSize);
        printf("Max threads per block: %d\n", p.maxThreadsPerBlock);
        printf("Max threads dim: (%d, %d, %d)\n",
            p.maxThreadsDim[0], p.maxThreadsDim[1], p.maxThreadsDim[2]);
        printf("Max grid size: (%d, %d, %d)\n",
            p.maxGridSize[0], p.maxGridSize[1], p.maxGridSize[2]);
        printf("Shared mem per block: %zu bytes\n", p.sharedMemPerBlock);
        printf("Total global mem: %.2f GB\n", (double)p.totalGlobalMem / (1024.0 * 1024.0 * 1024.0));
        printf("Mem clock rate: %d kHz\n", p.memoryClockRate);
        printf("Mem bus width: %d bits\n", p.memoryBusWidth);
    }
    return 0;
}

Parallel Programming in CUDA C

CUDA 并行编程模型：向量求和等示例，引出 grid/block/thread。

#include <cstdio>
#include <device_launch_parameters.h>
#include <cuda_runtime.h>
#include <source_location>
#include <iostream>
#include <memory>

inline void cuda_check(cudaError_t e, std::source_location loc = std::source_location::current()) {
    if (e != cudaSuccess) {
        //std::cerr << std::format("CUDA error {}:{}: {}\n", loc.file_name(), loc.line(), cudaGetErrorString(e));
        //std::exit(1);
        std::fprintf(stderr, "CUDA error %s:%u: %s\n",
            loc.file_name(), loc.line(), cudaGetErrorString(e));
        std::exit(1);
    }
}

__global__ void add1(int* a, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        a[i] += 1;
    }
}

int main() {
    constexpr const int n = 1 << 20;
    constexpr const int array_size = n * sizeof(int);
    auto h = std::make_unique<int[]>(n);
    
    for (int i = 0;i < n; ++i) {
        h[i] = i;
    }

    int* d = nullptr;
    
    cuda_check(cudaMalloc(&d, array_size));
    cuda_check(cudaMemcpy(d, h.get(), array_size, cudaMemcpyHostToDevice));

    int block = 256;
    int grid = (n + block - 1) / block;

    add1 <<< grid, block >>> (d, n);
    cuda_check(cudaGetLastError());
    cuda_check(cudaDeviceSynchronize());

    cuda_check(cudaMemcpy(h.get(), d, array_size, cudaMemcpyDeviceToHost));
    cuda_check(cudaFree(d));

    for (int i = 0;i < 10; ++i) {
        std::cout << h[i] << std::endl;
    }
    return 0;
}

做一件最简单的事： 在 CPU 上准备一个数组 → 拷到 GPU → GPU 每个元素 +1 → 再拷回 CPU → 打印前 10 个验证。

它覆盖了 CUDA 程序的四件套：

1. Host（CPU）内存准备
2. Device（GPU）内存分配
3. Host↔︎Device 数据传输
4. Kernel 启动与同步、再取回结果

__global__ 的意思

• 这个函数在 GPU 上执行
• 由 CPU 端用 <<< >>> 语法启动

索引计算

int i = blockIdx.x * blockDim.x + threadIdx.x;

第一个参数 grid：告诉 GPU 启动多少个块（Block）。

第二个参数 block：告诉 GPU 每个块里有多少个线程（Thread）。

含义：

• threadIdx.x：你这个 block 里的第几个线程（0..blockDim.x-1）
• blockIdx.x：这是第几个 block（0..gridDim.x-1）
• blockDim.x：每个 block 有多少线程

cudaGetLastError()：检查 kernel 启动配置是否有错 例如 block 太大、参数非法等（这类错误在 launch 这一刻就能发现）

cudaDeviceSynchronize()：等待 GPU 把 kernel 跑完 如果 kernel 里出现 非法访存 等运行时错误，必须同步后才能暴露出来

接下来可以改成vector相加的代码

记时

// 新建事件 
cudaEvent_t start, stop;
 cuda_check(cudaEventCreate(&start));
 cuda_check(cudaEventCreate(&stop));
// 记录
 cuda_check(cudaEventRecord(start));
 vec_add <<< grid, block >>> (da, db, n);
 cuda_check(cudaGetLastError());
 cuda_check(cudaEventRecord(stop));
 cuda_check(cudaDeviceSynchronize());
//计算时间
 float ms = 0.0f;
 cuda_check(cudaEventElapsedTime(&ms, start, stop));
//销毁
 cuda_check(cudaEventDestroy(start));
 cuda_check(cudaEventDestroy(stop));

第 5 章：Thread Cooperation

线程协作：拆分并行块、shared memory、同步；含 dot product、位图等例子。 pearson.de+1

第 6 章：Constant Memory and Events

常量内存（以 ray tracing 示例展开）、以及用 events 做性能测量。 pearson.de+1

第 7 章：Texture Memory

纹理内存概念与使用，用“热传导模拟”例子讲 texture（含 2D texture）。 pearson.de+1

第 8 章：Graphics Interoperability

与图形 API 互操作（OpenGL/DirectX 等）：图形互操作下的 ripple、热传导等。 pearson.de+1

第 9 章：Atomics

Compute Capability、原子操作概览，用直方图（histogram）讲 atomics。 pearson.de+1

第 10 章：Streams

Pinned（page-locked）主机内存、CUDA streams、单/多 stream 使用与调度。 pearson.de+1

第 11 章：CUDA C on Multiple GPUs

Zero-copy host memory、多 GPU 使用、portable pinned memory。 pearson.de+1

第 12 章：The Final Countdown

CUDA 工具（如 toolkit、库、SDK、调试/Profiler 等）与进一步阅读/代码资源。 pearson.de+1

Appendix A：Advanced Atomics

更高级的原子用法：dot product revisited、用原子锁实现 hash table 等。 pearson.de+1

PMPP

第一章

GPU 为什么适合做“海量并行”

CPU：低延迟（latency）优先 典型策略：复杂控制逻辑、深缓存层次、乱序执行等，尽量让“单线程尽快跑完”。 stevengong.co+1

GPU：高吞吐（throughput）优先 用大量相对简单的执行核心 + 海量线程，把硬件资源更多用于算术吞吐和带宽利用，让“单位时间完成更多总工作”。

现代 GPU 的关键组织方式

多 SM（Streaming Multiprocessor）/计算单元：GPU 由多个 SM 组成，每个 SM 里有多条执行通道/更小的核心（SP/ALU 等）共享部分控制与指令缓存等资源。 faculty.kfupm.edu.sa

以大量线程隐藏延迟：当一批线程遇到访存等待，硬件切换去跑另一批就绪线程（用并发度“盖住”内存延迟）。这是 GPU 吞吐思路的核心之一。 faculty.kfupm.edu.sa

全局显存（GDDR/Global Memory）：GPU 通常配套高带宽显存（书里常称 global memory），适合大吞吐数据流，但延迟仍然高，因此更需要并发与良好的访问模式

第三章 CUDA 编程

CUDA 程序把工作分成两类：

• Host（CPU）代码：负责控制流程、分配/回收设备内存、搬运数据、发起 kernel。
• Device（GPU）代码：用 kernel 在 GPU 上并行执行大量线程完成计算。

把某个计算搬到 GPU 上”通常遵循固定套路： (1) Host→Device 拷贝输入 → (2) 启动 kernel 计算 → (3) Device→Host 拷贝结果。

CUDA C 的函数类型与关键字

CUDA 通过在 C 函数声明上加限定词，区分函数运行位置/调用关系：Programming_Massively_Parallel_…

• __global__：kernel 函数（在 device 上跑、由 host 发起调用/launch）
• __device__：设备函数（在 device 上跑、只能由 device 调用）
• __host__：主机函数（在 host 上跑；不写默认就是 host）

Kernel launch 语法

执行配置最核心就是两件事：

• grid 里有多少个 block
• 每个 block 有多少个 thread（以及是否是 2D/3D 的组织）

第四章 CUDA线程

第 4 章叫 CUDA Threads，核心是在回答三件事： 1）线程在 CUDA 里怎么“编号/组织”；

2）怎么用这些编号把数据分给线程（让规模可扩展）；

3）线程在硬件上怎么被分配与调度（为什么 block 大小会影响性能）。