异步并行、底层控制与系统优化

存算重叠

• ncu 显示优化的最多的 warptiling 实现 Compute % 和 Memory % 都比较低

• 整体利用率都较低，因等待长延迟操作完成而被阻塞（Stall Long Scoreboard）
• 原因是计算在进行时访存没有被有效利用；而访存在被利用时，计算利用率也比较低

延迟隐藏

• 存算重叠：重叠计算和访存操作，计算的同时进行下次需要的访存

双缓冲（Double Buffering）

• 设置两个相同的 SMEM Buffers

当前计算第 1 块数据时，异步将第 2 块数据从 HBM 读取到共享内存；算完第 1 块，立刻算第 2 块，同时预取第 3 块

• 线程束特化（Warp Specialization/Spatial Partitioning)：不同的 Warp 可以充当不同的角色，进行不同的任务

  • 前一半 warp 负责访存，后一半 warp 负责计算

异步执行（Asynchronous Execution）

核函数发射后 Host 端代码继续执行 → 设备侧函数和计算与主机侧在异步执行

CUDA 提供了异步数据拷贝（Memcpy_async）和流水线（cuda::pipeline）原语

• 生产者-消费者模型：

  • Producer（生产者）：负责数据拷贝（从 Global Memory 到 Shared Memory）
  • Consumer（消费者）：负责计算（读取 Shared Memory 中的数据进行计算）

• 流水线生命周期：

  • (Create)：初始化流水线 make_pipeline
  • Acquire：producer_acquire()，申请一块共享内存用于写入
  • (Submit)：底层发起异步拷贝任务
  • Commit：producer_commit()，提交拷贝，通知流水线数据正在路上
  • Wait：consumer_wait()，消费者等待数据到达共享内存
  • (Compute)：执行 compute() 函数进行实际计算
  • Release：consumer_release()，释放共享内存，允许生产者复用这块内存

template <size_t stages_count = 2> // 默认流水线深度为 2
__global__ void with_staging(int *global_out, int const *global_in, size_t size,
                             size_t batch_sz) {
    auto grid = cooperative_groups::this_grid();
    auto block = cooperative_groups::this_thread_block();
    assert(size == batch_sz * grid.size());

    extern __shared__ int shared[];
    // stages_count * block.size() * sizeof(int) bytes
    size_t shared_offset[stages_count];
    for (int s = 0; s < stages_count; ++s)
        shared_offset[s] = s * block.size();

    __shared__ cuda::pipeline_shared_state<cuda::thread_scope::thread_scope_block,stages_count>
        shared_state;
    auto pipeline = cuda::make_pipeline(block, &shared_state);

    auto block_batch = [&](size_t batch) -> int {
        return block.group_index().x * block.size() + grid.size() * batch;
    };

    // compute_batch: next batch to process
    // fetch_batch: next batch to fetch from global memory
    for (size_t compute_batch = 0, fetch_batch = 0; compute_batch < batch_sz;
         ++compute_batch) {
        // The outer loop iterates over the computation of the batches
        for (;
             fetch_batch < batch_sz && fetch_batch < (compute_batch + stages_count);
             ++fetch_batch) {
            // This inner loop iterates over the memory transfers, making sure that
            // the pipeline is always full
            pipeline.producer_acquire();
            size_t shared_idx = fetch_batch % stages_count;
            size_t block_batch_idx = block_batch(fetch_batch);
            cuda::memcpy_async(block, shared + shared_offset[shared_idx],
                               global_in + block_batch_idx,
                               sizeof(int) * block.size(), pipeline);
            pipeline.producer_commit();
        }

        pipeline.consumer_wait();
        int shared_idx = compute_batch % stages_count;
        int block_batch_idx = block_batch(compute_batch); // 注：图片中此变量声明后似乎未显式使用，实际计算传参用了 block_batch(compute_batch)
        compute(global_out + block_batch(compute_batch),
                shared + shared_offset[shared_idx]);
        pipeline.consumer_release();
    }
}

Submit（发起异步拷贝）和 Compute（计算）是线程独立的操作，可以同时发生；而 Acquire、Commit、Wait、Release 是共同合作的操作，负责多线程间的同步协调

数据拷贝举例

同步加载

• 数据路径：全局内存（GMEM） -> 寄存器（Register） -> 共享内存（SMEM）
• 线程必须先把数据从全局内存读进自己的寄存器，然后再写入共享内存

