CUDA C++核心庫（CCCL）

文章目錄

- CUDA C++核心庫（CCCL）
- 核心庫介紹
- - CUDA C++ 開發工具的層級范圍
  - 各層級工具的具體內容
- Thrust
- - 自動內存管理
  - 類型安全
  - 自定義分配器（頁鎖定內存）
  - 高級API替代底層操作
  - thrust::transform
  - - 基本使用
    - 幾種執行策略
  - iterator
  - load_cs
  - 高效索引mdspan
- Libcu++
- - cuda::std::tuple
  - cuda::std::variant
  - cuda::std::pair

CUDA C++核心庫（CCCL）

一、標準 C++ 組成

C++ 語言核心
標準庫（Standard Library）
- 核心功能：
  ? 通用抽象（General Purpose Abstractions）
  ? 數據結構（Data Structures）
  ? 算法庫（Algorithms）
- 價值：簡化開發流程，提升代碼健壯性，避免底層重復實現

二、CUDA C++ 組成

基礎組件：
? C++ 語言核心
? 主機標準庫（Host Standard Library）
擴展組件：
? CUDA 語言擴展（CUDA Language Extensions）
? CUDA C++ 核心庫（CUDA C++ Core Libraries, CCCL）
CCCL 核心功能：
- 異構 C++ 標準庫（Heterogeneous C++ Standard Library）
  ? 支持 CPU/GPU 協同編程
  ? 提供統一內存模型抽象
- CUDA 基礎抽象（Fundamental CUDA Abstractions）
  ? 封裝線程管理（Thread Hierarchy）
  ? 設備內存操作（Memory Management）
- 高性能并行算法庫（High-Performance Parallel Algorithms）
  ? 矩陣運算加速
  ? 數據并行處理原語

核心庫介紹

CUDA C++ 開發工具的層級范圍

橫軸意義：從左到右，工具呈現從 “高層且易用（High-Level & Productive）” 到 “底層且更具控制力（Low-Level & More Control）” 的變化。
左側起點（綠色箭頭 “Start Here”）：建議從高層工具開始，如Thrust，因其易用性高、能提升開發效率。
右側終點（紅色箭頭 “Don’t Start Here”）：不建議直接從底層工具（如 PTX Wrappers）入手，這類工具復雜且維護難度大。

在這里插入圖片描述

各層級工具的具體內容

高層工具（High-Level & Productive）
- libcu++：提供 C++ 標準庫擴展，例如cuda::std::variant和cuda::std::optional，方便使用容器及抽象化功能。
- Thrust：提供 CPU/GPU 并行算法，適用于快速開發高層算法和數據處理。
中間層工具（中等抽象層次）
- 迭代器（Fancy Iterators）：如cuda::std::span和cuda::std::mdspan，用于處理復雜數據結構。
- 設備范圍算法（Device-wide Algorithms）：用于對設備內數據進行全局操作。
- 塊范圍算法（Block-Scope Algorithms）：如cuda::memcpy_async，適合更精細的塊級控制。
- Warp 范圍算法（Warp-Scope Algorithms）：通過cuda::atomic實現線程束間的同步與控制。
底層工具（Low-Level & More Control）
- PTX Wrappers：對 PTX 匯編代碼進行封裝，適用于需要極端性能優化的場景。
- CUB：提供低級 GPU 并行算法的實現，靈活性更高，但使用復雜度也隨之增加。

Thrust

自動內存管理

解釋：Thrust容器使用RAII（資源獲取即初始化）模式，在對象生命周期結束時自動釋放內存，避免內存泄漏。傳統CUDA編程需手動調用cudaFree，容易遺漏。

傳統方法：

int* d_data;
cudaMalloc(&d_data, N * sizeof(int));
// 使用d_data
cudaFree(d_data); // 必須手動釋放，否則內存泄漏

Thrust方法：

thrust::device_vector<int> d_data(N); // 自動分配內存
// 使用d_data
// 離開作用域時自動釋放內存

類型安全

解釋：Thrust容器在編譯期進行類型檢查，防止不匹配的數據操作。傳統cudaMemcpy僅檢查指針類型，不驗證實際數據類型。

