cuda编程笔记（7）--多GPU上的CUDA

零拷贝内存

在流中，我们介绍了cudaHostAlloc这个函数，它有一些标志，其中cudaHostAllocMapped允许内存映射到设备，也即GPU可以直接访问主机上的内存，不用额外再给设备指针分配内存

通过下面的操作，即可让设备指针也可访问主机内存

cudaHostAlloc((void**)&a, N * sizeof(float), cudaHostAllocWriteCombined | cudaHostAllocMapped);
cudaHostGetDevicePointer(&dev_a, a, 0); // 将主机指针映射为设备可用指针

由于GPU虚拟内存空间和CPU不同，不能直接使用指针a，必须调用cudaHostGetDevicePointer函数；这样 dev_a 就是设备端可以直接访问的 host 内存。

原理简介

• 在调用 cudaHostAllocMapped 时，CUDA 会在主机申请一块 页锁定内存（pinned memory）；

• 再通过 cudaHostGetDevicePointer 把这块主机内存映射为设备端地址空间中的指针；

• 当 GPU 访问 dev_a[i] 时，会通过 PCIe 总线从主机 RAM 中取数据，实现 零拷贝访问。

所以它虽然“看起来像显存指针”，但其实访问的是主机内存。

下面用该机制重写cuda编程笔记（2.5）--简易的应用代码-CSDN博客里的矢量点乘

#ifndef __CUDACC__
#define __CUDACC__
#endif
#include <cuda_runtime.h>
#include <device_launch_parameters.h>

#include <iostream>
#include<cstdio>


#define threadsPerBlock 256
const int Blocks = 32;
const int N = Blocks * threadsPerBlock;

void error_handling(cudaError_t res) {
    if (res !=cudaSuccess) {
        std::cout << "error!" << std::endl;
    }
}
__global__ void dot(float* a, float* b, float* c) {
    __shared__ float cache[threadsPerBlock];
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    int cacheIndex = threadIdx.x;
    float temp = 0;
    if (tid < N) temp = a[tid] * b[tid];
    cache[cacheIndex] = temp;
    __syncthreads();
    for (int stride = blockDim.x / 2; stride > 0; stride>>= 1) {
        if (cacheIndex < stride)
            cache[cacheIndex] += cache[cacheIndex + stride];
        __syncthreads();
    }
    // 将每个 block 的结果写入全局内存
    if (cacheIndex == 0) {
        c[blockIdx.x] = cache[0];
    }

}
int main() {
    cudaEvent_t start, stop;
    float* a, * b, c, * partial_c;
    float* dev_a, * dev_b, * dev_partial_c;
    float elapsedTime;
    error_handling(cudaEventCreate(&start));
    error_handling(cudaEventCreate(&stop));
    //在cpu上分配内存
    error_handling(cudaHostAlloc((void**)&a, N * sizeof(float),cudaHostAllocWriteCombined|cudaHostAllocMapped));
    error_handling(cudaHostAlloc((void**)&b, N * sizeof(float), cudaHostAllocWriteCombined | cudaHostAllocMapped));
    error_handling(cudaHostAlloc((void**)&partial_c, Blocks * sizeof(float), cudaHostAllocWriteCombined | cudaHostAllocMapped));
    for (int i = 0; i < N; i++) {
        a[i] = i;
        b[i] = i * 2;
    }
    error_handling(cudaHostGetDevicePointer(&dev_a, a, 0));
    error_handling(cudaHostGetDevicePointer(&dev_b, b, 0));
    error_handling(cudaHostGetDevicePointer(&dev_partial_c, partial_c, 0));
    error_handling(cudaEventRecord(start, 0));
    dot << < Blocks, threadsPerBlock >> > (dev_a, dev_b, dev_partial_c);
    error_handling(cudaDeviceSynchronize());
    error_handling(cudaEventRecord(stop, 0));
    error_handling(cudaEventSynchronize(stop));
    error_handling(cudaEventElapsedTime(&elapsedTime, start, stop));
    c = 0;
    for (int i = 0; i < Blocks; i++)
        c += partial_c[i];
    error_handling(cudaFreeHost(a));
    error_handling(cudaFreeHost(b));
    error_handling(cudaFreeHost(partial_c));
    error_handling(cudaEventDestroy(start));
    error_handling(cudaEventDestroy(stop));
    printf("Value calculated: %f\n", c);
    printf("Time consumed:%f\n", elapsedTime);
}

 使用条件

启动多GPU

使用多个线程，就可以同时启动多个 GPU 来并行计算，这是现代 CUDA 编程中非常推荐且常用的做法。 

CUDA 的执行模型是：

• 每个 CPU 线程 通过 cudaSetDevice(id) 绑定到某个 GPU

• 每个线程可以在绑定的 GPU 上：

  • 分配显存

  • 启动 kernel

  • 执行 memcpy

  • 做同步

CUDA runtime 为每个 CPU 线程维护独立的 GPU 上下文（context），所以 不同线程绑定不同 GPU，就可以各自独立调度、执行自己的 kernel。

#include <thread>
#include <iostream>

__global__ void kernel(int id) {
    printf("Hello from GPU %d, thread %d\n", id, threadIdx.x);
}

void gpu_task(int device_id) {
    cudaSetDevice(device_id);
    kernel<<<1, 4>>>(device_id);
    cudaDeviceSynchronize();  // 等待 GPU 完成
}

int main() {
    int num_devices = 0;
    cudaGetDeviceCount(&num_devices);

    std::vector<std::thread> threads;
    for (int i = 0; i < num_devices; ++i) {
        threads.emplace_back(gpu_task, i);  // 每个线程负责一个 GPU
    }

    for (auto& t : threads) t.join(); // 等待所有线程完成
    return 0;
}