template <typename T>
__global__ void sync_copy_kernel(T *global, T *output, uint64_t *clock,
                                 size_t copy_count, size_t total_elements) {
    // 动态分配共享内存，将其视为字符数组，然后重新解释为类型 T 的数组
    extern __shared__ char s[];
    T *shared = reinterpret_cast<T *>(s);

    // 计算当前线程块在全局内存中的起始偏移量
    // blockIdx.x 是当前块的索引，blockDim.x 是块内线程数，copy_count 是每个线程负责拷贝的元素个数
    const size_t block_offset = blockIdx.x * blockDim.x * copy_count;
    
    // 记录当前线程开始执行的时钟周期
    uint64_t clock_start = clock64();

    // 第一个循环：从全局内存(Global Memory)加载数据到共享内存(Shared Memory)
    for (size_t i = 0; i < copy_count; ++i) {
        // 计算共享内存中的局部索引
        const size_t local_idx = blockDim.x * i + threadIdx.x;
        // 计算全局内存中的全局索引
        const size_t global_idx = block_offset + local_idx;

        // 边界检查，防止越界访问
        if (global_idx < total_elements) {
            shared[local_idx] = global[global_idx]; // 数据搬运
        }
    }

    // 【关键】线程块内同步：等待块内所有线程都完成上述全局内存->共享内存的搬运
    __syncthreads();

    // 记录同步完成后的时钟周期
    uint64_t clock_end = clock64();

    // 只有每个线程块的 0 号线程负责将耗时累加到全局时钟计数器
    if (threadIdx.x == 0) {
        // 使用 atomicAdd 进行原子加操作，防止多线程写入冲突
        // (注：这里将 uint64_t* 强转为 unsigned long long* 以适配 atomicAdd)
        atomicAdd(reinterpret_cast<unsigned long long *>(clock),
                  clock_end - clock_start);
    }

    // 第二个循环：将共享内存中的数据写回全局内存(Output)
    for (size_t i = 0; i < copy_count; ++i) {
        const size_t local_idx = blockDim.x * i + threadIdx.x;
        const size_t global_idx = block_offset + local_idx;

        if (global_idx < total_elements) {
            output[global_idx] = shared[local_idx]; // 数据搬运
        }
    }
}

Long Scoreboard Stalls → 等待长延迟操作（全局内存访问）

异步加载

数据路径：全局内存（GMEM） -> 共享内存（SMEM）

__pipeline_memcpy_async() 会异步地将 GMEM 的数据直接拷贝到 SMEM

单次拷贝的字节数sizeof(T)必须是 4、8 或 16

生命周期：Issue → Commit → Wait

template <typename T>
__global__ void async_copy_kernel(T *global, T *output, uint64_t *clock,
                                  size_t copy_count, size_t total_elements) {
    // 动态分配共享内存，并将其重新解释为 T 类型的数组
    extern __shared__ char s[];
    T *shared = reinterpret_cast<T *>(s);

    // 计算当前线程块在全局内存中的起始偏移量
    const size_t block_offset = blockIdx.x * blockDim.x * copy_count;
    // 记录当前线程的开始时钟周期
    uint64_t clock_start = clock64();

    // 循环：发起异步数据拷贝
    for (size_t i = 0; i < copy_count; ++i) {
        const size_t local_idx = blockDim.x * i + threadIdx.x;
        const size_t global_idx = block_offset + local_idx;

        if (global_idx < total_elements) {
            // 【核心】异步拷贝：将全局内存(global)的数据直接拷贝到共享内存(shared)
            // 注意：不经过寄存器，直接由硬件 DMA 引擎搬运
            __pipeline_memcpy_async(&shared[local_idx], &global[global_idx],
                                    sizeof(T));
        }
    }

    // 【提交】提交所有之前发起的异步拷贝操作
    __pipeline_commit();
    // 【等待】等待之前提交的异步拷贝全部完成（0 表示等待所有已提交的拷贝）
    __pipeline_wait_prior(0);

    // 记录等待完成后的时钟周期
    uint64_t clock_end = clock64();

    // 统计耗时：由每个线程块的 0 号线程负责累加到全局时钟计数器
    if (threadIdx.x == 0) {
        // 使用 atomicAdd 避免多线程写入冲突
        atomicAdd(reinterpret_cast<unsigned long long *>(clock),
                  clock_end - clock_start);
    }

    // 循环：将数据从共享内存写回全局内存 Output
    for (size_t i = 0; i < copy_count; ++i) {
        const size_t local_idx = blockDim.x * i + threadIdx.x;
        const size_t global_idx = block_offset + local_idx;

        if (global_idx < total_elements) {
            output[global_idx] = shared[local_idx];
        }
    }
}

