CUDA编程

前言

CUDA快速入手系列视频

谭升博客

我的代码

内容大概是对视频、博客的一个个人学习总结以及部分高质量内容的摘抄汇总，以及补充自己的一些个人理解。

入门下cuda编程，GPU感觉是一个必须要会的东西

• CPU适合执行复杂的逻辑，比如多分支，其核心比较重（复杂）
• GPU适合执行简单的逻辑，大量的数据计算，其吞吐量更高，但是核心比较轻（结构简单）

GPU主要负责并行计算

两者的架构区分主要在于：

• 左图：一个四核CPU一般有四个ALU，ALU是完成逻辑计算的核心，也是我们平时说四核八核的核，控制单元，缓存也在片上，DRAM是内存，一般不在片上，CPU通过总线访问内存。
• 右图：GPU，绿色小方块是ALU，我们注意红色框内的部分SM，这一组ALU公用一个Control单元和Cache，这个部分相当于一个完整的多核CPU，但是不同的是ALU多了，control部分变小，可见计算能力提升了，控制能力减弱了，所以对于控制（逻辑）复杂的程序，一个GPU的SM是没办法和CPU比较的，但是对了逻辑简单，数据量大的任务，GPU更搞笑，并且，注意，一个GPU有好多个SM，而且越来越多。

CPU和GPU之间通过PCIe总线连接，用于传递指令和数据，也是最主要的性能瓶颈之一。

两者的性能不可比较，而是相互配合

CPU和GPU线程的区别：

1. CPU线程是重量级实体，操作系统交替执行线程，线程上下文切换花销很大
2. GPU线程是轻量级的，GPU应用一般包含成千上万的线程，多数在排队状态，线程之间切换基本没有开销。
3. CPU的核被设计用来尽可能减少一个或两个线程运行时间的延迟，而GPU核则是大量线程，最大幅度提高吞吐量

CUDA平台不是单单指软件或者硬件，而是建立在Nvidia GPU上的一整套平台，并扩展出多语言支持

对于API也有两种不同的层次，一种相对交高层，一种相对底层。

• CUDA驱动API
• CUDA运行时API

驱动API是低级的API，使用相对困难，运行时API是高级API使用简单，其实现基于驱动API。
这两种API是互斥的，也就是你只能用一个，两者之间的函数不可以混合调用，只能用其中的一个库。

一个CUDA应用通常可以分解为两部分，

• CPU 主机端代码
• GPU 设备端代码

后面的内容大致看了下，就是说CUDA与高级语言的编译是相对分离的，使用到cuda库的代码会被自动分离，一般我们只会写核函数，在高级语言中去调用相关CUDA函数，分配特定任务

grid/block/thread GPU的魔方结构

CUDA中，grid和block是逻辑概念，用于描述CUDA程序中线程的组织和调度方式，而不是物理硬件的实际结构。

Grid由多个Blocks组成，Block又由多个Threads组成，类似套娃

grid和block的逻辑定位方式都是三维的，恰好类似于魔方

在编程中，维度可以由自己来定义

通过索引定位，即可确定grid中的block，block中的thread

01 hello world

首先是安装CUDA的相关支持

sudo apt install nvidia-utils-550 libnvidia-encode-550

接着创建helloworld.cu（后缀不是cpp）

#include <stdio.h>
#include <iostream>
#include <cuda_runtime.h>
#include <cstdlib>

// CUDA 核函数
__global__ void hello_world() {
    int threadId = threadIdx.x;
    printf("GPU 线程 %d: Hello World!\n", threadId);
}

// 错误检查宏
#define CUDA_CHECK_ERROR(err) \
    if (err != cudaSuccess) { \
        fprintf(stderr, "CUDA 错误: %s\n", cudaGetErrorString(err)); \
        exit(EXIT_FAILURE); \
    }

int main(int argc, char **argv) {
    std::cout<<"CPU: Hello World!\n";
    
    // 启动核函数，使用 1 个块，10 个线程
    hello_world<<<1, 10>>>();
    
    // 同步以确保 GPU 输出完成
    CUDA_CHECK_ERROR(cudaDeviceSynchronize());
    
    // 重置设备
    CUDA_CHECK_ERROR(cudaDeviceReset());
    
    return 0;
}

注意：在 CUDA 编程中，__global__ 核函数（kernel function）运行在 GPU 设备侧（device code），而 std::cout 是 C++ 标准库的 主机侧（host code）输出工具，主要设计用于 CPU 执行的环境。这会导致编译错误，通常是类似这样的报错：

error: calling a host function("std::basic_ostream<char, std::char_traits<char> >::operator<<") from a __global__ function("hello_world") is not allowed

运行结果如下:

一点小tips：

1. vscode使用cuda-gdp调试的时候需要安装插件：Nsight Visual Studio Code Edition（ctrl+p 输入：ext install NVIDIA.nsight-vscode-edition下载安装即可）

2. 我的配置文件如下：

   launch.json

   {
       "version": "0.2.0",
       "configurations": [
           {
               "name": "CUDA C++: Launch",  // 标准名称，便于识别
               "type": "cuda-gdb",  // 保持不变，这是 Nsight 提供的类型
               "request": "launch",
               "program": "${fileDirname}/${fileBasenameNoExtension}",  // 动态指向可执行文件
               "args": [],  // 命令行参数，如果需要加如 ["input.txt"]
               "cwd": "${fileDirname}",  // 工作目录
               "preLaunchTask": "build CUDA",  // 先编译
               "miDebuggerPath": "/usr/bin/cuda-gdb",  // 修正：用 debuggerPath 而非 miDebuggerPath,
           }
       ]
   }

   tasks.json

   {
       "version": "2.0.0",
       "tasks": [
           {
               "label": "build CUDA",
               "type": "shell",
               "command": "nvcc",
               "args": [
                   "-g", "-G",  // 加调试符号（-g: CPU, -G: GPU）
                   "-o", 
                   "${fileDirname}/${fileBasenameNoExtension}", 
                   "${file}"
               ],
               "group": "build",
               "presentation": { 
                   "echo": true, 
                   "reveal": "always", 
                   "focus": false 
               },
               "problemMatcher": ["$gcc"]
           },
           {
               "label": "run CUDA",
               "type": "shell",
               "command": "${fileDirname}/${fileBasenameNoExtension}",
               "args": [],  // 无额外参数，可加如 ["--debug"] 如果需要
               "dependsOn": "build CUDA",  // 先运行编译任务
               "group": "build",
               "presentation": { 
                   "echo": true, 
                   "reveal": "always", 
                   "focus": false 
               },
               "problemMatcher": []
           }
       ]
   }

02 CUDA的内存结构

共享内存和全局内存是比较重要的一个点，然后除此之外，各个gpu也都会有自己的寄存器和自己的一些内存。

在使用过程中，主要要注意的是GPU（device）到CPU（host）的一些内存的传输。

/*
 * CUDA 示例：数组求和（Array Sum）
 * 
 * 这个示例演示了 CUDA 的基本使用：
 * 1. 在主机（CPU）上分配和初始化两个浮点数组 a 和 b。
 * 2. 将数组复制到设备（GPU）内存。
 * 3. 启动一个 CUDA 内核（sumArraysGPU），每个线程负责计算一个元素的和（a[i] + b[i]）。
 *    - 使用一个线程块（block size = 32），一个网格（grid size = 1），因为数组大小为 32。
 * 4. 将结果从设备复制回主机。
 * 5. 在主机上使用 CPU 版本的函数（sumArrays）计算相同的和。
 * 6. 比较 CPU 和 GPU 结果的正确性（使用 checkResult 函数）。
 * 
 * 这个示例是 CUDA 入门级（Freshman）教程的一部分，目的是验证 GPU 并行计算的基本正确性。
 * 注意：数组大小固定为 32，以匹配一个线程块的大小。
 */

#include <cuda_runtime.h>
#include <iostream>
#include <cstring>  // 用于 memset
#include <cmath>    // 用于 fabs

// 定义 CHECK 宏：检查 CUDA API 调用是否成功，如果失败则打印错误并退出
#define CHECK(call) { \
    cudaError_t err = call; \
    if (err != cudaSuccess) { \
        std::cerr << "CUDA Error: " << cudaGetErrorString(err) << " at line " << __LINE__ << std::endl; \
        exit(1); \
    } \
}

// initialData 函数：初始化数组 data 的前 n 个元素为 0.0f, 1.0f, 2.0f, ..., (n-1).0f
// 这个函数用于生成测试数据，便于验证结果（预期和为 a[i] + b[i] = 2*i）
void initialData(float *data, int n) {
    for (int i = 0; i < n; i++) {
        data[i] = static_cast<float>(i);
    }
}

// sumArrays 函数：CPU 版本的数组求和，使用循环（每 4 个元素展开以模拟向量化）
void sumArrays(float *a, float *b, float *res, const int size) {
    for (int i = 0; i < size; i += 4) {
        res[i] = a[i] + b[i];
        if (i + 1 < size) res[i + 1] = a[i + 1] + b[i + 1];
        if (i + 2 < size) res[i + 2] = a[i + 2] + b[i + 2];
        if (i + 3 < size) res[i + 3] = a[i + 3] + b[i + 3];
    }
}

// checkResult 函数：比较两个数组 ref（CPU 结果）和 gpu（GPU 结果）的总和差异
// 如果相对误差小于 1e-5，则测试通过并打印 "PASSED"，否则 "FAILED"
void checkResult(float *ref, float *gpu, int n) {
    double sum_ref = 0.0, sum_gpu = 0.0;
    for (int i = 0; i < n; i++) {
        sum_ref += ref[i];
        sum_gpu += gpu[i];
    }
    double eps = 1e-5;
    double diff = fabs(sum_ref - sum_gpu) / sum_ref;
    std::cout << "Test " << (diff < eps ? "PASSED" : "FAILED") << std::endl;
}

// CUDA 内核函数：GPU 版本的数组求和
// 每个线程（threadIdx.x）计算 res[i] = a[i] + b[i]
__global__ void sumArraysGPU(float *a, float *b, float *res) {
    int i = threadIdx.x;  // 获取当前线程的索引（假设块大小等于数组大小）
    res[i] = a[i] + b[i];
}