不安全示例（傳統方法）：

cuda::std::complex<float>* d_complex;
cudaMalloc(&d_complex, N * sizeof(cuda::std::complex<float>));
int* h_int = new int[N];
// 錯誤：類型不匹配，但cudaMemcpy不報錯
cudaMemcpy(d_complex, h_int, N * sizeof(int), cudaMemcpyHostToDevice);

安全示例（Thrust方法）：

thrust::device_vector<cuda::std::complex<float>> d_complex(N);
thrust::host_vector<int> h_int(N);
// 編譯錯誤：無法將int賦值給complex<float>
// d_complex = h_int; // 此操作會觸發編譯錯誤

自定義分配器（頁鎖定內存）

解釋：Thrust允許通過自定義分配器使用頁鎖定內存（pinned memory），提升主機與設備間的數據傳輸效率。

傳統方法：

float* h_data;
cudaHostAlloc(&h_data, N * sizeof(float), cudaHostAllocDefault); // 分配頁鎖定內存
// 使用h_data
cudaFreeHost(h_data); // 手動釋放

Thrust方法：

using pinned_allocator = thrust::cuda_cub::tagged_allocator<float, // 分配的元素類型thrust::cuda_cub::thread_safe_allocator_tag, // 線程安全標記thrust::cuda_cub::pinned_host_memory_resource_tag // 固定內存標記
>;// 使用頁鎖定內存的host_vector
thrust::host_vector<float, pinned_allocator> h_data(N);
// 使用h_data
// 自動釋放

高級API替代底層操作

解釋：Thrust提供高層算法（如fill_n）替代底層函數（如cudaMemset），提升代碼可讀性和安全性。

傳統方法：

float* d_data;
cudaMalloc(&d_data, N * sizeof(float));
cudaMemset(d_data, 0, N * sizeof(float)); // 字節級操作，需手動計算字節數
// 使用d_data
cudaFree(d_data);

Thrust方法：

thrust::device_vector<float> d_data(N);
thrust::fill(d_data.begin(), d_data.end(), 0.0f); // 類型安全，自動處理內存

thrust::transform

基本使用

傳統方法：

__global__ void gelu_kernel(float* out, const float* inp, int N) {int i = blockIdx.x * blockDim.x + threadIdx.x;if (i < N) {float x = inp[i];float cube = x * x * x;float cdf = 0.5f * (1.0f + tanh(0.79788456f * (x + 0.044715f * cube)));out[i] = x * cdf;}
}void gelu_forward(float* out, const float* inp, int N) {int block_size = 128;int grid_size = CEIL_DIV(N, block_size);gelu_kernel<<<grid_size, block_size>>>(out, inp, N);
}

thrust::transform方法:

使用thrust::transform方法，只需要關心每個元素的計算邏輯，不用手動管線程分工，代碼更加簡潔，突出”算法意圖“。
可以直接使用多種多種執行策略，控制計算在設備上的執行，如thrust::device、thrust::cuda::par_nosync和thrust::cuda::par_on(stream)等。