• 使用异步流水线加载比同步加载快出很多
• 异步的 Compute % 和 Memory % 都高出不少，尤其是 Memory %
• 存算重叠 → 单个任务内的访存和计算重叠

多流

不显式使用的话，CUDA 会给任务默认安排进默认流（Default Stream)，默认流传入为 0 (or nullptr/NULL) ，可以显式创建和使用 Stream，叫做Named Stream

• cudaStream_t：CUDA 流（结构）
• cudaStreamCreate()：创建一个 named stream
• cudaStreamDestroy()：销毁一个 named stream

// 定义一个函数，接收四个设备端（GPU）指针
void test_two_named_streams(int *d_a, int *d_b, int *d_c, int *d_d) {
    // 1. 声明两个 CUDA 流对象
    cudaStream_t stream1, stream2;
    
    // 2. 创建两个流（named streams）
    // CHECK_CUDA 是一个自定义的宏，用于检查 CUDA API 调用是否成功
    CHECK_CUDA(cudaStreamCreate(&stream1));
    CHECK_CUDA(cudaStreamCreate(&stream2));

    // 3. 计算网格（Grid）和线程块（Block）的维度
    // BLOCK_SIZE 通常是一个宏（如 256），N 是向量长度
    dim3 blockDim(BLOCK_SIZE);
    dim3 gridDim((N + BLOCK_SIZE - 1) / BLOCK_SIZE);

    // 4. 在 stream1 中启动向量加法核函数
    // 语法：<<<gridDim, blockDim, 共享内存大小, 流对象>>>
    // 注意这里的第 4 个参数指定了 stream1，表示这个任务属于 stream1
    vectorAdd<<<gridDim, blockDim, 0, stream1>>>(d_a, d_b, d_c);
    
    // 5. 在 stream2 中启动向量乘法核函数
    // 这个任务属于 stream2，与 stream1 中的任务是独立的
    vectorMul<<<gridDim, blockDim, 0, stream2>>>(d_a, d_b, d_d);

    // 6. 设备同步：阻塞 CPU，直到 GPU 上所有流中的所有任务都执行完毕
    CHECK_CUDA(cudaDeviceSynchronize());
    
    // 7. 销毁流，释放资源
    CHECK_CUDA(cudaStreamDestroy(stream1));
    CHECK_CUDA(cudaStreamDestroy(stream2));
}

显示流和默认流的区别：在 CUDA 中，默认流具有隐式的同步特性，当默认流中有任务在执行时，它会阻塞所有其他显式流中的任务，实际上变成了串行执行

测试三个情况：

1. 默认流与默认流
2. 显式流与显式流
3. 默认流与显式流

• 显式流与显式流在并行执行，用时最短

如果确实需要在多流环境下工作，或者需要使用默认流和显式流混合：

1. 避免混用：推荐全部使用显式流（Named Streams），彻底避开默认流的同步陷阱
2. 使用非阻塞流：可以在创建显式流时加上标志位，使其不受默认流同步影响：

cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking);

1. 启用 PTDS：在编译时加入 --default-stream per-thread 编译选项，可以让每个线程拥有自己的默认流

数据传输

之前使用的是 cudaMalloc() + cudaMemcpy() 的形式

cudaMalloc() + cudaMemcpy() 是显式且同步地进行设备侧内存的分配和 CPU→GPU 间的拷贝

异步拷贝

• cudaMallocAsync()：异步分配
• cudaMemcpyAsync()：异步传输
• cudaFreeAsync()：异步释放之前异步分配的内存

cudaStream_t stream;
// 创建一个非阻塞（Non-Blocking）的显式流
// cudaStreamNonBlocking 标志使得该流不会与默认流（Default Stream）发生隐式同步
cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking);

// cudaMallocAsync 是异步分配操作，立即返回，不阻塞 CPU
// 参数分别是：指向设备指针的指针、分配的字节大小、指定在哪个流上执行分配
cudaMallocAsync(&d_a, N * sizeof(float), stream);
cudaMallocAsync(&d_b, N * sizeof(float), stream);
cudaMallocAsync(&d_c, N * sizeof(float), stream);

// cudaMemcpyAsync 是异步拷贝，将数据从主机端(h_a, h_b)拷贝到设备端(d_a, d_b)
// cudaMemcpyHostToDevice 表示拷贝方向是主机到设备
// 最后一个参数指定了流，这些拷贝操作会在该流中排队执行
cudaMemcpyAsync(d_a, h_a, N * sizeof(float), cudaMemcpyHostToDevice, stream);
cudaMemcpyAsync(d_b, h_b, N * sizeof(float), cudaMemcpyHostToDevice, stream);