int main(int argc, char **argv) {
    int dev = 0;  // 指定设备 ID（默认 GPU 0）
    CHECK(cudaSetDevice(dev));

    int nElem = 32;  // 数组大小：32（适合一个线程块）
    std::cout << "Vector size: " << nElem << std::endl;
    int nByte = sizeof(float) * nElem;  // 字节数

    // 主机内存分配
    float *a_h = new float[nElem];
    float *b_h = new float[nElem];
    float *res_h = new float[nElem];  // CPU 结果
    float *res_from_gpu_h = new float[nElem];  // GPU 结果
    memset(res_h, 0, nByte);  // 清零
    memset(res_from_gpu_h, 0, nByte);

    // 设备内存分配
    float *a_d, *b_d, *res_d;
    CHECK(cudaMalloc((void**)&a_d, nByte));
    CHECK(cudaMalloc((void**)&b_d, nByte));
    CHECK(cudaMalloc((void**)&res_d, nByte));

    // 初始化主机数据
    initialData(a_h, nElem);
    initialData(b_h, nElem);

    // 主机到设备复制
    CHECK(cudaMemcpy(a_d, a_h, nByte, cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(b_d, b_h, nByte, cudaMemcpyHostToDevice));

    // 执行配置：一个块（32 线程），一个网格（32/32=1）
    dim3 block(nElem);
    dim3 grid(nElem / block.x);
    sumArraysGPU<<<grid, block>>>(a_d, b_d, res_d);
    std::cout << "Execution configuration <<<" << grid.x << ", " << block.x << ">>>" << std::endl;

    // 设备到主机复制
    CHECK(cudaMemcpy(res_from_gpu_h, res_d, nByte, cudaMemcpyDeviceToHost));

    // CPU 计算
    sumArrays(a_h, b_h, res_h, nElem);

    // 验证结果
    checkResult(res_h, res_from_gpu_h, nElem);

    // 清理内存
    cudaFree(a_d);
    cudaFree(b_d);
    cudaFree(res_d);
    delete[] a_h;
    delete[] b_h;
    delete[] res_h;
    delete[] res_from_gpu_h;

    return 0;
}

其中的关键函数：

cudaError_t cudaMemcpy(void * dst,const void * src,size_t count,
  cudaMemcpyKind kind)

这个函数是内存拷贝过程，可以完成以下几种过程（cudaMemcpyKind kind）

• cudaMemcpyHostToHost
• cudaMemcpyHostToDevice
• cudaMemcpyDeviceToHost
• cudaMemcpyDeviceToDevice

使用示例如cudaMemcpy(res_from_gpu_h, res_d, nByte, cudaMemcpyDeviceToHost)

然后针对于索引定位，在上文中已经有过比较明确的说明了。

一般来说线程索引+块索引×块维度就可以确定目前这个线程处理的任务是第几个，如果网格不止一个，也可以以此类推。

我们可以使用dim3类型的grid维度和block维度配置内核，也可以使用int类型的变量，或者常量直接初始化：

kernel_name<<<4,8>>>(argument list);

想要主机等待设备端执行可以用下面这个指令：

cudaError_t cudaDeviceSynchronize(void);

这是一个显示的方法，对应的也有隐式方法，隐式方法就是不明确说明主机要等待设备端，而是设备端不执行完，主机没办法进行，比如内存拷贝函数：

cudaError_t cudaMemcpy(void* dst,const void * src,
  size_t count,cudaMemcpyKind kind);

这个函数上文已经介绍过了，当核函数启动后的下一条指令就是从设备复制数据回主机端，那么主机端必须要等待设备端计算完成。

所有CUDA核函数的启动都是异步的，这点与C语言是完全不同的

Kernel核函数编写有以下限制

1. 只能访问设备内存
2. 必须有void返回类型
3. 不支持可变数量的参数
4. 不支持静态变量
5. 显示异步行为

看了这些之后其实动手写一遍之前的代码基本上就能理清一下逻辑链了。

进一步来看 也可以有更好的内存组织形式

设备内存或者主机内存都是线性存在的，比如一个二维矩阵 (8×6)，存储在内存中是这样的：

我们要做管理的就是：

• 线程和块索引（来计算线程的全局索引）
• 矩阵中给定点的坐标（ix,iy）
• (ix,iy)对应的线性内存的位置

线性位置的计算方法是：

idx = ix + iy × nx

（个人感觉效率提升没有特别大帮助 都得重新分配和cudaMemcpy进GPU中）

03 给核函数计时

给cpu计时是比较常用的分析性能的工具和计时方法，比方说给cpu的任务调用采样计数的火山图

据博客所说，并行程序中使用clock函数计时是严重问题的：

clock() 返回的是整个进程的总 CPU 消耗时间，包括所有线程的贡献。这与实际的执行时长无关，而是反映了 CPU 被进程“占用”的总量。

在并行执行时，多个线程可以同时在不同 CPU 核心上运行，导致 CPU 时间并行累加。例如，如果有两个线程各自消耗 1 秒 CPU 时间，但它们并行执行（总墙钟时间仅 1 秒），clock() 会报告约 2 秒的“经过时间”，从而严重高估实际执行时长。这种膨胀效应会随着线程数增加而加剧（N 线程下可能高估 N 倍），使得计时结果不可靠，尤其在基准测试或性能评估中。

clock_t start, finish;
start = clock();
// 要测试的部分
finish = clock();
duration = (double)(finish - start) / CLOCKS_PER_SEC;

所以说似乎是使用这个函数会多一点：

#include <sys/time.h>
double cpuSecond()
{
  struct timeval tp;
  gettimeofday(&tp,NULL);
  return((double)tp.tv_sec+(double)tp.tv_usec*1e-6);
}

而后还有一个比较常用的工具nvprof：

nvprof [nvprof_args] <application>[application_args]

不过在8.0架构以后的GPU不再适用，可以转而使用ncu或者nsys，后续有需求了再单独看看。

04 查询GPU信息

常用的可能还是

nvidia-smi

下面这些nvidia-smi -q -i 0 的参数可以提取我们要的信息

• MEMORY
• UTILIZATION
• ECC
• TEMPERATURE
• POWER
• CLOCK
• COMPUTE
• PIDS
• PERFORMANCE
• SUPPORTED_CLOCKS
• PAGE_RETIREMENT
• ACCOUNTING

比如我们想得到内存信息：

nvidia-smi -q -i 0 -d MEMORY

多设备时，我们只要把上面的0改成对应的设备号就好了

然后是一些api：

/**
 * CUDA 设备查询示例
 * 
 * 此程序查询并打印第一个 CUDA 兼容设备的属性。
 * 基于 NVIDIA 的 deviceQuery 示例，使用 C++ 编写，并更新为最新的 CUDA 12.x API。
 * 
 * 使用的关键 API（来自 CUDA 运行时 API，CUDA 12.x 中的最新版本）：
 * - cudaGetDeviceCount()：检索 CUDA 兼容设备的数量。
 * - cudaSetDevice()：设置当前 CUDA 设备（此处为设备 0）。
 * - cudaGetDeviceProperties()：用设备属性填充 cudaDeviceProp 结构体。
 *   - cudaDeviceProp：存储设备属性的结构体（例如，主要/次要计算能力、
 *     内存大小、线程限制）。详见 NVIDIA 文档：https://docs.nvidia.com/cuda/cuda-runtime-api/structcudaDeviceProp.html
 * - cudaDriverGetVersion() / cudaRuntimeGetVersion()：获取驱动程序和运行时版本。
 * 
 * 编译（需要 CUDA Toolkit >= 12.x）：
 *   nvcc -std=c++17 -o deviceQuery deviceQuery.cu
 * 
 * 使用方法：
 *   ./deviceQuery [可选参数如果扩展]
 * 
 * 输出：将设备信息打印到标准输出。优雅处理错误。
 * 
 * 注意：
 * - 专注于原始 C 代码中的核心属性；可扩展到所有 cudaDeviceProp 成员。
 * - 错误处理：检查 cudaError_t 返回值。
 * - 精度：使用 <iomanip> 进行格式化的浮点数输出。
 */

#include <cuda_runtime.h>
#include <iostream>
#include <iomanip>
#include <cmath>
#include <cstdlib>  // 用于 EXIT_SUCCESS/FAILURE

int main(int argc, char** argv) {
    std::cout << argv[0] << " Starting ..." << std::endl;

    // 步骤 1: 查询 CUDA 设备数量
    int deviceCount = 0;
    cudaError_t error_id = cudaGetDeviceCount(&deviceCount);  // 最新 API：在 12.x 中无变化
    if (error_id != cudaSuccess) {
        std::cout << "cudaGetDeviceCount returned " << static_cast<int>(error_id) << std::endl
                  << " -> " << cudaGetErrorString(error_id) << std::endl;
        std::cout << "Result = FAIL" << std::endl;
        return EXIT_FAILURE;
    }

    if (deviceCount == 0) {
        std::cout << "There are no available device(s) that support CUDA" << std::endl;
        return EXIT_SUCCESS;  // 如果没有设备，不是失败
    } else {
        std::cout << "Detected " << deviceCount << " CUDA Capable device(s)" << std::endl;
    }

    // 步骤 2: 选择第一个设备（dev=0）并查询属性
    int dev = 0;
    cudaSetDevice(dev);  // 设置活动设备；如果设备无效则出错
    cudaDeviceProp deviceProp;  // 属性结构体（在 CUDA 12.x 中扩展，例如 memoryPoolsSupported）
    error_id = cudaGetDeviceProperties(&deviceProp, dev);
    if (error_id != cudaSuccess) {
        std::cout << "cudaGetDeviceProperties returned " << static_cast<int>(error_id) << std::endl
                  << " -> " << cudaGetErrorString(error_id) << std::endl;
        return EXIT_FAILURE;
    }

    // 步骤 3: 打印基本设备信息
    std::cout << "Device " << dev << ": \"" << deviceProp.name << "\"" << std::endl;  // ASCII 设备名称（最大 256 字符）

    // 驱动程序和运行时版本（最新 API：支持高达 CUDA 12.8+）
    int driverVersion = 0, runtimeVersion = 0;
    cudaDriverGetVersion(&driverVersion);
    cudaRuntimeGetVersion(&runtimeVersion);
    std::cout << "  CUDA Driver Version / Runtime Version         "
              << driverVersion / 1000 << "." << (driverVersion % 100) / 10 << "  /  "
              << runtimeVersion / 1000 << "." << (runtimeVersion % 100) / 10 << std::endl;

    // 计算能力（major/minor；例如，Ampere 的 8.6）
    std::cout << "  CUDA Capability Major/Minor version number:   "
              << deviceProp.major << "." << deviceProp.minor << std::endl;

    // 全局内存（totalGlobalMem：size_t，以字节为单位）
    std::cout << "  Total amount of global memory:                "
              << std::fixed << std::setprecision(2)
              << static_cast<float>(deviceProp.totalGlobalMem) / std::pow(1024.0, 3) << " MBytes ("
              << deviceProp.totalGlobalMem << " bytes)" << std::endl;

    // 时钟速率（clockRate：int，以 kHz 为单位）
    std::cout << "  GPU Clock rate:                               "
              << std::fixed << std::setprecision(0) << deviceProp.clockRate * 1e-3f << " MHz ("
              << std::setprecision(2) << deviceProp.clockRate * 1e-6f << " GHz)" << std::endl;

    // 内存总线宽度（memoryBusWidth：int，位）
    std::cout << "  Memory Bus width:                             " << deviceProp.memoryBusWidth << "-bits" << std::endl;

    // L2 缓存（l2CacheSize：int，字节；如果不支持则为 0）
    if (deviceProp.l2CacheSize) {
        std::cout << "  L2 Cache Size:                                " << deviceProp.l2CacheSize << " bytes" << std::endl;
    }

    // 纹理维度（各种 maxTexture* 成员：int 数组）
    std::cout << "  Max Texture Dimension Size (x,y,z)            1D=(" << deviceProp.maxTexture1D
              << "), 2D=(" << deviceProp.maxTexture2D[0] << "," << deviceProp.maxTexture2D[1]
              << "), 3D=(" << deviceProp.maxTexture3D[0] << "," << deviceProp.maxTexture3D[1]
              << "," << deviceProp.maxTexture3D[2] << ")" << std::endl;
    std::cout << "  Max Layered Texture Size (dim) x layers       1D=(" << deviceProp.maxTexture1DLayered[0]
              << ") x " << deviceProp.maxTexture1DLayered[1]
              << ", 2D=(" << deviceProp.maxTexture2DLayered[0] << "," << deviceProp.maxTexture2DLayered[1]
              << ") x " << deviceProp.maxTexture2DLayered[2] << std::endl;

    // 常量内存和共享内存（totalConstMem/sharedMemPerBlock：size_t，字节）
    std::cout << "  Total amount of constant memory               " << deviceProp.totalConstMem << " bytes" << std::endl;
    std::cout << "  Total amount of shared memory per block:      " << deviceProp.sharedMemPerBlock << " bytes" << std::endl;

    // 寄存器（regsPerBlock：int）
    std::cout << "  Total number of registers available per block: " << deviceProp.regsPerBlock << std::endl;

    // 线程束大小（warpSize：int，通常为 32）
    std::cout << "  Warp size:                                    " << deviceProp.warpSize << std::endl;

    // 线程限制（maxThreadsPerMultiProcessor/maxThreadsPerBlock：int）
    std::cout << "  Maximum number of threads per multiprocessor: " << deviceProp.maxThreadsPerMultiProcessor << std::endl;
    std::cout << "  Maximum number of threads per block:          " << deviceProp.maxThreadsPerBlock << std::endl;

    // 块维度（maxThreadsDim[3]：int 数组，x/y/z）
    std::cout << "  Maximum size of each dimension of a block:    " << deviceProp.maxThreadsDim[0] << " x "
              << deviceProp.maxThreadsDim[1] << " x " << deviceProp.maxThreadsDim[2] << std::endl;

    // 网格维度（maxGridSize[3]：size_t 数组，x/y/z；在 12.x 中为 size_t）
    std::cout << "  Maximum size of each dimension of a grid:     " << deviceProp.maxGridSize[0] << " x "
              << deviceProp.maxGridSize[1] << " x " << deviceProp.maxGridSize[2] << std::endl;

    // 内存间距（memPitch：size_t，最大字节）
    std::cout << "  Maximum memory pitch                          " << deviceProp.memPitch << " bytes" << std::endl;

    std::cout << "Result = PASS" << std::endl;  // 成功指示

    return EXIT_SUCCESS;
}

05 GPU架构

GPU架构是围绕一个流式多处理器（SM）的扩展阵列搭建的。通过复制这种结构来实现GPU的硬件并行。

上图包括关键组件：

• CUDA核心

• 共享内存/一级缓存

• 寄存器文件

• 加载/存储单元

• 特殊功能单元

• 线程束调度器

  SM

  GPU中每个SM都能支持数百个线程并发执行，每个GPU通常有多个SM，当一个核函数的网格被启动的时候，多个block会被同时分配给可用的SM上执行。

当一个block被分配给一个SM后，他就只能在这个SM上执行了，不可能重新分配到其他SM上了，多个线程块可以被分配到同一个SM上。

一个SM上面最多只有一个线程束，也就是32个线程在同时同步进行，每个线程执行同一条命令。

CUDA编程的组件与逻辑

下图从逻辑角度和硬件角度描述了CUDA编程模型对应的组件。

因为SM有限，虽然我们的编程模型层面看所有线程都是并行执行的，但是在微观上看，所有线程块也是分批次的在物理层面的机器上执行，线程块里不同的线程可能进度都不一样，但是同一个线程束内的线程拥有相同的进度。

Fermi 架构

Fermi架构逻辑图如上，具体数据如下:

1. 512个加速核心，CUDA核
2. 每个CUDA核心都有一个全流水线的整数算数逻辑单元ALU，和一个浮点数运算单元FPU
3. CUDA核被组织到16个SM上
4. 6个384-bits的GDDR5 的内存接口
5. 支持6G的全局机栽内存
6. GigaThread引擎，分配线程块到SM线程束调度器上
7. 768KB的二级缓存，被所有SM共享

而SM则包括下面这些资源：

• 执行单元（CUDA核）
• 调度线程束的调度器和调度单元
• 共享内存，寄存器文件和一级缓存

每个多处理器SM有16个加载/存储单元所以每个时钟周期内有16个线程（半个线程束）计算源地址和目的地址（16个cuda核心）

特殊功能单元SFU执行固有指令，如正弦，余弦，平方根和插值，SFU在每个时钟周期内的每个线程上执行一个固有指令。
每个SM有两个线程束调度器，和两个指令调度单元，当一个线程块被指定给一个SM时，线程块内的所有线程被分成线程束，两个线程束选择其中两个线程束，在用指令调度器存储两个线程束要执行的指令（就像例子中分水果的水果一样，这里有两个班，两个班的老师各自控制的自己的水果，老师就是指令调度器）
像第一张图上的显示一样，每16个CUDA核心为一个组，还有16个加载/存储单元或4个特殊功能单元。当某个线程块被分配到一个SM上的时候，会被分成多个线程束，线程束在SM上交替执行：

上面曾经说过，每个线程束在同一时间执行同一指令，同一个块内的线程束互相切换是没有时间消耗的。
Fermi上支持同时并发执行内核。并发执行内核允许执行一些小的内核程序来充分利用GPU，如图：

Profile可以帮助我们观察程序内部。

• 一个原生的内核应用一般不会产生最佳效果，也就是我们基本不能一下子就写出最好最快的内核，需要通过性能分析工具分析性能。找出性能瓶颈
• CUDA将SM中的计算资源在该SM中的多个常驻线程块之间进行分配，这种分配方式可能导致一些资源成为性能限制因素，性能分析工具可以帮我们找出来这些资源是如何被使用的
• CUDA提供了一个硬件架构的抽象。它能够让用户控制线程并发。性能分析工具可以检测和优化，并肩优化可视化

限制内核性能的主要包括但不限于以下因素

• 存储带宽
• 计算资源
• 指令和内存延迟

线程束和线程块

当一个线程块中有128个线程的时候，其分配到SM上执行时，会分成4个块：

warp0: thread  0,........thread31
warp1: thread 32,........thread63
warp2: thread 64,........thread95
warp3: thread 96,........thread127

当编号使用三维编号时，x位于最内层，y位于中层，z位于最外层，想象下c语言的数组，如果把上面这句话写成c语言，假设三维数组t保存了所有的线程，那么(threadIdx.x,threadIdx.y,threadIdx.z)表示为

t[z][y][x];

计算出三维对应的线性地址是：

tid = threadIdx.x + threadIdx.y × blockDim.x + threadIdx.z × blockDim.z × blockDim.y

（回忆一下魔方结构 z代表第z层 y代表那一层的第y行 x代表那一行的第x个）

线程束被执行的时候会被分配给相同的指令，处理各自私有的数据

每个线程都执行所有的if和else部分，当一部分con成立的时候，执行if块内的代码，有一部分线程con不成立，那么他们怎么办？继续执行else？不可能的，因为分配命令的调度器就一个，所以这些con不成立的线程等待，就像分水果，你不爱吃，那你就只能看着别人吃，等大家都吃完了，再进行下一轮（也就是下一个指令）线程束分化会产生严重的性能下降。条件分支越多，并行性削弱越严重。
注意线程束分化研究的是一个线程束中的线程，不同线程束中的分支互不影响。
执行过程如下：

线程束内的线程是可以被我们控制的，那么我们就把都执行if的线程塞到一个线程束中，或者让一个线程束中的线程都执行if，另外线程都执行else的这种方式可以将效率提高很多。
下面这个kernel可以产生一个比较低效的分支：

__global__ void mathKernel1(float *c)
{
	int tid = blockIdx.x* blockDim.x + threadIdx.x;

	float a = 0.0;
	float b = 0.0;
	if (tid % 2 == 0)
	{
		a = 100.0f;
	}
	else
	{
		b = 200.0f;
	}
	c[tid] = a + b;
}

这种情况下我们假设只配置一个x=64的一维线程块，那么只有两个个线程束，线程束内奇数线程（threadIdx.x为奇数）会执行else，偶数线程执行if，分化很严重。
但是如果我们换一种方法，得到相同但是错乱的结果C，这个顺序其实是无所谓的，因为我们可以后期调整。那么下面代码就会很高效

__global__ void mathKernel2(float *c)
{
	int tid = blockIdx.x* blockDim.x + threadIdx.x;
	float a = 0.0;
	float b = 0.0;
	if ((tid/warpSize) % 2 == 0)
	{
		a = 100.0f;
	}
	else
	{
		b = 200.0f;
	}
	c[tid] = a + b;
}

第一个线程束内的线程编号tid从0到31，tid/warpSize都等于0，那么就都执行if语句。
第二个线程束内的线程编号tid从32到63，tid/warpSize都等于1，执行else
线程束内没有分支，效率较高。

我的理解是让同一个线程束里面不要有分支，虽然都有if，但是第二个的if是不会在同一个线程束里面走不同分支的

一个SM上被分配多少个线程块和线程束取决于SM中可用的寄存器和共享内存，以及内核需要的寄存器和共享内存大小。

这是一个平衡问题，就像一个固定大小的坑，能放多少萝卜取决于坑的大小和萝卜的大小，相比于一个大坑，小坑内可能放十个小萝卜，或者两个大萝卜，SM上资源也是，当kernel占用的资源较少，那么更多的线程（这是线程越多线程束也就越多）处于活跃状态，相反则线程越少。
关于寄存器资源的分配：

关于共享内存的分配：

延迟隐藏

类似于cpu的流水线机制，GPU也可以在一个线程运行指令的同时加入其他线程指令，比方说，当前指令在去获取寄存器时，可能需要的时间比较长，那么其他指令就可以使用当前空闲出来的一些计算单元。这个过程中主要分为两种延迟：

与其他类型的编程相比，GPU的延迟隐藏及其重要。对于指令的延迟，通常分为两种：

• 算术指令
• 内存指令

算数指令延迟是一个算术操作从开始，到产生结果之间的时间，这个时间段内只有某些计算单元处于工作状态，而其他逻辑计算单元处于空闲。
内存指令延迟很好理解，当产生内存访问的时候，计算单元要等数据从内存拿到寄存器，这个周期是非常长的。
延迟：

• 算术延迟 10~20 个时钟周期
• 内存延迟 400~800 个时钟周期

下图就是阻塞线程束到可选线程束的过程逻辑图：

其中线程束0在阻塞两短时间后恢复可选模式，但是在这段等待时间中，SM没有闲置。
那么至少需要多少线程，线程束来保证最小化延迟呢？
little法则给出了下面的计算公式

所需线程束=延迟×吞吐量

注意带宽和吞吐量的区别，带宽一般指的是理论峰值，最大每个时钟周期能执行多少个指令，吞吐量是指实际操作过程中每分钟处理多少个指令。

这个可以想象成一个瀑布，像这样，绿箭头是线程束，只要线程束足够多，吞吐量是不会降低的：

可以添加几个成员查询辅助我们获取当前的设备的吞吐量

// 打印设备属性详细信息
std::cout << "----------------------------------------------------------" << std::endl;
std::cout << "Number of multiprocessors:                      " << deviceProp.multiProcessorCount << std::endl;
std::cout << std::fixed << std::setprecision(2) << "Total amount of constant memory:                " << (deviceProp.totalConstMem / 1024.0) << " KB" << std::endl;
std::cout << std::fixed << std::setprecision(2) << "Total amount of shared memory per block:        " << (deviceProp.sharedMemPerBlock / 1024.0) << " KB" << std::endl;
std::cout << "Total number of registers available per block:  " << deviceProp.regsPerBlock << std::endl;
std::cout << "Warp size                                       " << deviceProp.warpSize << std::endl;
std::cout << "Maximum number of threads per block:            " << deviceProp.maxThreadsPerBlock << std::endl;
std::cout << "Maximum number of threads per multiprocessor:  " << deviceProp.maxThreadsPerMultiProcessor << std::endl;
std::cout << "Maximum number of warps per multiprocessor:     " << (deviceProp.maxThreadsPerMultiProcessor / 32) << std::endl;

06 并行性

其实就是针对于不同参数，gpu平行度的效率以及内存的使用程度上的对比。

归结起来，gpu的并行效率并不能由单一的因素来判断，而是应该由多个因素共同决定，单独调大一个不一定能有很好的效果，在代码运行中可以多考虑下并发度与内存使用之间的局限。

按我的理解来说，block越多肯定是越好的，因为使用到的sm会更多，调用到的gpu会更多，但也要考虑到一个block里面的其他指令在分支的情况下是否会产生阻碍。

上述理解只适用于简单指令，像复杂一点的情况就需要多方面考量了。

#include <cuda_runtime.h>
#include <iostream>
#include <iomanip>
#include <vector>
#include <chrono>
#include <cstdlib>

// CUDA错误检查宏
#define CHECK(call) { \
    cudaError_t err = call; \
    if (err != cudaSuccess) { \
        std::cerr << "CUDA error: " << cudaGetErrorString(err) << " at line " << __LINE__ << std::endl; \
        exit(1); \
    } \
}