struct gelu_functor {__host__ __device__float operator()(const float& x) const {  // 運算符重載operator()，將結構體轉變為函數對象float cube = x * x * x;return x * 0.5f * (1.0f + tanh(0.79788456f * (x + 0.044715f * cube)));}
};void gelu_forward(float* out, const float* inp, int N) {thrust::device_ptr<const float> d_inp(inp);thrust::device_ptr<float> d_out(out);// d_inp輸入數據的起始位置；d_inp+N輸入數據的結束位置；d_out輸出數據的起始位置thrust::transform(thrust::device, d_inp, d_inp + N, d_out, gelu_functor());
}

幾種執行策略

thrust::device

它指定 Thrust 算法應該在 GPU 設備上執行。它是一種相對通用的設備端執行策略，使用默認的 CUDA 流
在默認情況下，使用thrust::device執行完 Thrust 算法后，主機端代碼如果繼續訪問與該操作相關的設備內存，會隱式地等待操作完成

thrust::cuda::par_nosync

指示 Thrust 算法在 GPU 設備上運行，但它強調不進行隱式同步。這意味著 Thrust 操作提交到 GPU 后，主機端代碼會立即繼續執行，而不會等待 GPU 操作完成。
利用這種策略可以實現主機端和設備端的重疊計算，提高整體的計算效率。例如，在一些復雜的計算流程中，主機端可能需要在設備端計算的同時進行一些其他的預處理或后處理操作，此時thrust::cuda::par_nosync就非常有用。
由于沒有隱式同步，如果主機端后續需要訪問設備端操作的結果，必須手動調用 CUDA 的同步函數（如cudaDeviceSynchronize() ）來確保 GPU 操作已經完成，否則會導致錯誤

thrust::cuda::par_on(stream)

通過使用特定的 CUDA 流，可以更好地管理 GPU 資源，實現更細粒度的并行。比如，在一個應用中有多個不同優先級或者不同類型的計算任務，可以分別放在不同的 CUDA 流中，利用thrust::cuda::par_on(stream)將相應的 Thrust 操作分配到合適的流上。
并發與同步管理：可以結合不同流之間的同步原語（如cudaStreamSynchronize() ）來控制不同流中操作的執行順序和依賴關系。同時，不同流中的操作在硬件支持下能夠并發執行，進一步提高 GPU 利用率。

iterator

傳統方法：

__global__ void unpermute_kernel(float* inp, float *out, int B, int N, int NH, int d) {int idx = blockIdx.x * blockDim.x + threadIdx.x;if (idx < B * NH * N * d) {int b = idx / (NH * N * d);int rest = idx % (NH * N * d);int nh_ = rest / (N * d);rest = rest % (N * d);int n = rest / d;int d_ = rest % d;int other_idx = (b * NH * N * d) + (n * NH * d) + (nh_ * d) + d_;out[other_idx] = __ldcs(&inp[idx]);}
}// 核函數調用示例
num_blocks = CEIL_DIV(B * T * C, block_size);
unpermute_kernel<<<num_blocks, block_size>>>(vaccum, out, B, T, NH, HS);

優化方法：

// Thrust庫迭代器方式實現
auto map = thrust::make_transform_iterator(thrust::make_counting_iterator(0),  // 創建從 0 開始的計數迭代器，提供連續的索引值 idx[=] __host__ __device__ (int idx) {   // 通過 lambda 表達式作為變換操作，[=] 表示按值捕獲外部變量auto [b, n, nh_, d_] = i2n(idx, NH, T, HS);return (b * NH * T * HS) + (n * NH * HS) + (nh_ * HS) + d_;}
);cub::CacheModifiedInputIterator<cub::LOAD_CS, float> vaccumcs(vaccum);
// 該操作會根據 map 提供的索引，把 vaccumcs 指向的數據按序分散到 out 數組對應位置
thrust::scatter(thrust::device, vaccumcs, vaccumcs + B * T * C, map, out);

load_cs

thrust的高級抽象不會犧牲底層控制能力。可以通過cub::CacheModifiedInputIterator<cub::LOAD_CS, float>來等價cuda kernel中直接調用底層的__ldcs指令。

傳統方法：

// CUDA 核函數，實現兩個數組對應元素相加，結果存入輸出數組
__global__ void residual_forward_kernel(float* out, float* inp1, float* inp2, int N) {int idx = blockIdx.x * blockDim.x + threadIdx.x;if (idx < N) {out[idx] = __ldcs(&inp1[idx]) + __ldcs(&inp2[idx]);}
}// 主機端函數，用于調用上述 CUDA 核函數
void residual_forward(float* out, float* inp1, float* inp2, int N) {const int block_size = 256;const int grid_size = CEIL_DIV(N, block_size); residual_forward_kernel<<<grid_size, block_size>>>(out, inp1, inp2, N);cudaCheck(cudaGetLastError()); 
}

thrust：

void residual_forward(float* out, float* inp1, float* inp2, int N) {cub::CacheModifiedInputIterator<cub::LOAD_CS, float> inp1cs(inp1);cub::CacheModifiedInputIterator<cub::LOAD_CS, float> inp2cs(inp2);thrust::transform(thrust::device,inp1cs, inp1cs + N, inp2cs, out, thrust::plus<float>());
}

高效索引mdspan

使用cuda::std::mdspan管理多維數據，從而將傳統的按照一維數組訪問的方式，轉變為以多維方式進行訪問。

傳統方法：

// permute_kernel函數實現矩陣重排列操作
__global__ void permute_kernel(float* q, float* k, float* v,const float* inp, int B, int N, int NH, int d) {// 計算當前線程的全局索引int idx = blockIdx.x * blockDim.x + threadIdx.x;// 原始代碼中的矩陣重排列計算// dlb[nh][n][d_] = inp[b][n][nh][d_]// 計算輸入張量的各個維度索引int b = idx / (NH * N * d);      // batch維度int rest = idx % (NH * N * d);   // 剩余部分int nh = rest / (N * d);         // head維度rest = rest % (N * d);           // 繼續分解剩余部分int n = rest / d;                // 序列長度維度int d_ = rest % d;               // 特征維度// 計算輸入張量的線性索引int inp_idx = (b * N * NH * d) +           // batch偏移(n * NH * d) +               // 序列長度偏移(nh * d) +                   // head偏移d_;                          // 特征維度偏移// 執行張量重排列操作q[idx] = __ldcs(&inp[inp_idx]); // 使用__ldcs進行緩存優化的內存讀取k[idx] = __ldcs(&inp[inp_idx + NH * d]);v[idx] = __ldcs(&inp[inp_idx + 2 * NH * d]);
}// attention_forward函數實現注意力前向傳播
void attention_forward(float* out, float* veccum, float* qkv, float* presft, float* att,int B, int T, int C, int NH) {const int block_size = 256;              // CUDA線程塊大小const int softmax_block_size = 256;      // Softmax操作的線程塊大小int HS = C / NH;                        // 每個head的維度大小// 計算每個head的維度大小float *q, *k, *v;q = qkv;                                 // 查詢矩陣Q的起始位置k = qkv + B * T * C;                     // 鍵矩陣K的起始位置v = qkv + 2 * B * T * C;                 // 值矩陣V的起始位置// 計算需要的CUDA線程塊數量int total_threads = B * NH * T * HS;int num_blocks = CEIL_DIV(total_threads, block_size);// 啟動permute_kernel進行張量重排列permute_kernel<<<num_blocks, block_size>>>(q, k, v, qkv, B, T, NH, HS);
}

優化方法：

void attention_forward(float* out, float* vaccum, float* qkvr, float* prestt, float* att,float* inp, int B, int T, int C, int NH) {// 設置CUDA塊大小常量const int block_size = 256;const int softmax_block_size = 256;// 計算每個注意力頭的維度大小int HS = C / NH;  // head size// 設置Q、K、V矩陣的指針，它們在內存中是連續存儲的float *q, *k, *v;q = qkvr + 0 * B * T * C;      // Q矩陣起始位置k = qkvr + 1 * B * T * C;      // K矩陣起始位置v = qkvr + 2 * B * T * C;      // V矩陣起始位置// 使用CUDA動態內存分配constexpr auto dyn = cuda::std::dynamic_extent;  // 這是一個編譯時常量，表示數組維度的大小在運行時確定using ext_t = cuda::std::extent<int, dyn, dyn, 3, dyn, dyn>;  // 定義了一個 5 維數組的維度結構using mds_t = cuda::std::mdspan<const float, ext_t>;  // 不擁有內存，只是提供多維索引到一維內存的映射；提高代碼可讀性和維護性，避免手動計算偏移量// 創建多維數組視圖，用于更方便地訪問數據ext_t extents(B, T, NH, HS);   // 只傳入動態維度，extents仍然是5維的mds_t inp_md(inp, extents);  // 將一維內存指針 inp 映射為多維視圖/*** 示例訪問方式：* 訪問批次b、時間步t、第0個矩陣(對應Q)、頭nh_、維度hs的數據* float value = inp_md(b, t, 0, nh_, hs);*/// 使用thrust庫創建迭代器，用于并行處理auto begin = thrust::make_counting_iterator(0);auto end = begin + B * NH * T * T;// 原始重排列操作的注釋：Q[b][nh][t][d_] = inp[b][t][nh][d_]// 使用thrust并行處理每個元素// [=]：捕獲方式為值捕獲（By Value），表示 Lambda 內部可以使用外部作用域的所有變量（如B, T, C, NH, q, k, v, inp_md等）thrust::for_each(thrust::cuda::par,begin, end,[=] __device__ (int idx) {// 計算當前處理位置的各個維度索引auto [b, t, nh_, hs] = idx2(idx, NH, T, HS);// 執行Q、K、V矩陣的數據重排列q[idx] = inp_md(b, t, 0, nh_, hs);  // Q矩陣賦值k[idx] = inp_md(b, t, 1, nh_, hs);  // K矩陣賦值v[idx] = inp_md(b, t, 2, nh_, hs);  // V矩陣賦值});
}

Libcu++

cuda::std::tuple

傳統方法：

__global__ void permute_kernel(float* q, float* k, float* v,const float* inp, int B, int N, int NH, int d) {int idx = blockIdx.x * blockDim.x + threadIdx.x;if (idx < B * NH * N * d) {int b = idx / (NH * N * d);int rest = idx % (NH * N * d);int nh_ = rest / (N * d);rest = rest % (N * d);int n = rest / d;int d_ = rest % d;//...}
}__global__ void unpermute_kernel(float* inp, float *out, int B,int N, int NH, int d) {int idx = blockIdx.x * blockDim.x + threadIdx.x;if (idx < B * NH * N * d) {int b = idx / (NH * N * d);int rest = idx % (NH * N * d);int nh_ = rest / (N * d);rest = rest % (N * d);int n = rest / d;int d_ = rest % d;//...}
}

優化方法：

使用cuda::std::tuple減少了代碼冗余

__host__ __device__ 
cuda::std::tuple<int, int, int, int> 
idx2n(int idx, int E1, int E2, int E3) {int b = idx / (E1 * E2 * E3);int rest = idx % (E1 * E2 * E3);int nh_ = rest / (E2 * E3);rest = rest % (E2 * E3);int t = rest / E3;int hs = rest % E3;return {b, t, nh_, hs};
}__global__ void permute_kernel(float* q, float* k, float* v,const float* inp, int B, int N, int NH, int d) {int idx = blockIdx.x * blockDim.x + threadIdx.x;if (idx < B * NH * N * d) {auto [b, n, nh_, d_] = idx2n(idx, NH, N, d);//...}
}__global__ void unpermute_kernel(float* inp, float *out, int B,int N, int NH, int d) {int idx = blockIdx.x * blockDim.x + threadIdx.x;if (idx < B * NH * N * d) {auto [b, n, nh_, d_] = idx2n(idx, NH, N, d);//...}
}

cuda::std::variant

cuda::std::variant 是一種類型安全的聯合（tagged union ）。它可以在同一時間存儲一個值，這個值的類型可以是在聲明 variant 時指定的幾種類型中的任意一種。與傳統的 union 不同，variant 知道當前存儲的值具體是哪種類型，并且提供了安全的訪問方式。

__global__ void variant_kernel() {// 聲明一個可以存儲 int 或 float 類型的 variantcuda::std::variant<int, float> var; var = 42;  // 存儲 int 類型的值if (auto *i = cuda::std::get_if<int>(&var)) {// 安全地獲取 int 類型的值std::cout << "The value is an int: " << *i << std::endl; }var = 3.14f;  // 存儲 float 類型的值if (auto *f = cuda::std::get_if<float>(&var)) {std::cout << "The value is a float: " << *f << std::endl; }
}

cuda::std::pair

cuda::std::pair 是一個簡單的模板類，用于將兩個不同類型的對象組合成一個單一的對象。它有兩個公共成員變量 first 和 second，可以方便地訪問存儲的兩個值。

__global__ void pair_kernel() {// 創建一個 pair，存儲 int 和 float 類型的值cuda::std::pair<int, float> my_pair(10, 2.5f); std::cout << "First value: " << my_pair.first << std::endl;std::cout << "Second value: " << my_pair.second << std::endl;
}

主要參考：