// cudaFreeAsync 是异步释放操作，将之前 cudaMallocAsync 分配的显存释放
// 释放操作同样被放入指定的流中排队执行，确保不会过早释放还在使用的内存
cudaFreeAsync(d_a, stream);
cudaFreeAsync(d_b, stream);
cudaFreeAsync(d_c, stream);

• 同步分配与传输：会阻塞，调用开始会一直持续到完成，之后的操作会等待
• 异步分配与传输：不会阻塞，发射完就进行下一项操作，流中依次执行

Pinned Memory

**Pinned Memory：**通过 cudaMallocHost() 或 cudaHostAlloc() 分配的特殊主机内存

• 这里不会被换页，进行映射后 GPU 可以直接通过 DMA 进行数据获取

Zero-Copy

• 映射固定内存 (Mapped)，没有 cudaMemcpy 这个动作，GPU 核函数拿着 d_a 指针，直接读主机内存
• cudaMallocHost(): 分配至 Pinned Memory，带上 cudaHostAllocMapped 标志会映射至 GPU
• cudaHostGetDevicePointer(): 获得设备侧的对应指针
• cudaFreeHost(): 释放之前分配至 Pinned Memory 的内存

// 分配主机固定内存（Pinned Memory），并带上 cudaHostAllocMapped 标志
// 这个标志的作用是：将这块主机内存映射到 GPU 的虚拟地址空间，使 GPU 可以直接访问
cudaHostAlloc(&h_a, N * sizeof(float), cudaHostAllocMapped);
cudaHostAlloc(&h_b, N * sizeof(float), cudaHostAllocMapped);
cudaHostAlloc(&h_c, N * sizeof(float), cudaHostAllocMapped);

// 获取映射到设备（GPU）侧的对应指针
// d_a 是接收设备指针的变量，h_a 是主机侧指针，最后一个参数 0 是保留标志位
// 之后在 GPU 核函数里，就可以直接使用 d_a 来访问 h_a 所指向的主机内存
cudaHostGetDevicePointer(&d_a, h_a, 0);
cudaHostGetDevicePointer(&d_b, h_b, 0);
cudaHostGetDevicePointer(&d_c, h_c, 0);

// 释放之前用 cudaHostAlloc 分配的主机固定内存
cudaFreeHost(h_a);
cudaFreeHost(h_b);
cudaFreeHost(h_c);

普通的 Pinned + MemcpyAsync 和之前的差不多，因为本质上都是把数据从主机“搬运”到 GPU 显存

但 Zero-Copy 的性能提升了不少，数据不再需要通过 PCIe 总线搬运到 GPU 显存里，而是 GPU 计算核心直接通过 PCIe 总线去读取主机内存，Zero-Copy 看不到显式的 Memcpy

是否存在能省去CPU与GPU互相显示拷贝的操作，实现CUDA自行迁移呢？

Managed Memory

Unified Memory（统一内存）：CPU 与 GPU 看起来是共享一块内存

• cudaMallocManaged 分配的是统一内存（Unified Memory），只需一个指针（h_a），同时可以在 CPU 和 GPU 上使用，操作系统和 CUDA 驱动会自动管理这块内存在 CPU 和 GPU 之间的迁移

cudaMallocManaged(&h_a, N * sizeof(float));
cudaMallocManaged(&h_b, N * sizeof(float));
cudaMallocManaged(&h_c, N * sizeof(float));

• 释放统一内存，统一使用 cudaFree 即可，不需要像 Pinned Memory 那样使用 cudaFreeHost

cudaFree(h_a);
cudaFree(h_b);
cudaFree(h_c);

• 因为统一内存迁移需要 CUDA 自己判断什么时候怎么迁移，而不是我们自己明确指定，有一定开销
• 缺页中断：统一内存通过 Page Fault 来进行换页迁移

1. 调用 cudaMallocManaged 分配了统一内存，并在 CPU 端填入了数据
2. GPU 核函数（add_kernel）启动，它拿着指针去访问数据，却发现数据不在显存中
3. 硬件触发中断，通知 CUDA 驱动
4. CUDA 驱动暂停 GPU 的执行，临时启动 PCIe 的 DMA，把数据从主机 RAM 通过 PCIe 搬到 GPU 显存
5. 等 1.83 秒的数据搬运完成后，add_kernel 终于可以开始计算了

• 从 nsys 中能看出，确实是通过 Page fault 来进行的，这也导致第一个 kernel 的实际执行时间很长

其他api：

