零拷貝內存
在流中，我們介紹了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);
}
| 優點 | 說明 |
|---|---|
減少顯式 cudaMemcpy 調用 | 主機 → 設備零拷貝 |
| 避免重復申請/釋放顯存 | 數據只分配一次 |
| 簡化代碼結構 | 多個內核之間共享同一 host 指針 |
| 適合小規模、實時更新場景 | 如 GUI 控件、攝像頭圖像 |
| 缺點 | 說明 |
|---|---|
| 訪問速度遠慢于 global memory | 因為要通過 PCIe |
| 僅適用于某些 GPU（如支持 UVA） | 非所有設備支持 |
| 最佳性能只在小數據量/零延遲訪問場景 | 比如小型圖像處理、調試等 |
| 受限于 CPU 內存頁 | 頁大小影響效率，不能高并發 |
?使用條件
| 要點 | 說明 |
|---|
| GPU 必須支持 UVA（統一虛擬地址空間） | 可用 cudaGetDeviceProperties() 查詢 unifiedAddressing 是否為 1 |
最好配合 WriteCombined | 適合只寫不讀場景（如從主機寫入，GPU 讀取） |
| 不適合大規模數據訓練/推理 | 會嚴重拖慢 GPU 性能，PCIe 帶寬遠小于顯存帶寬 |
啟動多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); // 推薦帶標志創建更輕量
| 標志 | 含義 | 說明 |
|---|---|---|
cudaEventDefault | 默認行為 | 會記錄耗時，可用于性能計時 |
cudaEventDisableTiming | 禁用計時功能 | 更輕量，推薦用于同步控制 |
cudaEventInterprocess | 可用于多進程共享事件 | 不常用于多 GPU 同步（屬于高級功能） |
cudaEventRecord
表示 之前所有 stream中的操作都完成時，該事件被標記完成。?
cudaStreamWaitEvent
cudaError_t cudaStreamWaitEvent(cudaStream_t stream, cudaEvent_t event, unsigned int flags);
| 參數 | 類型 | 含義 |
|---|---|---|
stream | cudaStream_t | 要等待事件的 CUDA 流。這個 stream 將在 event 被觸發后才開始執行其后續任務。 |
event | cudaEvent_t | 要等待的事件。這個事件應該在其他設備或流上通過 cudaEventRecord 創建。 |
flags | unsigned int | 當前必須設為 0。（CUDA 12.4 以前不支持其他選項） |
![《[系統底層攻堅] 張冬〈大話存儲終極版〉精讀計劃啟動——存儲架構原理深度拆解之旅》-系統性學習筆記（適合小白與IT工作人員）](http://pic.xiahunao.cn/《[系統底層攻堅] 張冬〈大話存儲終極版〉精讀計劃啟動——存儲架構原理深度拆解之旅》-系統性學習筆記（適合小白與IT工作人員）)
)
