C 標準函式庫、Lambda 與多型包裝器 (C Standard Library, Lambdas and Polymorphic Wrappers)

重點總覽

device code 可用一小組 C standard library 函式（clock()/printf()/memcpy()/malloc()/alloca() 等），各有 CUDA 專屬語意與限制。Lambda 的執行空間由最內層 enclosing function 的執行空間決定；要把 lambda 傳進 kernel 必須是 __device__ 或 __host__ __device__，且常需 extended lambda（--extended-lambda 旗標）。Extended lambda 有大量限制（this 捕捉、__CUDA_ARCH__、捕捉方式、return type 內省等）。nvstd::function 提供 host/device 皆可用的多型函式包裝器，但不可跨 host/device 邊界傳遞。

device 可用的 C standard library functions

device code 只能用一小組 C 標準函式，且多數有 CUDA 專屬語意。建議優先用 <cuda/std/...> 標頭內的 cuda::std:: 對應版本（較安全、可攜）。

device printf()

__host__ __device__ __tile__ int printf(const char* format[, arg, ...]);

• 行為類似標準 C printf()，但每個 thread 各自執行：multi-threaded kernel 中每個遇到 printf() 的 thread 都會在 host stream 印出一份輸出。
• 回傳值與標準 C 不同：回傳解析的引數數（非印出的字元數）。無引數回傳 0；format 為 NULL 回傳 -1；內部錯誤回傳 -2。
• 內部用共享資料結構，可能改變 thread 執行順序；但 CUDA 本就不保證執行順序（除 __syncthreads() 屏障外）。

Format specifiers：%[flags][width][.precision][size]type

tile code 中的差異：引數可為 tile（每元素依對應 specifier 印出）；回傳值永遠是提供的引數數（即使 NULL 或內部錯誤）；format string 必須是 literal；引數數與 specifier 數不符會報錯（SIMT device code 中只是 warning）。

Host-side buffer：固定大小、環狀緩衝，預設 1 MB（cudaLimitPrintfFifoSize 可調）。滿了會覆蓋舊輸出。只在下列事件 flush：

flush 時機：
  ├─ kernel launch（<<<>>> 或 cuLaunchKernel）啟動時
  │   └─ 若 CUDA_LAUNCH_BLOCKING=1，啟動結束時也 flush
  ├─ 同步：cudaDeviceSynchronize / cuCtxSynchronize / cudaStreamSynchronize
  │        / cuStreamSynchronize / cudaEventSynchronize / cuEventSynchronize
  ├─ blocking memcpy：cudaMemcpy* / cuMemcpy*
  ├─ module load/unload：cuModuleLoad / cuModuleUnload
  ├─ context 銷毀：cudaDeviceReset / cuCtxDestroy
  └─ 執行 cudaLaunchHostFunc / cuLaunchHostFunc 加入的 stream callback 之前
程式結束時 ⇒ 不會自動 flush

device malloc() / free() 與 heap