• cudaMemAdvise(): 给 CUDA 一些 heuristics，告诉 CUDA 使用什么样的策略进行迁移
• cudaMemPrefetchAsync(): 让 CUDA 进行异步预取

int device = -1;
cudaGetDevice(&device); // 获取当前设备 ID

// 给 CUDA 提供使用建议（Usage hints）
// 告知 CUDA：h_a 和 h_b 主要是只读的（ReadMostly），计算时通常只需要读
cudaMemAdvise(h_a, N * sizeof(float), cudaMemAdviseSetReadMostly, device);
cudaMemAdvise(h_b, N * sizeof(float), cudaMemAdviseSetReadMostly, device);
// 告知 CUDA：h_c 的首选位置在 GPU 上（PreferredLocation），即结果尽量待在显存里
cudaMemAdvise(h_c, N * sizeof(float), cudaMemAdviseSetPreferredLocation, device);

// 在计算前，主动把数据从主机内存“预取”到 GPU 显存中
cudaMemPrefetchAsync(h_a, N * sizeof(float), device, stream);
cudaMemPrefetchAsync(h_b, N * sizeof(float), device, stream);

• Unified Memory w/ Prefetch + Advise：1.2161 ms
• 预取机制允许 GPU 一边计算，一边在后台异步预取后续可能用到的数据

CUDA Graph

CUDA Graph 可以将一系列 CUDA 操作（如内核调用、内存拷贝等）组织成有向图 Graph 结构，可以实了高效的任务调度和执行

• 执行流程：捕获 -> 创图 -> 实例化 -> 执行
• 减少启动开销：批量提交操作，将多次启动的开销合并为一次
• 增强并行性和调度：图结构明确任务间的依赖关系，调度器能更好地识别可并行执行的任务
• 避免重复工作：对于重复执行的相同操作序列，图只需构建一次，便可多次执行

• 常规流程：按顺序逐个调用（发射）核函数

void regular_approach(float *d_data, int n, int block_size, cudaStream_t stream) {
    int num_blocks = (n + block_size - 1) / block_size;
    // 依次发射5个核函数
    sinKernel<<<num_blocks, block_size, 0, stream>>>(d_data, n);
    addOneKernel<<<num_blocks, block_size, 0, stream>>>(d_data, n);
    sqrtKernel<<<num_blocks, block_size, 0, stream>>>(d_data, n);
    multiplyKernel<<<num_blocks, block_size, 0, stream>>>(d_data, n);
    subtractKernel<<<num_blocks, block_size, 0, stream>>>(d_data, n);
}

• 没用 CUDA Graph 的时候核函数之间有空泡（bubble)

• CUDA Graph 可以通过流捕获（Stream Capture）的方式进行图的创建

void create_graph(cudaGraphExec_t *graph_exec, float *d_data, int n, int block_size) {
    cudaStream_t capture_stream;
    cudaGraph_t graph;
    cudaStreamCreate(&capture_stream);

    // 1. 开始捕获流
    cudaStreamBeginCapture(capture_stream, cudaStreamCaptureModeGlobal);

    // 2. 运行一遍常规流程（此时不会真正执行，而是被记录为图节点）
    regular_approach(d_data, n, block_size, capture_stream);

    // 3. 结束捕获并实例化
    cudaStreamEndCapture(capture_stream, &graph);
    cudaGraphInstantiate(graph_exec, graph, NULL, NULL, 0);

    // 清理
    cudaGraphDestroy(graph);
    cudaStreamDestroy(capture_stream);
}

void execute_graph(cudaGraphExec_t graph_exec, cudaStream_t stream) {
    // 4. 执行图
    cudaGraphLaunch(graph_exec, stream);
}

• 发射的核函数越多，提升越大

多GPU协同

节点内协同（Intra-node）

• CPU 与 GPU：通过 PCIe 连接，呈树状结构（Tree）

• GPU 与 GPU：

  • 如果 GPU 之间支持 NVLink 且已激活，它们通过 NVSwitch 构成全互联的网状结构（Mesh）

  • 如果显卡没有 NVLink，或 NVSwitch 未激活，GPU 之间只能通过 PCIe 总线通信

    • 通常需要经过 CPU 的中转，带宽受限

• GPU 与 NIC（网卡）：GPU 数据流向网络接口，通过 PCIe 总线

High-Performance GPU Server Hardware Topology and Cluster Networking