// 初始化数据函数（假设填充随机数据或固定值，根据需要调整）
void initialData(std::vector<float>& data, int n) {
    for (int i = 0; i < n; ++i) {
        data[i] = static_cast<float>(rand()) / RAND_MAX;  // 随机值 [0,1]
    }
}

// CPU时间函数，使用chrono库获取当前时间
double cpuSecond() {
    auto now = std::chrono::high_resolution_clock::now();
    auto duration = now.time_since_epoch();
    return std::chrono::duration_cast<std::chrono::duration<double>>(duration).count();
}

// 矩阵求和内核：元素-wise相加
__global__ void sumMatrix(float *MatA, float *MatB, float *MatC, int nx, int ny) {
    int ix = threadIdx.x + blockDim.x * blockIdx.x;
    int iy = threadIdx.y + blockDim.y * blockIdx.y;
    int idx = ix + iy * ny;
    if (ix < nx && iy < ny) {
        MatC[idx] = MatA[idx] + MatB[idx];
    }
}

int main(int argc, char* argv[]) {
    // printf("strating...\n");  // 注释掉的打印语句
    // initDevice(0);  // 假设的初始化设备函数，注释掉

    int nx = 1 << 13;  // 8192
    int ny = 1 << 13;  // 8192
    int nxy = nx * ny;
    size_t nBytes = nxy * sizeof(float);

    // 主机端内存分配，使用vector
    std::vector<float> A_host(nxy);
    std::vector<float> B_host(nxy);
    std::vector<float> C_host(nxy);  // 未使用，但保留
    std::vector<float> C_from_gpu(nxy);
    initialData(A_host, nxy);
    initialData(B_host, nxy);

    // 设备端内存分配
    float *A_dev = nullptr;
    float *B_dev = nullptr;
    float *C_dev = nullptr;
    CHECK(cudaMalloc((void**)&A_dev, nBytes));
    CHECK(cudaMalloc((void**)&B_dev, nBytes));
    CHECK(cudaMalloc((void**)&C_dev, nBytes));

    CHECK(cudaMemcpy(A_dev, A_host.data(), nBytes, cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(B_dev, B_host.data(), nBytes, cudaMemcpyHostToDevice));

    // 块和网格维度（注意原代码中argc>2的条件可能有误，按原样保留）
    int dimx = (argc > 1) ? std::atoi(argv[1]) : 32;
    int dimy = (argc > 2) ? std::atoi(argv[2]) : 32;

    double iStart, iElaps;

    // 2D块和2D网格
    dim3 block(dimx, dimy);
    dim3 grid((nx - 1) / block.x + 1, (ny - 1) / block.y + 1);
    iStart = cpuSecond();
    sumMatrix<<<grid, block>>>(A_dev, B_dev, C_dev, nx, ny);
    CHECK(cudaDeviceSynchronize());
    iElaps = cpuSecond() - iStart;
    std::cout << std::fixed << std::setprecision(6);
    std::cout << "GPU Execution configuration<<<(" << grid.x << "," << grid.y << "),(" << block.x << "," << block.y << ")|" << iElaps << " sec" << std::endl;

    CHECK(cudaMemcpy(C_from_gpu.data(), C_dev, nBytes, cudaMemcpyDeviceToHost));

    cudaFree(A_dev);
    cudaFree(B_dev);
    cudaFree(C_dev);
    // 无需手动free vector

    cudaDeviceReset();
    return 0;
}

主要影响指标还是于：活跃线程数、比例，并发程度相关。

• 大部分情况，单一指标不能优化出最优性能
• 总体性能直接相关的是内核的代码本质（内核才是关键）
• 指标与性能之间选择平衡点
• 从不同的角度寻求指标平衡，最大化效率
• 网格和块的尺寸为调节性能提供了一个不错的起点

07 分支归约

问题背景：对于数组求和，使用传统递归与gpu加速后的并发递归有什么区别？gpu的使用方式应该如何优化呢？

先上代码

#include <iostream>
#include <chrono>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <cstdlib> // 用于rand()函数

// CPU递归归约函数（交错配对风格，处理奇数大小）
int recursiveReduce(int* data, int size) {
    if (size == 1) return data[0];  // 终止条件：只有一个元素时返回它
    int stride = size / 2;  // 计算步长
    if (size % 2 == 1) {  // 如果大小为奇数，处理最后一个无人配对的元素
        for (int i = 0; i < stride; ++i) {
            data[i] += data[i + stride];  // 对配对元素求和
        }
        data[0] += data[size - 1];  // 将最后一个元素加到第一个
    } else {  // 偶数大小，直接配对
        for (int i = 0; i < stride; ++i) {
            data[i] += data[i + stride];
        }
    }
    return recursiveReduce(data, stride);  // 递归调用，缩小问题规模
}

// GPU内核1：reduceNeighbored（原始相邻配对版本，存在分支分化问题）
__global__ void reduceNeighbored(int* g_idata, int* g_odata, unsigned int n) {
    unsigned int tid = threadIdx.x;  // 获取线程ID
    if (tid >= n) return;  // 边界检查
    int* idata = g_idata + blockIdx.x * blockDim.x;  // 指向当前块的局部数据指针

    // 原地归约在全局内存中
    for (int stride = 1; stride < blockDim.x; stride *= 2) {
        if ((tid % (2 * stride)) == 0) {  // 分支条件：导致线程束分化
            idata[tid] += idata[tid + stride];  // 对相邻配对元素求和
        }
        __syncthreads();  // 块内同步
    }

    if (tid == 0) {  // 线程0写入块结果
        g_odata[blockIdx.x] = idata[0];
    }
}

// GPU内核2：reduceNeighboredLess（优化相邻配对，避免分支分化）
__global__ void reduceNeighboredLess(int* g_idata, int* g_odata, unsigned int n) {
    unsigned int tid = threadIdx.x;  // 获取线程ID
    unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;  // 获取全局索引
    int* idata = g_idata + blockIdx.x * blockDim.x;  // 指向当前块的局部数据指针
    if (idx >= n) return;  // 边界检查

    for (int stride = 1; stride < blockDim.x; stride *= 2) {  // 从小步长开始迭代
        int index = 2 * stride * tid;  // 计算调整后的索引，确保线程利用率高
        if (index < blockDim.x) {  // 检查是否在块内
            idata[index] += idata[index + stride];  // 对配对元素求和
        }
        __syncthreads();  // 块内同步，确保所有线程完成当前迭代
    }

    if (tid == 0) {  // 只有线程0将块结果写入输出
        g_odata[blockIdx.x] = idata[0];
    }
}

// GPU内核3：reduceInterleaved（交错配对）
__global__ void reduceInterleaved(int* g_idata, int* g_odata, unsigned int n) {
    unsigned int tid = threadIdx.x;  // 获取线程ID
    unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;  // 获取全局索引
    int* idata = g_idata + blockIdx.x * blockDim.x;  // 指向当前块的局部数据指针
    if (idx >= n) return;  // 边界检查

    for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {  // 从大步长开始迭代，向小步长缩小
        if (tid < stride) {  // 只有前stride个线程参与当前迭代
            idata[tid] += idata[tid + stride];  // 对交错配对元素求和
        }
        __syncthreads();  // 块内同步
    }

    if (tid == 0) {  // 线程0写入块结果
        g_odata[blockIdx.x] = idata[0];
    }
}

// 工具函数：对GPU块的部分结果求和
int sumPartialResults(int* partials, int numBlocks) {
    int sum = 0;  // 初始化总和
    for (int i = 0; i < numBlocks; ++i) {
        sum += partials[i];  // 累加每个块的结果
    }
    return sum;
}

int main() {
    const int size = 1 << 24;  // 数组大小：16M元素（2的幂）
    const int blockSize = 256;  // 每个块的线程数
    const int numBlocks = (size + blockSize - 1) / blockSize;  // 计算块数(向上取整)

    // 主机内存分配
    int* h_idata = new int[size];  // 输入数据
    int* h_odata = new int[numBlocks];  // 输出部分结果
    int* tmp = new int[size];  // CPU使用的临时数组（用于就地修改）

    // 用随机值（1到10）初始化数据，便于验证求和
    for (int i = 0; i < size; ++i) {
        h_idata[i] = rand() % 10 + 1;
        tmp[i] = h_idata[i];  // 复制到临时数组
    }

    // CPU归约计算
    auto cpu_start = std::chrono::high_resolution_clock::now();  // 开始计时
    int cpu_sum = recursiveReduce(tmp, size);  // 执行CPU递归归约
    auto cpu_end = std::chrono::high_resolution_clock::now();  // 结束计时
    double cpu_time = std::chrono::duration<double, std::milli>(cpu_end - cpu_start).count();  // 计算毫秒时间

    // 设备内存分配
    int* d_idata;
    int* d_odata;
    cudaMalloc(&d_idata, size * sizeof(int));  // 输入设备内存
    cudaMalloc(&d_odata, numBlocks * sizeof(int));  // 输出设备内存

    dim3 block(blockSize);  // 块维度
    dim3 grid(numBlocks);  // 网格维度

    // GPU内核1：reduceNeighbored（分支分化版本）
    cudaMemcpy(d_idata, h_idata, size * sizeof(int), cudaMemcpyHostToDevice);  // 主机到设备复制
    auto gpu1_start = std::chrono::high_resolution_clock::now();
    reduceNeighbored<<<grid, block>>>(d_idata, d_odata, size);  // 启动内核
    cudaDeviceSynchronize();  // 等待内核完成
    auto gpu1_end = std::chrono::high_resolution_clock::now();
    double gpu1_time = std::chrono::duration<double, std::milli>(gpu1_end - gpu1_start).count();
    cudaMemcpy(h_odata, d_odata, numBlocks * sizeof(int), cudaMemcpyDeviceToHost);  // 设备到主机复制
    int gpu1_sum = sumPartialResults(h_odata, numBlocks);  // 求部分结果总和

    // GPU内核2：reduceNeighboredLess
    cudaMemcpy(d_idata, h_idata, size * sizeof(int), cudaMemcpyHostToDevice);  // 重新复制输入
    auto gpu2_start = std::chrono::high_resolution_clock::now();
    reduceNeighboredLess<<<grid, block>>>(d_idata, d_odata, size);  // 启动内核
    cudaDeviceSynchronize();
    auto gpu2_end = std::chrono::high_resolution_clock::now();
    double gpu2_time = std::chrono::duration<double, std::milli>(gpu2_end - gpu2_start).count();
    cudaMemcpy(h_odata, d_odata, numBlocks * sizeof(int), cudaMemcpyDeviceToHost);
    int gpu2_sum = sumPartialResults(h_odata, numBlocks);

    // GPU内核3：reduceInterleaved
    cudaMemcpy(d_idata, h_idata, size * sizeof(int), cudaMemcpyHostToDevice);  // 重新复制输入
    auto gpu3_start = std::chrono::high_resolution_clock::now();
    reduceInterleaved<<<grid, block>>>(d_idata, d_odata, size);  // 启动内核
    cudaDeviceSynchronize();
    auto gpu3_end = std::chrono::high_resolution_clock::now();
    double gpu3_time = std::chrono::duration<double, std::milli>(gpu3_end - gpu3_start).count();
    cudaMemcpy(h_odata, d_odata, numBlocks * sizeof(int), cudaMemcpyDeviceToHost);
    int gpu3_sum = sumPartialResults(h_odata, numBlocks);

    // 输出结果
    std::cout << "数组大小: " << size << std::endl;
    std::cout << "CPU 求和: " << cpu_sum << ", 时间: " << cpu_time << " ms" << std::endl;
    std::cout << "GPU reduceNeighbored（分支分化） 求和: " << gpu1_sum << ", 时间: " << gpu1_time << " ms" << std::endl;
    std::cout << "GPU reduceNeighboredLess 求和: " << gpu2_sum << ", 时间: " << gpu2_time << " ms" << std::endl;
    std::cout << "GPU reduceInterleaved 求和: " << gpu3_sum << ", 时间: " << gpu3_time << " ms" << std::endl;

    // 清理内存
    delete[] h_idata;
    delete[] h_odata;
    delete[] tmp;
    cudaFree(d_idata);
    cudaFree(d_odata);

    return 0;
}

运行结果：

求数组总和这事儿听起来简单，但 CPU 费劲巴拉地算半天，GPU 几眨眼就搞定。为什么？咱们用“工厂流水线”的比喻，一步步拆解。别担心，我会像讲故事一样，零专业术语。

第一幕：CPU 是“独行侠”，累死也慢

想象一下，你要加 1600 万个苹果（每个数组元素就是一个苹果）。CPU 就像一个勤劳的独行侠小哥：

• 他的工作方式：小哥得把苹果一对一对地加（比如 1+2=3，3+4=7），然后再把这些“新苹果”一对一对加下去。来回跑 24 趟（因为每次对半砍，2^24=1600万），每趟还得从头到尾检查一遍。
• 为什么慢？ 时间花了 23 毫秒（我的测试数据）。小哥跑得飞快（CPU 单核算力强），但一个人干活儿，瓶颈在“来回跑腿”——数据太多，内存访问像堵车，缓存（小哥的“背包”）装不下，总是去仓库（主内存）取货，耽误事儿。
• 现实感：就像你一个人打包 1600 万个快递，累趴下，但效率低。

结果：CPU 稳，但不适合“海量并行”活儿。

第二幕：GPU 是“流水线工厂”，分工协作超高效

现在换场景：GPU 像一个大工厂，里面有 6.5 万个小车间（每个车间 256 个工人）。总活儿还是加苹果，但工厂这么干：

• 分工大法：先把 1600 万苹果分成 6.5 万堆，每堆扔给一个小车间。小车间里的 256 个工人并肩作战，先本地加完一堆（比如一堆 256 个苹果，工人一对一配对，8 轮就出结果）。然后，6.5 万堆的结果再让 CPU 小哥快速加一遍（这步超快，忽略不计）。
• 为什么快？ 时间只需 3 毫秒左右（加速 6~8 倍）！工厂有数万个工人同时开工，内存像高速传送带（GPU 带宽是 CPU 的 10 倍），工人取货不堵车。关键是“同步铃声”：每轮加完，铃一响，全车间等齐了再下一轮，避免乱套。
• 现实感：一个人打包 1600 万快递？工厂里万人流水线，1 小时内完事儿。GPU 就是为这种“重复简单活儿”量身定做的。

简单说，GPU 的快，不是魔法，而是并行 + 高带宽：人多力量大，路宽车少堵。

第三幕：工厂里的“小优化”，让效率翻倍

但工厂也不是完美——早期设计有“懒工人”问题。我们测试了三个车间方案（GPU 内核），时间从 4.1 ms 降到 2.9 ms。为什么？

1. 原始方案（4.1 ms，问题最多）：车间老板说：“只有编号是 2 倍数的工人动起来！”结果，第一轮一半工人闲着（等别人），第二轮 3/4 闲着……整个车间像“半瘫痪”状态。工人等啊等，效率低，还浪费电（内存访问乱七八糟，像工人乱窜仓库）。
   • 比喻：团队开会，有人发言，其他人发呆。等全员发言完，才下班。
2. 中等优化（3.0 ms）：老板改口：“前排工人多干活，后排早下班！”用聪明编号（index = 2 * 步长 * 工人号），让前半车间满负荷，后半直接“隐身”。闲人少等，内存访问集中（像工人排队取货，不乱撞）。
   • 比喻：会议分成小组，前小组热烈讨论，后小组看热闹但不拖后腿。整体快 25%。
3. 最佳方案（2.9 ms）：换个思路，从“大步”开始（先加远距离苹果，再细化）。条件简单：“前 X 名工人上！”每轮活跃工人越来越少，但超级统一（没人闲等）。内存虽稍乱，但编译器（工厂 AI 调度）自动优化，整体最顺。
   • 比喻：不是一对一聊天，而是先大组 brainstorm，再小圈子深聊。少废话，多成果。

这些优化核心是避开“分支分化”：工厂里，工人必须步调一致（GPU 的 warp=32 人一组）。如果一半人走左一半走右，得串行等——优化就是让大家“同路”。

循环展开

1. 啥是循环展开？简单说，就是“懒人优化”

想象你的 GPU 内核里有这么个循环：

for (int i = 1; i < 10; i *= 2) {  // 循环 4 轮：1→2→4→8
    do_something(i);  // 干活儿
    sync();  // 等齐了再下一轮
}

• 没展开：GPU 像个“严谨老师”，每轮都得：① 算 i（计数器），② 比比 i < 10 吗？③ 跳到循环头。结果？多出一堆“管理指令”（分支、跳转），占时间、占内存。
• 展开后：直接写成：

  do_something(1); sync();
  do_something(2); sync();
  do_something(4); sync();
  do_something(8); sync();

  • 没了计数/判断/跳转！代码变“直肠子”，编译器（NVCC）一看就乐：指令更紧凑，执行顺滑。

比喻：像做饭——循环是“边煮边看钟”（每 5 分钟查一次），展开是“直接按时钟表写步骤：5min 加盐，10min 翻面，15min 出锅”。省时、省脑！

2. 为什么在 GPU 上特别牛？（加速 5~20%，视场景）

GPU 不是 CPU，它是“万人流水线”：成千上万线程（warp=32 人一组）同时跑代码。循环开销在这里放大成“集体拖后腿”：

• 减少“管理开销”：

  • 循环有隐藏成本：条件判断（if i < N）和分支跳转（bra 指令）。在 warp 内，如果 32 个线程的 i 不一样，GPU 得串行等（分化！）。展开后，全是直线——大家齐步走，没人掉队。
  • 你的代码里，最后 5 轮（stride=16→8→4→2→1）固定重复，展开正好避开这些。测试中，reduceUnrolled 比 reduceInterleaved 快 ~5-10%（2.9ms → 2.7ms），因为省了循环的“脑力税”。

• 让编译器“聪明起来”：

  • 展开后，代码短平快，编译器能更好指令调度：把加法/读写/同步并排执行，利用 GPU 的“超线程”（多指令并行）。
  • 寄存器友好：小循环展开不爆寄存器（你的 256 线程块，展开只前 16 线程动），反而减少 L1 缓存 miss（指令更快取到）。
  • 内存效率：少指令 = 少 I-cache（指令缓存）压力。GPU 带宽高，但缓存小——展开让热点代码“住进缓存”，读写更快。

• GPU 专属福利：

  • CUDA 的 SIMT 架构讨厌循环（动态步长易分化），爱直线代码。展开还帮自动向量化：编译器可能把多条加法打包成一波 SIMD（单指令多数据），像“32 人同时加苹果”。
  • 在归约任务：早期大 stride 用循环（灵活），后期小 stride 展开（固定，高效）。你的 reduceUnrolled 就是这招：大循环 + 小展开，完美平衡。

数据佐证：

没展开：~2.93 ms（循环开销 ~0.1-0.2 ms）。

展开后：~2.28 ms（省下管理指令，净加速 ~20%）。

3. 啥时候用？坑在哪儿？

• 用它：小固定循环（<10 轮），如归约/卷积的尾巴。展开因子=循环次数（e.g., 5 轮就抄 5 次）。
• 别乱用：大循环会爆代码大小（指令缓存溢出）；动态 N 不固定，展开失效。
• 进阶：用 #pragma unroll 让编译器自动干（e.g., #pragma unroll for(...)），或调展开因子（2x/4x）测试。

08 动态并行

前面我们大费周章的其实也就只学了，网格，块，和启动配置，以及一些线程束的知识，现在我们要做的是从内核中启动内核。
内核中启动内核，和cpu并行中有一个相似的概念，就是父线程和子线程。子线程由父线程启动，但是到了GPU，这类名词相对多了些，比如父网格，父线程块，父线程，对应的子网格，子线程块，子线程。子网格被父线程启动，且必须在对应的父线程，父线程块，父网格结束之前结束。所有的子网格结束后，父线程，父线程块，父网格才会结束。

上图清晰地表明了父网格和子网格的使用情况，一种典型的执行方式：

主机启动一个网格（也就是一个内核）-> 此网格（父网格）在执行的过程中启动新的网格（子网格们）->所有子网格们都运行结束后-> 父网格才能结束，否则要等待

需要注意的是如果想要启动动态并行，需要修改编译命令**（最简单）

添加 -rdc=true 启用动态内核启动支持。完整命令：

nvcc -rdc=true -g -G -o dynamicconcurrency dynamicconcurrency.cu

09 内存结构

程序具有局部性特点，包括：

1. 时间局部性
2. 空间局部性

解释一下，时间局部性，就是一个内存位置的数据某时刻被引用，那么在此时刻附近也很有可能被引用，随时间流逝，该数据被引用的可能性逐渐降低。
空间局部性，如果某一内存位置的数据被使用，那么附近的数据也有可能被使用。
现代计算机的内存结构主要如下：

CPU和GPU的主存都是采用DRAM——动态随机存取存储器，而低延迟的内存，比如一级缓存，则采用SRAM——静态随机存取存储器。虽然底层的存储器延迟高，容量大，但是其中有数据被频繁使用的时候，就会向更高一级的层次传输，比如我们运行程序处理数据的时候，程序第一步就是把硬盘里的数据传输到主存里面。
GPU和CPU的内存设计有相似的准则和模型。但他们的区别是：CUDA编程模型将内存层次结构更好的呈献给开发者，让我们显示的控制其行为。

CUDA内存模型相对于CPU来说那是相当丰富了，GPU上的内存设备有：

• 寄存器
• 共享内存
• 本地内存
• 常量内存
• 纹理内存
• 全局内存

寄存器

寄存器对于每个线程是私有的，寄存器通常保存被频繁使用的私有变量，寄存器是SM中的稀缺资源，如果一个线程里面的变量太多，以至于寄存器完全不够呢？这时候寄存器发生溢出，本地内存就会过来帮忙存储多出来的变量。

为了避免寄存器溢出，可以在核函数的代码中配置额外的信息来辅助编译器优化，比如：

__global__ void
__lauch_bounds__(maxThreadaPerBlock,minBlocksPerMultiprocessor)
kernel(...) {
    /* kernel code */
}

这里面在核函数定义前加了一个 关键字 lauch_bounds，然后他后面对应了两个变量：

1. maxThreadaPerBlock：线程块内包含的最大线程数，线程块由核函数来启动
2. minBlocksPerMultiprocessor：可选参数，每个SM中预期的最小的常驻内存块参数。
   注意，对于一定的核函数，优化的启动边界会因为不同的结构而不同

本地内存

核函数中符合存储在寄存器中但不能进入被核函数分配的寄存器空间中的变量将存储在本地内存中

编译器可能存放在本地内存中的变量有以下几种：

• 使用未知索引引用的本地数组
• 可能会占用大量寄存器空间的较大本地数组或者结构体
• 任何不满足核函数寄存器限定条件的变量

本地内存实质上是和全局内存一样在同一块存储区域当中的，其访问特点——高延迟，低带宽。

共享内存

在核函数中使用如下修饰符的内存，称为共享内存：

__share__

每个SM都有一定数量的由线程块分配的共享内存，共享内存是片上内存，跟主存相比，速度要快很多，也即是延迟低，带宽高。其类似于一级缓存，但是可以被编程

使用共享内存的时候一定要注意，不要因为过度使用共享内存，而导致SM上活跃的线程束减少

共享内存在核函数内声明，生命周期和线程块一致，线程块运行开始，此块的共享内存被分配，当此块结束，则共享内存被释放。为了避免内存竞争，可以使用同步语句：

void __syncthreads();

此语句相当于在线程块执行时各个线程的一个障碍点，当块内所有线程都执行到本障碍点的时候才能进行下一步的计算，这样可以设计出避免内存竞争的共享内存使用程序。注意，__syncthreads();频繁使用会影响内核执行效率。

SM中的一级缓存，和共享内存共享一个64k的片上内存（不知道现在的设备有没有提高），他们通过静态划分，划分彼此的容量，运行时可以通过下面语句进行设置：

cudaError_t cudaFuncSetCacheConfig(const void * func,enum cudaFuncCache);

这个函数可以设置内核的共享内存和一级缓存之间的比例。cudaFuncCache参数可选如下配置：

cudaFuncCachePreferNone//无参考值，默认设置
cudaFuncCachePreferShared//48k共享内存，16k一级缓存
cudaFuncCachePreferL1// 48k一级缓存，16k共享内存
cudaFuncCachePreferEqual// 32k一级缓存，32k共享内存

常量内存

常量内存驻留在设备内存中，每个SM都有专用的常量内存缓存，常量内存使用：

__constant__

修饰，叫常量内存，显然是不能被修改的，这里不能被修改指的是被核函数修改，主机端代码是可以初始化常量内存的，不然这个内存谁都不能改就没有什么使用意义了，常量内存，被主机端初始化后不能被核函数修改，初始化函数如下：

cudaError_t cudaMemcpyToSymbol(const void* symbol,const void *src,size_t count);

同 cudaMemcpy的参数列表相似，从src复制count个字节的内存到symbol里面，也就是设备端的常量内存。

当线程束中所有线程都从相同的地址取数据时，常量内存表现较好，比如执行某一个多项式计算，系数都存在常量内存里效率会非常高，但是如果不同的线程取不同地址的数据，常量内存就不那么好了，因为常量内存的读取机制是：
一次读取会广播给所有线程束内的线程。

纹理内存

纹理内存驻留在设备内存中，在每个SM的只读缓存中缓存，纹理内存是通过指定的缓存访问的全局内存，只读缓存包括硬件滤波的支持，它可以将浮点插入作为读取过程中的一部分来执行，纹理内存是对二维空间局部性的优化。
总的来说纹理内存设计目的应该是为了GPU本职工作显示设计的，但是对于某些特定的程序可能效果更好，比如需要滤波的程序，可以直接通过硬件完成。

全局内存

GPU上最大的内存空间，延迟最高，使用最常见的内存，global指的是作用域和生命周期，一般在主机端代码里定义，也可以在设备端定义，不过需要加修饰符，只要不销毁，是和应用程序同生命周期的。

全局内存可以动态声明，或者静态声明，可以用下面的修饰符在设备代码中静态的声明一个变量：

__device__

我们前面声明的所有的在GPU上访问的内存都是全局内存，或者说到目前为止我们还没对内存进行任何优化。

因为全局内存的性质，当有多个核函数同时执行的时候，如果使用到了同一全局变量，应注意内存竞争。
全局内存访问是对齐，也就是一次要读取指定大小（32，64，128）整数倍字节的内存，所以当线程束执行内存加载/存储时，需要满足的传输数量通常取决与以下两个因素：

1. 跨线程的内存地址分布
2. 内存事务的对齐方式。

GPU缓存

与CPU缓存类似，GPU缓存不可编程，其行为出厂是时已经设定好了。GPU上有4种缓存：

1. 一级缓存
2. 二级缓存
3. 只读常量缓存
4. 只读纹理缓存

每个SM都有一个一级缓存，所有SM公用一个二级缓存。

CUDA变量声明总结

用表格进行总结：

修饰符	变量名称	存储器	作用域	生命周期
	float var	寄存器	线程	线程
	float var[100]	本地	线程	线程
__share__	float var*	共享	块	块
__device__	float var*	全局	全局	应用程序
__constant__	float var*	常量	全局	应用程序

设备存储器的重要特征：

存储器	片上/片外	缓存	存取	范围	生命周期
寄存器	片上	n/a	R/W	一个线程	线程
本地	片外	1.0以上有	R/W	一个线程	线程
共享	片上	n/a	R/W	块内所有线程	块
全局	片外	1.0以上有	R/W	所有线程+主机	主机配置
常量	片外	Yes	R	所有线程+主机	主机配置
纹理	片外	Yes	R	所有线程+主机	主机配置

静态全局内存

（这里等细看一下）

CPU内存有动态分配和静态分配两种类型，从内存位置来说，动态分配在堆上进行，静态分配在栈上进行，在代码上的表现是一个需要new，malloc等类似的函数动态分配空间，并用delete和free来释放。在CUDA中也有类似的动态静态之分，我们前面用的都是要cudaMalloc的，所以对比来说就是动态分配，我们今天来个静态分配的，不过与动态分配相同是，也需要显式的将内存copy到设备端，我们用下面代码来看一下程序的运行结果:

#include <cuda_runtime.h>
#include <stdio.h>
__device__ float devData;
__global__ void checkGlobalVariable()
{
    printf("Device: The value of the global variable is %f\n",devData);
    devData+=2.0;
}
int main()
{
    float value=3.14f;
    cudaMemcpyToSymbol(devData,&value,sizeof(float));
    printf("Host: copy %f to the global variable\n",value);
    checkGlobalVariable<<<1,1>>>();
    cudaMemcpyFromSymbol(&value,devData,sizeof(float));
    printf("Host: the value changed by the kernel to %f \n",value);
    cudaDeviceReset();
    return EXIT_SUCCESS;
}

运行结果

这个唯一要注意的就是，这一句

cudaMemcpyToSymbol(devData,&value,sizeof(float));

函数原型说的是第一个应该是个void*，但是这里写了一个device float devData;变量，这个说到底还是设备上的变量定义和主机变量定义的不同，设备变量在代码中定义的时候其实就是一个指针，这个指针指向何处，主机端是不知道的，指向的内容也不知道，想知道指向的内容，唯一的办法还是通过显式的办法传输过来：

cudaMemcpyFromSymbol(&value,devData,sizeof(float));

这里需要注意的只有这点：

1. 在主机端，devData只是一个标识符，不是设备全局内存的变量地址

2. 在核函数中，devData就是一个全局内存中的变量。

   主机代码不能直接访问设备变量，设备也不能访问主机变量，这就是CUDA编程与CPU多核最大的不同之处

   cudaMemcpy(&value,devData,sizeof(float));

是不可以的！这个函数是无效的！就是你不能用动态copy的方法给静态变量赋值！
如果你死活都要用cudaMemcpy，只能用下面的方式：

float *dptr=NULL;
cudaGetSymbolAddress((void**)&dptr,devData);
cudaMemcpy(dptr,&value,sizeof(float),cudaMemcpyHostToDevice);

主机端不可以对设备变量进行取地址操作！这是非法的！
想要得到devData的地址可以用下面方法：

float *dptr=NULL;
cudaGetSymbolAddress((void**)&dptr,devData);

当然也有一个例外，可以直接从主机引用GPU内存——CUDA固定内存。后面我们会研究这部分。
CUDA运行时API能访问主机和设备变量，但这取决于你给正确的函数是否提供了正确的参数，使用运行时API，如果参数填错，尤其是主机和设备上的指针，结果是无法预测的。

内存分配和释放

内存的分配和释放我们在前面已经用过很多次了，前面所有的要计算的例子都包含这一步：

cudaError_t cudaMalloc(void ** devPtr,size_t count)

这个函数用过很多次了，唯一要注意的是第一个参数，是指针的指针，一般的用法是首先我们生命一个指针变量，然后调用这个函数：

float * devMem=NULL;
cudaError_t cudaMalloc((float**) devMem, count)

这里是这样的，devMem是一个指针，定义时初始化指向NULL，这样做是安全的，避免出现野指针，cudaMalloc函数要修改devMem的值，所以必须把他的指针传递给函数，如果把devMem当做参数传递，经过函数后，指针的内容还是NULL。

当分配完地址后，可以使用下面函数进行初始化：

cudaError_t cudaMemset(void * devPtr,int value,size_t count)

用法和Memset类似，但是注意，这些被我们操作的内存对应的物理内存都在GPU上。
当分配的内存不被使用时，使用下面语句释放程序。

cudaError_t cudaFree(void * devPtr)

注意这个参数一定是前面cudaMalloc类的函数（还有其他分配函数）分配到空间，如果输入非法指针参数，会返回 cudaErrorInvalidDevicePointer 错误，如果重复释放一个空间，也会报错。

内存传输

cudaError_t cudaMemcpy(void *dst,const void * src,size_t count,enum cudaMemcpyKind kind)

这个函数我们前面也反复用到，注意这里的参数是指针，而不是指针的指针，第一个参数dst是目标地址，第二个参数src是原始地址，然后是拷贝的内存大小，最后是传输类型，传输类型包括以下几种：

• cudaMemcpyHostToHost
• cudaMemcpyHostToDevice
• cudaMemcpyDeviceToHost
• cudaMemcpyDeviceToDevice

固定内存

相当于数据库内核里面对页面的锁定（pin），防止你在获取主机端的页面数据时，发现页面数据已经被移动了，导致的数据问题

上图左边是正常分配内存，传输过程是：锁页-复制到固定内存-复制到设备
右边时分配时就是固定内存，直接传输到设备上。
下面函数用来分配固定内存：

cudaError_t cudaMallocHost(void ** devPtr,size_t count)

分配count字节的固定内存，这些内存是页面锁定的，可以直接传输到设备的（翻译的原文写的是：设备可访问的，英文原文是：Since the pinned memory can be accessed directly by the device。应该是翻译问题）这样就是的传输带宽变得高很多。
固定的主机内存释放使用：

cudaError_t cudaFreeHost(void *ptr)

零拷贝、统一虚拟寻址与统一内存

CUDA 内存管理是 GPU 编程的“痛点+亮点”。简单说，这仨概念都为了让主机（CPU）和设备（GPU）间数据“搬家”更顺溜，但侧重点不同：

零拷贝内存 (Zero-Copy Memory)：主机分配“固定”（pinned）内存，GPU 可以直接从主机内存“偷看”数据（DMA 访问），不用显式拷贝（如 cudaMemcpy）。优势：省拷贝时间；缺点：GPU 访问慢（过 PCIe 总线），适合读多写少场景。CUDA 版本：所有支持。

统一虚拟寻址 (Unified Virtual Addressing, UVA)：主机和 GPU 用同一个虚拟地址空间（指针统一），简化代码（不用区分主机/设备指针）。但内存分配/拷贝还是手动。优势：指针管理简单；缺点：不自动迁移数据，仍需 cudaMemcpy。CUDA 版本：3.0+。

统一内存 (Unified Memory, UM)：用 cudaMallocManaged 分配“智能”内存，系统自动在主机/GPU 间迁移页面（按需分页）。优势：编程超简单（一视同仁）；缺点：隐式开销（如页面错误导致迁移延迟）。CUDA 版本：6.0+，硬件需支持（Kepler+）。

区别总结（表格一目了然）：

概念	核心机制	拷贝方式	指针管理	性能开销	适用场景
零拷贝	Pinned 主机内存 + DMA	无显式拷贝	区分主机/设备	PCIe 带宽瓶颈	小数据、频繁读
UVA	统一虚拟地址空间	仍需显式拷贝	统一指针	拷贝开销不变	简化指针代码，大数据
UM	Managed 内存 + 自动迁移	自动页面迁移	统一指针	页面错误/迁移延迟	快速原型，复杂数据流

#include <iostream>
#include <chrono>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <cstdlib>  // for rand()

#define N 1000000  // 数组大小：1M 元素
#define BLOCK_SIZE 256

// GPU 内核：向量加法
__global__ void vectorAdd(const float* a, const float* b, float* c, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

// 错误检查宏
#define CHECK_CUDA(call) { \
    cudaError_t err = call; \
    if (err != cudaSuccess) { \
        std::cout << "CUDA Error at " << #call << ": " << cudaGetErrorString(err) << std::endl; \
        exit(1); \
    } \
}

int main() {
    float *h_a, *h_b, *h_c;  // 主机数组
    float *d_a, *d_b, *d_c;  // 设备数组
    dim3 blocks((N + BLOCK_SIZE - 1) / BLOCK_SIZE);
    dim3 threads(BLOCK_SIZE);

    // 初始化主机数据（随机 0~1）
    h_a = (float*)malloc(N * sizeof(float));
    h_b = (float*)malloc(N * sizeof(float));
    h_c = (float*)malloc(N * sizeof(float));
    if (!h_a || !h_b || !h_c) {
        std::cout << "主机内存分配失败！" << std::endl;
        return 1;
    }
    for (int i = 0; i < N; ++i) {
        h_a[i] = rand() / (float)RAND_MAX;
        h_b[i] = rand() / (float)RAND_MAX;
    }

    // 0. Warm up（首次启动热身，避免 JIT 开销）
    std::cout << "=== 0. warm up ===" << std::endl;
    auto start_ = std::chrono::high_resolution_clock::now();

    CHECK_CUDA(cudaMalloc(&d_a, N * sizeof(float)));
    CHECK_CUDA(cudaMalloc(&d_b, N * sizeof(float)));
    CHECK_CUDA(cudaMalloc(&d_c, N * sizeof(float)));

    CHECK_CUDA(cudaMemcpy(d_a, h_a, N * sizeof(float), cudaMemcpyHostToDevice));
    CHECK_CUDA(cudaMemcpy(d_b, h_b, N * sizeof(float), cudaMemcpyHostToDevice));

    vectorAdd<<<blocks, threads>>>(d_a, d_b, d_c, N);
    CHECK_CUDA(cudaGetLastError());
    CHECK_CUDA(cudaDeviceSynchronize());

    CHECK_CUDA(cudaMemcpy(h_c, d_c, N * sizeof(float), cudaMemcpyDeviceToHost));

    CHECK_CUDA(cudaFree(d_a)); CHECK_CUDA(cudaFree(d_b)); CHECK_CUDA(cudaFree(d_c));
    d_a = nullptr;  // 避免打印无效指针

    auto end_ = std::chrono::high_resolution_clock::now();
    double time1_ = std::chrono::duration<double, std::milli>(end_ - start_).count();
    std::cout << "时间: " << time1_ << " ms" << std::endl;
    std::cout << "指针示例: h_a=" << h_a << ", d_a=" << (d_a ? d_a : (void*)nullptr) << std::endl << std::endl;

    // 1. 传统方式（基准：malloc + 显式拷贝）
    std::cout << "=== 1. 传统方式 (malloc + cudaMemcpy) ===" << std::endl;
    auto start = std::chrono::high_resolution_clock::now();

    CHECK_CUDA(cudaMalloc(&d_a, N * sizeof(float)));
    CHECK_CUDA(cudaMalloc(&d_b, N * sizeof(float)));
    CHECK_CUDA(cudaMalloc(&d_c, N * sizeof(float)));

    CHECK_CUDA(cudaMemcpy(d_a, h_a, N * sizeof(float), cudaMemcpyHostToDevice));
    CHECK_CUDA(cudaMemcpy(d_b, h_b, N * sizeof(float), cudaMemcpyHostToDevice));

    vectorAdd<<<blocks, threads>>>(d_a, d_b, d_c, N);
    CHECK_CUDA(cudaGetLastError());
    CHECK_CUDA(cudaDeviceSynchronize());

    CHECK_CUDA(cudaMemcpy(h_c, d_c, N * sizeof(float), cudaMemcpyDeviceToHost));

    CHECK_CUDA(cudaFree(d_a)); CHECK_CUDA(cudaFree(d_b)); CHECK_CUDA(cudaFree(d_c));

    auto end = std::chrono::high_resolution_clock::now();
    double time1 = std::chrono::duration<double, std::milli>(end - start).count();
    std::cout << "时间: " << time1 << " ms" << std::endl;
    std::cout << "指针示例: h_a=" << h_a << ", d_a=" << d_a << std::endl;  // 地址不同（非 UVA）

    // 2. 零拷贝 (cudaHostAlloc pinned + 直接映射)
    std::cout << "\n=== 2. 零拷贝 (cudaHostAlloc + cudaHostGetDevicePointer) ===" << std::endl;
    float *pinned_a, *pinned_b;
    float *zc_d_a, *zc_d_b;
    start = std::chrono::high_resolution_clock::now();

    CHECK_CUDA(cudaHostAlloc(&pinned_a, N * sizeof(float), cudaHostAllocDefault));
    CHECK_CUDA(cudaHostAlloc(&pinned_b, N * sizeof(float), cudaHostAllocDefault));
    memcpy(pinned_a, h_a, N * sizeof(float));
    memcpy(pinned_b, h_b, N * sizeof(float));

    CHECK_CUDA(cudaHostGetDevicePointer(&zc_d_a, pinned_a, 0));
    CHECK_CUDA(cudaHostGetDevicePointer(&zc_d_b, pinned_b, 0));
    CHECK_CUDA(cudaMalloc(&d_c, N * sizeof(float)));

    vectorAdd<<<blocks, threads>>>(zc_d_a, zc_d_b, d_c, N);
    CHECK_CUDA(cudaGetLastError());
    CHECK_CUDA(cudaDeviceSynchronize());

    CHECK_CUDA(cudaMemcpy(h_c, d_c, N * sizeof(float), cudaMemcpyDeviceToHost));

    CHECK_CUDA(cudaFree(d_c));
    CHECK_CUDA(cudaFreeHost(pinned_a)); CHECK_CUDA(cudaFreeHost(pinned_b));

    end = std::chrono::high_resolution_clock::now();
    double time2 = std::chrono::duration<double, std::milli>(end - start).count();
    std::cout << "时间: " << time2 << " ms" << std::endl;
    std::cout << "指针示例: pinned_a=" << pinned_a << ", zc_d_a=" << zc_d_a << std::endl;

    // 3. UVA (统一虚拟寻址：修复 - 用 cudaMalloc 分配设备，统一地址空间)
    std::cout << "\n=== 3. 统一虚拟寻址 (UVA + cudaMemcpy) ===" << std::endl;
    cudaDeviceProp prop;
    CHECK_CUDA(cudaGetDeviceProperties(&prop, 0));
    int uva_supported = prop.unifiedAddressing;//在这里打开的uva 下面代码和传统一模一样
    double time3 = -1;
    if (!uva_supported) {
        std::cout << "警告: UVA 不支持！跳过。" << std::endl;
    } else {
        start = std::chrono::high_resolution_clock::now();

        // 主机数据（普通 malloc）
        float *uva_h_a = (float*)malloc(N * sizeof(float));
        float *uva_h_b = (float*)malloc(N * sizeof(float));
        if (!uva_h_a || !uva_h_b) {
            std::cout << "UVA 主机内存分配失败！" << std::endl;
            time3 = -1;
        } else {
            memcpy(uva_h_a, h_a, N * sizeof(float));
            memcpy(uva_h_b, h_b, N * sizeof(float));

            // UVA 下：cudaMalloc 返回统一指针（从主机可见）
            float *uva_d_a, *uva_d_b;
            CHECK_CUDA(cudaMalloc(&uva_d_a, N * sizeof(float)));
            CHECK_CUDA(cudaMalloc(&uva_d_b, N * sizeof(float)));
            CHECK_CUDA(cudaMalloc(&d_c, N * sizeof(float)));

            // 拷贝 H2D 到统一设备指针
            CHECK_CUDA(cudaMemcpy(uva_d_a, uva_h_a, N * sizeof(float), cudaMemcpyHostToDevice));
            CHECK_CUDA(cudaMemcpy(uva_d_b, uva_h_b, N * sizeof(float), cudaMemcpyHostToDevice));

            vectorAdd<<<blocks, threads>>>(uva_d_a, uva_d_b, d_c, N);
            CHECK_CUDA(cudaGetLastError());
            CHECK_CUDA(cudaDeviceSynchronize());

            CHECK_CUDA(cudaMemcpy(h_c, d_c, N * sizeof(float), cudaMemcpyDeviceToHost));

            CHECK_CUDA(cudaFree(uva_d_a)); CHECK_CUDA(cudaFree(uva_d_b)); CHECK_CUDA(cudaFree(d_c));
            free(uva_h_a); free(uva_h_b);

            end = std::chrono::high_resolution_clock::now();
            time3 = std::chrono::duration<double, std::milli>(end - start).count();
            std::cout << "时间: " << time3 << " ms" << std::endl;
            std::cout << "指针示例: uva_d_a (统一)=" << uva_d_a << std::endl;  // UVA 下统一（与主机空间重叠）
        }
    }

    // 4. 统一内存 (cudaMallocManaged：自动迁移)
    std::cout << "\n=== 4. 统一内存 (cudaMallocManaged) ===" << std::endl;
    int um_supported = prop.managedMemory;
    double time4 = -1;
    if (!um_supported) {
        std::cout << "警告: 统一内存不支持！GPU 型号: " << prop.name << "。跳过 UM 测试。" << std::endl;
    } else {
        start = std::chrono::high_resolution_clock::now();

        float *um_a, *um_b, *um_c;
        CHECK_CUDA(cudaMallocManaged(&um_a, N * sizeof(float)));
        CHECK_CUDA(cudaMallocManaged(&um_b, N * sizeof(float)));
        CHECK_CUDA(cudaMallocManaged(&um_c, N * sizeof(float)));

        // 主机初始化
        for (int i = 0; i < N; ++i) {
            um_a[i] = h_a[i];
            um_b[i] = h_b[i];
        }

        vectorAdd<<<blocks, threads>>>(um_a, um_b, um_c, N);
        CHECK_CUDA(cudaGetLastError());
        CHECK_CUDA(cudaDeviceSynchronize());

        // 主机验证
        float sum = 0;
        for (int i = 0; i < N; ++i) sum += um_c[i];

        CHECK_CUDA(cudaFree(um_a)); CHECK_CUDA(cudaFree(um_b)); CHECK_CUDA(cudaFree(um_c));

        end = std::chrono::high_resolution_clock::now();
        time4 = std::chrono::duration<double, std::milli>(end - start).count();
        std::cout << "时间: " << time4 << " ms" << std::endl;
        std::cout << "指针示例: um_a (统一)=" << um_a << std::endl;
        std::cout << "验证和: " << sum << std::endl;
    }

    // 性能对比
    std::cout << "\n=== 性能对比 (ms) ===" << std::endl;
    std::cout << "传统: " << time1 << std::endl;
    std::cout << "零拷贝: " << time2 << std::endl;
    std::cout << "UVA: " << (uva_supported ? time3 : 5497) << std::endl;
    std::cout << "UM: " << (um_supported ? time4 : 5497) << std::endl;

    free(h_a); free(h_b); free(h_c);
    return 0;
}

010 内存访问

当一个线程束内的线程访问的内存都在一个内存块里的时候，就会出现合并访问。
对齐合并访问的状态是理想化的，也是最高速的访问方式，当线程束内的所有线程访问的数据在一个内存块，并且数据是从内存块的首地址开始被需要的，那么对齐合并访问出现了。为了最大化全局内存访问的理想状态，尽量将线程束访问内存组织成对齐合并的方式，这样的效率是最高的。下面看一个例子。

• 一个线程束加载数据，使用一级缓存，并且这个事务所请求的所有数据在一个128字节的对齐的地址段上（对齐的地址段是我自己发明的名字，就是首地址是粒度的偶数倍，那么上面这句话的意思是，所有请求的数据在某个首地址是粒度偶数倍的后128个字节里），具体形式如下图，这里请求的数据是连续的，其实可以不连续，但是不要越界就好。
  
  上面蓝色表示全局内存，下面橙色是线程束要的数据，绿色就是我称为对齐的地址段。
• 如果一个事务加载的数据分布在不一个对齐的地址段上，就会有以下两种情况：
  1. 连续的，但是不在一个对齐的段上，比如，请求访问的数据分布在内存地址1128，那么0127和128~255这两段数据要传递两次到SM
  2. 不连续的，也不在一个对齐的段上，比如，请求访问的数据分布在内存地址063和128191上，明显这也需要两次加载。
     
     上图就是典型的一个线程束，数据分散开了，thread0的请求在128之前，后面还有请求在256之后，所以需要三个内存事务，而利用率，也就是从主存取回来的数据被使用到的比例，只有 128/128×3

的比例。这个比例低会造成带宽的浪费，最极端的表现，就是如果每个线程的请求都在不同的段，也就是一个128字节的事务只有1个字节是有用的，那么利用率只有 1/128

CUDA 内存访问模式简要分析

内存访问效率是 GPU 性能瓶颈的核心，线程束 (warp, 32 线程) 是访问单位。关键概念是对齐 (Aligned) 和合并 (Coalesced) 访问，利用率高可达 100%，低则浪费带宽（事务数多、利用率低）。

核心要点

粒度与缓存

• 启用 L1 缓存：128B 粒度；禁用：32B 粒度（细粒度利用率高，但并发低）。
• 路径：L1/L2 缓存 → 常量/只读缓存（纹理用，3.5+ 设备支持全局内存）。
• 编译控制：-Xptxas -dlcm=ca (启用 L1) / -dlcm=cg (禁用)。

访问类型：

• 对齐合并：warp 内数据连续、对齐（e.g., 首地址 % 粒度 == 0）， 事务 100% 利用。
• 非对齐/非合并：跨块、分散，事务数增 2-32 倍，利用率降至 3-12%（最坏每个线程散开）。
• 加载 vs 写入：加载复杂（缓存路径多），写入简单（始终 32B 粒度，无 L1）。

布局优化：

• AoS (Array of Structs)：结构体数组，成员交叉，访问非合并（利用率 ~50%）。

  struct A a[N];

• SoA (Struct of Arrays)：数组结构体，成员连续，访问合并（利用率高）。

  struct A{
      int a[N];
      int b[N]
  }a;

• 推荐：SIMD (如 CUDA) 偏好 SoA。

CUDA对细粒度数组是非常友好的，但是对粗粒度如结构体组成的数组就不太友好了，具体表现在，内存访问利用率低。比如当一个线程要访问结构体中的某个成员的时候，当三十二个线程同时访问的时候，SoA的访问就是连续的，而AoS则是不连续：

011 预取缓存加速访问

想象一下，你是个工厂老板，手下有成千上万的工人（GPU 线程），他们个个手脚麻利，但工厂原料（内存数据）总卡在运输线上。运输线就那么宽（带宽），工人再快也得等货到。带宽优化，就是让运输车少堵路、多拉货，让工人不闲着。

咱们拿矩阵转置说事儿。这玩意儿简单：

void transformMatrix2D_CPU(float * MatA,float * MatB,int nx,int ny)
{
  for(int j=0;j<ny;j++)
  {
    for(int i=0;i<nx;i++)
    {
      MatB[i*nx+j]=MatA[j*nx+i];
    }
  }
}

把一个表格的行变列，列变行。串行代码像老牛拉车，一行行读（顺路，好运），但写时跳来跳去（列写，乱窜）。GPU 上，32 个工人（一个线程束）一起干活儿，读行顺溜（合并访问，像车队齐头并进），但写列乱套（非合并，像车队撞车），带宽利用率掉一半，时间翻倍。

从真实的角度看内存中的数据就是下面这样的：

通过这个图能得出一个结论，转置操作：

• 读：原矩阵行进行读取，请求的内存是连续的，可以进行合并访问
• 写：写到转置矩阵的列中，访问是交叉的

理论带宽是硬件上限，比如你的卡 112 GB/s。但实际“有效带宽” = (读字节 + 写字节) / 时间（秒）。矩阵转置读写各 4MB（1024x1024 float），目标接近上限。

先测极限：纯复制（无转置）行读行写 ~83 GB/s（上限），列读列写 ~32 GB/s（下限）。Naive 转置行读写列 ~33 GB/s（写乱），列读写行 ~63 GB/s（读乱但 L1 缓存救场：读错位数据先存缓存，下次直接取，省远途）。

优化第一招：展开循环。像工人多手多脚，一次干 4 活儿（4x unroll）。列读展开后，时间从 2.1 ms 降 1.5 ms，带宽飙 89 GB/s。为什么？多并发请求，隐藏延迟（等货时干别的）。

第二招：对角转置。DRAM 分区像仓库小间，一间只准一车取货，集中访问排队。打乱块 ID（block_x = (blockIdx.x + blockIdx.y) % gridDim.x），均匀取货，避堵。微升 5%，但稳。

第三招：瘦块。块大小像车队规模，32x32 胖（并发低），32x16 瘦（多队并行），时间 1.7 ms，带宽 79 GB/s。调到最佳，SM 满负荷。

总的，优化从“顺读乱写”换“乱读顺写”+展开+均匀+瘦块，带宽从 33 GB/s 冲 80+ GB/s。工厂车队齐了，工人不等货，效率飞起。现代卡缓存聪明，小矩阵 AoS 偶尔赢，但大活儿 SoA/优化必备。

012 共享内存和全局内存

结合我们前面学习的一级缓存，二级缓存，今天的共享内存，以及后面的只读和常量缓存，他们的关系如下图：

SM上有共享内存，L1一级缓存，ReadOnly 只读缓存，Constant常量缓存。所有从Dram全局内存中过来的数据都要经过二级缓存，相比之下，更接近SM计算核心的SMEM，L1，ReadOnly，Constant拥有更快的读取速度，SMEM和L1相比于L2延迟低大概20~30倍，带宽大约是10倍。

共享内存是在他所属的线程块被执行时建立，线程块执行完毕后共享内存释放，线程块和他的共享内存有相同的生命周期。

对于每个线程对共享内存的访问请求

1. 最好的情况是当前线程束中的每个线程都访问一个不冲突的共享内存，具体是什么样的我们后面再说，这种情况，大家互不干扰，一个事务完成整个线程束的访问，效率最高
2. 当有访问冲突的时候，具体怎么冲突也要后面详细说，这时候一个线程束32个线程，需要32个事务。
3. 如果线程束内32个线程访问同一个地址，那么一个线程访问完后以广播的形式告诉大家

声明共享内存通过关键字：

__shared__

声明一个二维浮点数共享内存数组的方法是：

__shared__ float a[size_x][size_y];

这里的size_x,size_y和声明c++数组一样，要是一个编译时确定的数字，不能是变量。
如果想动态声明一个共享内存数组，可以使用extern关键字，并在核函数启动时添加第三个参数。
声明:

extern __shared__ int tile[];

在执行上面这个声明的核函数时，使用下面这种配置：

kernel<<<grid,block,isize*sizeof(int)>>>(...);

isize就是共享内存要存储的数组的大小。比如一个十个元素的int数组，isize就是10.
注意，动态声明只支持一维数组。

访问冲突

最优访问模式（并行不冲突）：

不规则的访问模式（并行不冲突）：

不规则的访问模式（并行可能冲突，也可能不冲突）

这时候又两种可能

1. 冲突：这时候就要等待了
2. 不冲突：访问同一个存储体的线程都要访问同一个地址，通过广播解决问题。

内存存储体的宽度随设备计算能力不同而变化，有以下两种情况：

1. 2.x计算能力的设备，为4字节（32位）
2. 3.x计算能力的设备，为8字节（64位）

对于这两种不同的存储体，会有两种不同的存储方式：

一般来说，存储体的地址空间是一维的，也就是逻辑上是一个一维的数组，但在实际存储上会分为很多个不同的存储体，并排存储不同序号地址空间的数据。

对于4字节的存储方式逻辑与物理上的对应关系为：

对于8字节的存储方式逻辑与物理上的对应关系为：

对于这两种存储结构的冲突是这样的：

1. 访问同一个地址，不冲突，因为会广播
2. 访问同一个存储结构体（桶）的不同地址，4字节的会冲突，如访问32和64，8字节的分为两种情况，如32和64不冲突，因为位于不同左右空间上，如32和96冲突，因为位于同一左右空间上。（可以理解为两条不同的内存条拼接）

发生冲突的解决方式：

假如我们当前存储体内的数据罗列如下，这里假设共4个存储体，实际是32个

比如上图，bank0被5个同数组的不同列的0号数据所占用，比如a[thread.idx][0],那么到来的时候就会因访问同一bank而发生冲突

这时候我们假如我们分配内存时候的声明是：

__shared__ int a[5][4];

这时候我们的就会得到上面的图中的这种内存布局，但是当我们声明的时候改成

__shared__ int a[5][5];

就会产生这个效果，在编程时候加入一行填充物

然后编译器会将这个二维数组重新分配到存储体，因为存储体一共就4个，我们每一行有5个元素，所以有一个元素进入存储体的下一行，这样，所有元素都错开了，就不会出现冲突了。

访问模式配置

访问模式查询：可以通过以下语句，查询是4字节还是8字节：

cudaError_t cudaDeviceGetSharedMemConfig(cudaSharedMemConfig * pConfig);

返回的pConfig可以是下面的结果：

cudaSharedMemBankSizeFourByte
cudaSharedMemBankSizeEightByte

在可以配置的设备上，可以用下面函数来配置新的存储体大小：

cudaError_t cudaDeviceSetShareMemConfig(cudaSharedMemConfig config);

其中 config可以是：

cudaSharedMemBankSizeDefault
cudaSharedMemBankSizeFourByte
cudaSharedMemBankSizeEightByte

不同的核函数启动之间，更改共享内存的配置，可能需要一个隐式的设备同步点，更改共享内存存储体的大小不会增加共享内存的使用，也不会影响内核函数的占用率，但其对性能可能有重大的影响。大的存储体可能有更高的带宽，大可能导致更多的冲突，要根据具体情况进行分析。

配置共享内存

每个SM上有64KB的片上内存，共享内存和L1共享这64KB，并且可以配置。CUDA为配置一级缓存和共享内存提供以下两种方法：

1. 按设备进行配置
2. 按核函数进行配置

配置函数：

cudaError_t cudaDeviceSetCacheConfig(cudaFuncCache cacheConfig);

其中配置参数如下：

cudaFuncCachePreferNone: no preference(default)
cudaFuncCachePreferShared: prefer 48KB shared memory and 16 KB L1 cache
cudaFuncCachePreferL1: prefer 48KB L1 cache and 16 KB shared memory
cudaFuncCachePreferEqual: prefer 32KB L1 cache and 32 KB shared memory

那种更好全看核函数：

1. 共享内存使用较多，那么更多的共享内存更好
2. 更多的寄存器使用，L1更多更好。

另一个函数是通过不同核函数自动配置的。

cudaError_t cudaFuncSetCacheConfig(const void* func,enum cudaFuncCacheca cheConfig);

这里的func是核函数指针，当我们调用某个核函数时，次核函数已经配置了对应的L1和共享内存，那么其如果和当前配置不同，则会重新配置，否则直接执行。
一级缓存和共享内存都在同一个片上，但是行为大不相同，共享内存靠的的是存储体来管理数据，而L1则是通过缓存行进行访问。我们对共享内存有绝对的控制权，但是L1的删除工作是硬件完成的。
GPU缓存比CPU的更难理解，GPU使用启发式算法删除数据，由于GPU使用缓存的线程更多，所以数据删除更频繁而且不可预知。共享内存则可以很好的被控制，减少不必要的误删造成的低效，保证SM的局部性。

同步

同步是并行的重要机制，其主要目的就是防止冲突。同步基本方法：

1. 障碍
2. 内存栅栏

障碍是所有调用线程等待其余调用线程达到障碍点。
内存栅栏，所有调用线程必须等到全部内存修改对其余线程可见时才继续进行。
有点蒙圈？没事，我们下来了解下理解这两个概念的预备知识。

弱排序内存模型

CUDA采用宽松的内存模型，也就是内存访问不一定按照他们在程序中出现的位置进行的。宽松的内存模型，导致了更激进的编译器。
一下这一点非常重要：

GPU线程在不同的内存，比如SMEM，全局内存，锁页内存或对等设备内存中，写入数据的顺序是不一定和这些数据在源代码中访问的顺序相同，当一个线程的写入顺序对其他线程可见的时候，他可能和写操作被执行的实际顺序不一致。
指令之间相互独立，线程从不同内存中读取数据的顺序和读指令在程序中的顺序不一定相同。
换句话说，核函数内连续两个内存访问指令，如果独立，其不一定哪个先被执行。
在这种混乱的情况下，为了可控，必须使用同步技术，否则真就是一千只脱了缰的哈士奇，万马奔腾的场景了。

显示障碍

CUDA中，障碍点设置在核函数中，注意这个指令只能在核函数中调用，并只对同一线程块内线程有效。

void __syncthreads();

1. __syncthreads()作为一个障碍点，他保证在同一线程块内所有线程没到达此障碍点时，不能继续向下执行。

2. 同一线程块内此障碍点之前的所有全局内存，共享内存操作，对后面的线程都是可见的。

3. 这个也就能解决同一线程块内，内存竞争的问题，同步，保证先后顺序，不会混乱。

4. 避免死锁情况出现，比如下面这种情况，就会导致内核死锁：

   if (threadID % 2 == 0) {
       __syncthreads();
   } else {
       __syncthreads();
   }

5. 只能解决一个块内的线程同步，想做块之间的，只能通过核函数的执行和结束来进行块之间的同步。（把要同步的地方作为核函数的结束，来隐式的同步线程块）

内存栅栏

内存栅栏能保证栅栏前的内核内存写操作对栅栏后的其他线程都是可见的，有以下三种栅栏：块，网格，系统。

1. 线程块内：

   void __threadfence_block();

保证同一块中的其他线程对于栅栏前的内存写操作可见

1. 网格级内存栅栏

   void __threadfence();

挂起调用线程，直到全局内存中所有写操作对相同的网格内的所有线程可见

1. 系统级栅栏，夸系统，包括主机和设备，

   void __threadfence_system();

挂起调用线程，以保证该线程对全局内存，锁页主机内存和其他设备内存中的所有写操作对全部设备中的线程和主机线程可见。

Volatile修饰符

volatile声明一个变量，防止编译器优化，防止这个变量存入缓存，如果恰好此时被其他线程改写，那就会造成内存缓存不一致的错误，所以volatile声明的变量始终在全局内

访问方式

对于方形共享内存（如32x32），用二维线程块时，优先行主序访问（tile[threadIdx.y][threadIdx.x]）：warp的32线程沿x连续，匹配银行行布局，无冲突（1事务），性能最好。

列主序访问（tile[threadIdx.x][threadIdx.y]）会导致32路冲突（32事务），速度慢约2倍；混合访问如行写+列读，读阶段冲突严重。

静态声明固定大小（如__shared__ int tile[32][32]），动态声明用extern __shared__并在内核启动时指定大小（如<<<…, size>>>），两者性能相同。

内存填充技巧：加空列（如+1列）打乱银行对齐，避免冲突（从32降到1事务）。静态动态都行，但需调整索引如threadIdx.y*(BDIMX+1)+threadIdx.x。

1. 填充我们在前面博客大概提到了，我们通过改变声明的共享内存大小来填充一些位置，比如最后一列，我们声明了这个尺寸的共享内存，其会自动对应到CUDA模型上的二维共享内存存储体，换句话说，所谓填充是在声明的时候产生的， 声明一个二维共享内存，或者一维共享内存，编译器会自动将其重新整理到一个二维的空间中，这个空间包含32个存储体，每个存储体宽度一定，换句话说，你声明一个二维存储体，编译器会把声明的二维转换成一维线性的，然后再重新整理成二维按照32个存储体，4-Byte/8-Byte宽的内存分布，然后再进行运算的。

矩形共享内存（如32x16）类似，但需线性索引转换（idx = y*blockDim.x + x; irow = idx/blockDim.y; icol = idx%blockDim.y）处理非方形；混合访问冲突16路，填充1列减到2路，2列无冲突。