• malloc()（device 端）、cuda::std::malloccalloc( 從 device heap 配置至少 size bytes，回傳對齊 16-byte 的指標；不足回傳 NULL。
• __nv_aligned_device_malloc() / cuda::std::aligned_alloc() 回傳 align 倍數位址，align 須為非零 2 的次方。
• 配置的記憶體在 CUDA context 生命週期內持續存在（除非被 free() 釋放），可被其他 thread、甚至後續 kernel launch 使用；任何 thread 都可釋放他人配置的記憶體，但同一指標不可重複 free（UB）。

常見配置模式：per-thread（每 thread 各自 malloc/free）、per-block（thread 0 配置、經 shared memory 分享指標、僅一 thread free）、跨 kernel 持續（__device__ 全域指標陣列保存配置，多次 kernel 重用後再 free）。

alloca()

__host__ __device__ void* alloca(size_t size);

• 在 caller 的 stack frame 配置 size bytes，回傳指標；device 端起始位址 16-byte 對齊。caller 回傳時自動釋放。

Lambda 的執行空間與傳給 kernel

compiler 依 lambda 所在最內層 enclosing function scope 的執行空間決定 lambda（closure type, C++11）的執行空間；若無 enclosing function scope，則為 __host__。也可用 extended lambda 語法明確標註。

傳給 __global__ 函式參數

• lambda/closure 只能在執行空間為 __device__ 或 __host__ __device__ 時當作 __global__ 函式的引數。
• global 或 namespace scope 的 lambda 不可當 __global__ 引數（即使是 extended lambda）。

template <typename T> __global__ void kernel(T input) {}
void host_function() {
  kernel<<<1, 1>>>([] __device__() {});            // 正確，extended lambda
  kernel<<<1, 1>>>([] __host__ __device__() {});   // 正確，extended lambda
  // kernel<<<1, 1>>>([](){});                     // 錯誤，host 執行空間的 closure
  // kernel<<<1, 1>>>(global_lambda);              // 錯誤，extended lambda 但在 global scope
}

Extended Lambdas 與其限制

--extended-lambda 旗標允許在 lambda 明確標註執行空間（標註置於 lambda introducer 之後、optional declarator 之前）。指定後 nvcc 定義巨集 __CUDACC_EXTENDED_LAMBDA__。

• extended lambda 必須定義在 __host__ 或 __host__ __device__ 函式的（巢狀）block scope 內。
• extended device lambda：標 __device__；extended host-device lambda：標 __host__ __device__。
• 與一般 lambda 不同，extended lambda 可當 __global__ 函式的 type 引數。

型別 trait

主要限制（精選）

CUDA compiler 在送給 host 編譯器前，會把 extended lambda 換成 namespace scope 的placeholder type，其 template 引數需取得 enclosing function 的位址。由此衍生大量限制：

1. extended lambda 不可定義在另一個 extended lambda 內。
2. 不可定義在 generic lambda（帶 auto 參數）內。
3. 若巢狀在多層 lambda 內，最外層 lambda 必須定義在某（非 lambda）函式的 block scope 內。
4. enclosing function 必須有名稱且位址可取得；若為 class member：所有 enclosing class 須有名稱，且 member function 與 enclosing class 不可為 private/protected。
5. 取 enclosing routine 位址須能無歧義（例如 alias 宣告遮蔽同名 template 型別引數時會出錯）。
6. 不可定義在函式內 local 的 class 中。
7. enclosing function 不可有 deduced return type（auto 回傳）。
8. host-device extended lambda 不可是 generic lambda（不可有 auto/auto... 參數）。
9. enclosing function 為 template 實例時：最多一個 variadic 參數且須在最後、template 參數須具名、實例化引數型別不可為 function-local 或 private/protected member。
10. MSVC host 編譯器：enclosing function 須有 external linkage。
11. MSVC：extended lambda 不可定義在 if constexpr block body 內。
12. 捕捉限制（重要）：
    • 只能 by value 捕捉，不可 by reference。
    • 陣列型別變數維度 > 7 不可捕捉；陣列元素型別須 default-constructible 且 copy-assignable。
    • variadic argument pack 的元素不可捕捉。
    • 捕捉變數型別不可 function-local（extended lambda closure 除外）或 private/protected member。
    • init-capture：host-device extended lambda 不支援；device extended lambda 支援，但 initializer 不可為陣列或 std::initializer_list。
    • extended lambda 的 operator() 非 constexpr，closure type 非 literal type；不可用 constexpr/consteval。
    • 巢狀於 extended lambda 內的 if constexpr block 中，變數不可首次隱式捕捉（除非已在外部隱式捕捉或在明確 capture list）。

void host_function() {
  auto l1 = [x = 1] __device__ () { return x; };           // 正確：device-only 可 init-capture
  // auto l2 = [x = 1] __host__ __device__ () { return x; }; // 錯誤：host-device 不可 init-capture
  int a = 1;
  // auto l3 = [&a] __device__ () { return a; };            // 錯誤：不可 by-reference 捕捉
}

__CUDA_ARCH__ 與 return type 內省限制

Host-Device Lambda 優化注意

• 與 device-only lambda 不同，extended host-device lambda 可從 host code 呼叫。CUDA compiler 會把 host code 中的 extended lambda 換成具名的 placeholder type，而 extended host-device lambda 的 placeholder type 是以**間接函式呼叫（indirect function call）**去呼叫原 lambda 的 operator()。
• 正因這個間接呼叫，host 編譯器對 extended host-device lambda 的優化會比隱式或明確的純 __host__ lambda 差：純 host lambda 的 body 可被輕易 inline 進呼叫點，但遇到 extended host-device lambda 時，host 編譯器難以將原 lambda body inline。
• device-only lambda 不可從 host 呼叫，故沒有這個 host 端 inline 的問題。

*this 捕捉 by-value（C++17）

依 C++11/14 規則，lambda 引用 class member 時捕捉的是 this 指標（by value），而非 member 本身。若 extended device/host-device lambda 在 host 函式定義、於 GPU 執行，且 this 指向 host memory，存取該 member 會造成 runtime error。

C++17 引入 *this 捕捉模式：複製 *this 指向的物件本身而非指標。

struct MyStruct {
  int var;
  __host__ __device__ MyStruct() : var(10) {}
  void run() {
    auto lambda1 = [=, *this] __device__ {   // 複製 *this，GPU 存取 copy_of_star_this->var
      return var + 1;
    };
    foo<<<1, 1>>>(lambda1);                  // kernel launch 成功
    cudaDeviceSynchronize();
  }
};

• CUDA 支援 *this 捕捉用於：__device__/__global__ 函式內的 lambda，以及 host code 內的 extended device-only lambda（需 --extended-lambda）。
• host code 內的非標註 lambda 或 extended host-device lambda，除非語言方言已啟用 *this 捕捉，否則不允許。

Polymorphic Function Wrappers function

<nvfunctional> 提供多型函式包裝器 class template nvstd::function，可儲存、複製、呼叫任意 callable target（如 lambda），host 與 device code 皆可用。

#include <nvfunctional>
__global__ void kernel(int* result) {
  nvstd::function<int()> fn1 = device_function;        // __device__
  nvstd::function<int()> fn2 = host_device_function;   // __host__ __device__
  nvstd::function<int()> fn3 = [](){ return 10; };
  *result = fn1() + fn2() + fn3();
}

nvstd::function 的成員（建構/解構/assignment/swap/operator bool/比較）多為 __device__ __host__，但函式呼叫運算子 operator()(ArgTypes...) 僅 __device__。

考試/測驗重點

Related Notes

• 05-Technical-Appendices/03-Cpp-Language-Features
• 05-Technical-Appendices/05-Cpp-Language-Restrictions
• 05-Technical-Appendices/11-Math-Functions-and-Intrinsics
• 02-Programming-GPUs/01-CUDA-Cpp-Kernels-and-Launch
• 05-Technical-Appendices/Practice-Technical-Appendices
• 00-Dashboard/Exam-Traps