• DMA (Direct Memory Access)：允许硬件子系统（如网卡、GPU）直接访问主存，而不需要 CPU 的干预
• PCIe Peer-to-Peer (P2P)：允许 GPU 直接访问另一块 GPU 的显存，而无需将数据拷贝回 CPU 主存，大幅降低延迟和 CPU 开销
• 在大模型训练中，张量并行（Tensor Parallelism, TP） 对节点内通信要求极高, NVLink 和 NVSwitch 的存在，使得 TP 组通常被限制在同一个节点内

节点间协同（Inter-node）

• 网络通信带宽：通常使用 单向 bit/s（比特每秒） 作为基础单位（如 100 Gb/s、400 Gb/s）

  • 网络链路通常是全双工的，光纤的一根芯负责发，另一根芯负责收，所以厂商习惯按单向速率标称

• 内存/显存带宽：通常使用 双向 Byte/s（字节每秒）或 T/s（太字节每秒） 作为基础单位

  • 显存控制器在同一物理总线上进行读写操作，带宽是双向共享的

• RDMA（Remote Direct Memory Access，远程直接内存访问）

  • 允许跨节点绕过 CPU 传输。网卡直接读取本机内存并发送到远端网卡，写入远端内存

• 节点间物理网络：

  • InfiniBand：专为 RDMA 设计的高性能网络，是目前 AI 集群的主流选择
  • RoCEv2：在标准以太网上实现 RDMA，成本较低，但需要无损网络（PFC/ECN）配置，调优难度大

High-Performance GPU Server Hardware Topology and Cluster Networking

拓扑信息

查看机器上 GPU 之间、GPU 与网卡（NIC）之间的物理互联结构：nvidia-smi topo -m

• NV#（NVLink）：性能最强。通过 NVLink 连接（图中为 NV12，表示 12 根 NVLink 通道）
• PIX（PCIe x1）：跨越最多一个 PCIe 桥，通常指两块卡挂在同一个 PCIe Switch 下，带宽较高
• PXB（PCIe x2）：跨越多个 PCIe 桥，但不跨越 PCIe Host Bridge。在同一 NUMA 节点内，不同 PCIe Switch 之间的连接
• PHB（PCIe Host Bridge）：跨越 PCIe Host Bridge，通常发生在同一 NUMA 节点内的 CPU 与 PCIe 设备之间
• NODE：性能较差。跨越了 PCIe Host Bridge，同时也跨越了 NUMA 节点之间的互联，这意味着两块 GPU 需要经过 CPU 间的通信通道
• SYS（System）：性能最差。跨越了 PCIe 和 SMP（对称多处理）互联（跨 NUMA 节点），通常需要经过 QPI/UPI 总线。跨 NUMA 访问延迟极高

• 同机器内：NV12 意味着 GPU0 到 GPU7 之间通信带宽极高，做 张量并行（TP） 非常合适

• 跨节点通信：节点间依赖网卡

  • 图中显示 NIC0 和 NIC1 偏向 NUMA 0（GPU0-3），NIC2 和 NIC3 偏向 NUMA 1（GPU4-7）
  • 组建分布式集群时，应该让 GPU0-3 走 NIC0/1，GPU4-7 走 NIC2/3，避免 SYS 的跨 NUMA 慢速路径

协同多GPU

显式跨 GPU 数据拷贝

让 GPU0 和 GPU1 直接通过 PCIe/NVLink 进行数据交换（DMA 方式），避免经过 CPU 主存中转

• cudaSetDevice()：设定当前 CPU 线程正在控制的 GPU 设备
• cudaDeviceCanAccessPeer()：查询两张 GPU 之间是否具备 P2P（Peer-to-Peer）直接通信的能力
• cudaDeviceEnablePeerAccess()：激活 P2P 访问权限。一旦激活，GPU0 的显存地址可以直接被 GPU1 的内核或 DMA 引擎访问
• cudaMemcpyPeer()：执行阻塞式的跨设备数据拷贝。它底层利用 DMA 引擎，不占用 CPU 算力，但 CPU 会卡在在这一行等待传输完成

cudaSetDevice(0);
cudaMalloc(&d0, SIZE * sizeof(int)); // 在 GPU0 分配显存
cudaSetDevice(1);
cudaMalloc(&d1, SIZE * sizeof(int)); // 在 GPU1 分配显存
cudaMalloc(&d_errors, SIZE * sizeof(int)); // 分配错误标志

// Host-staged memory fallback (CPU 中转缓存)
h_stage = (int *)malloc(SIZE * sizeof(int));

// 在 GPU0 上初始化数据
initKernel<<<blocks, threads>>>(d0, SIZE, 1000);
cudaDeviceSynchronize();