多 GPU 场景下共享主机内存

cudaHostAlloc中当flags传入cudaHostAllocPortable时

就意味着：

✅ 分配出的主机内存是跨 GPU 可见（portable）的，不属于某个特定的 GPU 上下文。

为什么多 GPU 编程中需要 cudaHostAllocPortable？

在默认情况下（无 cudaHostAllocPortable）：

• 使用 cudaHostAlloc() 分配的内存只绑定到当前 GPU 上下文；

• 如果你在另一个 GPU 上使用该内存（比如调用 cudaMemcpyAsync），就会报错或性能下降。

加上 cudaHostAllocPortable 后：

• 这块页锁定内存在所有 GPU 上都能直接访问（只要硬件支持 UVA）。

典型用法：多 GPU + Portable 内存

float *host_ptr;
cudaHostAlloc((void**)&host_ptr, N * sizeof(float), cudaHostAllocPortable);

 然后每个线程可以这样操作：

void run_on_device(int device_id, float* shared_host) {
    cudaSetDevice(device_id);
    float *dev_ptr;
    cudaMalloc(&dev_ptr, N * sizeof(float));

    // 每个 GPU 从共享主机内存拷贝数据
    cudaMemcpy(dev_ptr, shared_host, N * sizeof(float), cudaMemcpyHostToDevice);

    kernel<<<blocks, threads>>>(dev_ptr);
    cudaDeviceSynchronize();
    cudaFree(dev_ptr);
}

这样，每个 GPU 都能用同一块主机内存 shared_host 来做数据初始化、写回、交换数据等操作。

常见组合：

cudaHostAllocPortable | cudaHostAllocWriteCombined

GPU A 写结果，GPU B 读取验证

GPU A 写入 shared host memory，GPU B 读取验证是完全可能出现同步问题的

线程之间需要加同步

#include <cuda_runtime.h>
#include <iostream>
#include <thread>
#include <vector>
#include <cassert>

#define N 16

__global__ void write_kernel(int *data, int val) {
    int idx = threadIdx.x;
    if (idx < N) {
        data[idx] = val * 100 + idx;
    }
}

__global__ void read_kernel(int *data) {
    int idx = threadIdx.x;
    if (idx < N) {
        printf("GPU 1 reads: data[%d] = %d\n", idx, data[idx]);
    }
}

// GPU 0 线程函数：写入共享主机内存
void gpu0_writer(int *host_data, cudaEvent_t write_done_event) {
    cudaSetDevice(0);
    cudaStream_t stream;
    cudaStreamCreate(&stream);

    int *dev_data;
    cudaMalloc(&dev_data, N * sizeof(int));

    write_kernel<<<1, N, 0, stream>>>(dev_data, 1);

    // 将数据从设备拷贝到共享主机内存
    cudaMemcpyAsync(host_data, dev_data, N * sizeof(int), cudaMemcpyDeviceToHost, stream);

    // 记录写入完成事件
    cudaEventRecord(write_done_event, stream);

    cudaStreamSynchronize(stream);
    cudaFree(dev_data);
    cudaStreamDestroy(stream);
    std::cout << "[GPU 0] 写入完成\n";
}

// GPU 1 线程函数：等待事件后读取共享主机内存
void gpu1_reader(int *host_data, cudaEvent_t write_done_event) {
    cudaSetDevice(1);
    cudaStream_t stream;
    cudaStreamCreate(&stream);

    // 等待 GPU 0 写入完成
    cudaStreamWaitEvent(stream, write_done_event, 0);

    int *dev_data;
    cudaMalloc(&dev_data, N * sizeof(int));

    // 从共享主机内存拷贝到 GPU 1 上的显存
    cudaMemcpyAsync(dev_data, host_data, N * sizeof(int), cudaMemcpyHostToDevice, stream);

    read_kernel<<<1, N, 0, stream>>>(dev_data);

    cudaStreamSynchronize(stream);
    cudaFree(dev_data);
    cudaStreamDestroy(stream);
    std::cout << "[GPU 1] 读取完成\n";
}

int main() {
    int gpu_count = 0;
    cudaGetDeviceCount(&gpu_count);
    if (gpu_count < 2) {
        std::cerr << "需要至少两个 GPU！\n";
        return -1;
    }

    // 分配共享主机内存（portable）
    int *shared_host_data;
    cudaHostAlloc((void**)&shared_host_data, N * sizeof(int), cudaHostAllocPortable);

    // 创建用于跨 GPU 通信的事件
    cudaEvent_t write_done_event;
    cudaEventCreateWithFlags(&write_done_event, cudaEventDisableTiming); // faster event

    // 启动两个线程
    std::thread t0(gpu0_writer, shared_host_data, write_done_event);
    std::thread t1(gpu1_reader, shared_host_data, write_done_event);

    t0.join();
    t1.join();

    cudaEventDestroy(write_done_event);
    cudaFreeHost(shared_host_data);

    return 0;
}

cudaEventCreateWithFlags

事件创建：cudaEventCreateWithFlags

cudaEvent_t evt;
cudaEventCreateWithFlags(&evt, cudaEventDisableTiming); // 推荐带标志创建更轻量

cudaEventRecord

表示 之前所有 stream中的操作都完成时，该事件被标记完成。 

cudaStreamWaitEvent

cudaError_t cudaStreamWaitEvent(cudaStream_t stream, cudaEvent_t event, unsigned int flags);