// 检查 GPU0 和 GPU1 之间是否支持 P2P 访问
int canAccessPeer;
cudaDeviceCanAccessPeer(&canAccessPeer, 0, 1);

if (canAccessPeer) {
    // 激活 P2P 访问
    cudaSetDevice(0);
    cudaDeviceEnablePeerAccess(1, 0); // GPU0 允许访问 GPU1
    cudaSetDevice(1);
    cudaDeviceEnablePeerAccess(0, 0); // GPU1 允许访问 GPU0

    // 直接 P2P 传输（阻塞式）
    cudaMemcpyPeer(d1, 1, d0, 0, SIZE * sizeof(int));
} else {
    // 不支持 P2P，回退到 Host 中转
    cudaSetDevice(0);
    cudaMemcpy(h_stage, d0, SIZE * sizeof(int), cudaMemcpyDeviceToHost);
    cudaSetDevice(1);
    cudaMemcpy(d1, h_stage, SIZE * sizeof(int), cudaMemcpyHostToDevice);
}

但是，同步的 cudaMemcpyPeer 会阻塞 CPU，浪费时间，使用需要异步版本，配合Event来实现跨设备同步

• cudaEvent_t：这是一个轻量级的同步对象。主要作用是异步管理、流间同步、计时、状态查询
• cudaEventRecord()：在流中“打时间戳”或“插旗”。当流执行到该事件时，事件状态变为“已完成”
• cudaStreamWaitEvent()：跨设备/跨流同步的桥梁。让当前流（stream1）停下，一直等到指定的 Event（kernelDone）完成，才继续执行后面的任务
• cudaMemcpyPeerAsync()：非阻塞的 P2P 拷贝。它被排入 stream1 中，由 DMA 引擎在后台执行，CPU 不需要等待

cudaStream_t stream0, stream1;
cudaSetDevice(0); cudaStreamCreate(&stream0);
cudaSetDevice(1); cudaStreamCreate(&stream1);

cudaEvent_t kernelDone = NULL;
cudaEvent_t copyDone = NULL;

// 在 GPU0 的 stream0 上启动初始化内核
cudaSetDevice(0);
initKernel<<<blocks, threads, 0, stream0>>>(d0, SIZE, 1000);

if (canAccessPeer) {
    // 记录 GPU0 上内核执行完毕的事件
    cudaSetDevice(0);
    cudaEventCreate(&kernelDone);
    cudaEventRecord(kernelDone, stream0);

    // 让 GPU1 的 stream1 等待 kernelDone 事件
    cudaSetDevice(1);
    cudaStreamWaitEvent(stream1, kernelDone, 0);
    
    // 异步 P2P 传输（非阻塞）
    cudaMemcpyPeerAsync(d1, 1, d0, 0, SIZE * sizeof(int), stream1);
} else {
    // 主机中转异步版本（Staged Async Transfer）
    cudaSetDevice(0);
    cudaMemcpyAsync(h_stage, d0, SIZE * sizeof(int), cudaMemcpyDeviceToHost, stream0);
    cudaEventCreate(&copyDone);
    cudaEventRecord(copyDone, stream0); // 记录 D2H 完成

    cudaSetDevice(1);
    cudaStreamWaitEvent(stream1, copyDone, 0); // GPU1 等待 H2D 完成
    cudaMemcpyAsync(d1, h_stage, SIZE * sizeof(int), cudaMemcpyHostToDevice, stream1);
}

统一内存跨 GPU 数据拷贝

统一内存同节点内跨设备可用，不需要显式使用数据拷贝操作

但是只能节点内传输，如要进行更复杂的操作则代码会复杂很多，且性能/效率大概率不高

NCCL

NCCL (NVIDIA Collective Communications Library) 是 NVIDIA 专为多 GPU 和多节点通信打造的集合通信库，支持常用的通信操作、支持各类常见的互联技术、性能优良、使用简便、支持多种并行模型（单线程，多线程，多进程）

• Rank（进程/设备标识）：NCCL 中每个 GPU 拥有唯一的 ID（0, 1, 2…）
• ncclCommInitAll()：单线程初始化指定设备的通信组，适用于单机多卡场景
• ncclGroupStart() / ncclGroupEnd()：组合一系列通信操作，使它们作为一个整体并发执行，避免死锁，并最大化带宽利用

int devices[2] = {0, 1};
ncclComm_t comms[2];
cudaStream_t streams[2];

// 1. 初始化 NCCL 通信组（单线程即可初始化）
ncclCommInitAll(comms, 2, devices);

// 2. 为每个 GPU 创建流和分配显存
for (int i = 0; i < 2; i++) {
    cudaSetDevice(devices[i]);
    cudaMalloc(&d_data[i], bytes);
    cudaStreamCreate(&streams[i]);
}

// 3. 准备数据（GPU 0 为例）
cudaSetDevice(devices[0]);
cudaMemcpy(d_data[0], h_src.data(), bytes, cudaMemcpyHostToDevice);

// 4. 发起 P2P 发送/接收操作
ncclGroupStart(); // 开始组合通信操作
// NCCL P2P Send/Recv: 将 GPU0 的数据发送给 GPU1
ncclSend(d_data[0], N, ncclFloat, 0, comms[0], streams[0]);
ncclRecv(d_data[1], N, ncclFloat, 0, comms[1], streams[1]);
ncclGroupEnd(); // 结束组合

// 5. 流同步（等待数据传输完成）
for (int i = 0; i < 2; i++) {
    cudaSetDevice(devices[i]);
    cudaStreamSynchronize(streams[i]);
}

• 点对点通信（P2P）只能一对一，如果 8 张卡需要多个线程同时互相协作，用 P2P 写会极其复杂且效率低下
• 集合通信（Collective Communication）：通信组内的所有进程必须参与同一个集体性操作

查看相关表格

AllReduce

AllReduce 是分布式训练中最常见的操作，尤其在**数据并行（DP）**中用于梯度同步，是将各个 rank 中对应的输入数组进行规约操作，并将规约后的结果存入各 rank 的输出数组中

• AllReduce = Reduce + Broadcast：先规约到一个节点，再广播给所有节点

• cudaGetDeviceCount(): 获取 cc 2.0+ 的设备数
• ncclGetUniqueId(): 获取/初始化一个 unique id
• ncclCommInitRank(): 多线程/多进程初始化给定的通信组，所有组内线程/进程均需调用
• ncclAllReduce(): 进行 AllReduce sum 操作
• ncclCommFinalize(): 确保所有组内涉及的操作均已完成且网络相关的资源被 flush 和释放
• ncclCommDestroy(): 如已调用过ncclCommFinalize() 则仅销毁本地组内相关资源；如没调用，则会调用 ncclCommFinalize() ，此为集合操作，因为所有同节点内的 rank 均需调用

int numDevices;
cudaGetDeviceCount(&numDevices);

ncclComm_t comms[NUM_GPUS];
cudaStream_t streams[NUM_GPUS];
int *d_data[NUM_GPUS], *d_errors[NUM_GPUS];
ncclUniqueId id;

// 1. 获取/初始化唯一 ID (用于多进程/多线程建立通信)
ncclGetUniqueId(&id);

// 2. 使用多线程，每个线程绑定并控制一张 GPU
std::vector<std::thread> threads;
for (int i = 0; i < NUM_GPUS; i++) {
    threads.emplace_back([&, i]() {
        cudaSetDevice(i);
        cudaStreamCreate(&streams[i]);
        cudaMalloc(&d_data[i], DATA_SIZE * sizeof(int));
        cudaMalloc(&d_errors[i], sizeof(int));
        cudaMemset(d_errors[i], 0, sizeof(int));

        // 初始化数据: GPU0 填 1000+, GPU1 填 2000+ ...
        initKernel<<<1, DATA_SIZE, 0, streams[i]>>>(d_data[i], 1000 + i * 1000);
        cudaStreamSynchronize(streams[i]);

        // 初始化 NCCL 通信组 (多线程/多进程，需要指定 rank 和 uniqueId)
        ncclCommInitRank(&comms[i], NUM_GPUS, id, i);
    });
}
for (auto &t : threads) t.join();
threads.clear();

for (int i = 0; i < NUM_GPUS; i++) {
    threads.emplace_back([&, i]() {
        cudaSetDevice(i);
        // 执行 AllReduce Sum 操作 (in-place 操作)
        ncclAllReduce(d_data[i], d_data[i], DATA_SIZE, ncclInt, ncclSum, comms[i], streams[i]);
        cudaStreamSynchronize(streams[i]);

        // 将结果拷回主机验证
        cudaMemcpyAsync(h_result[i], d_data[i], DATA_SIZE * sizeof(int), cudaMemcpyDeviceToHost, streams[i]);
        cudaStreamSynchronize(streams[i]);
    });
}
for (auto &t : threads) t.join();
threads.clear();

// 清理 NCCL 资源
for (int i = 0; i < NUM_GPUS; i++) {
    ncclCommFinalize(comms[i]); // 确保所有操作完成，释放网络资源
    ncclCommDestroy(comms[i]);  // 销毁本地组内资源
}