C/C++ 語言限制 (C/C++ Language Restrictions)

重點總覽

CUDA C++ 在 device 端不是完整的 ISO C++：RTTI、例外、long double、recursion（__global__/__tile_global__/__tile__）皆不支援，且跨 execution space 的指標解參考、取址、變數宣告都有嚴格限制。除了通用限制外，每個 C++ 標準版本（C++11/14/17/20/23）還疊加各自的特殊規則。核心心法是：host 與 device 是兩個獨立執行空間，凡是「跨空間引用記憶體/位址/型別」幾乎都是 undefined behavior（UB）。

通用不支援功能與 namespace 保留

Unsupported Features

Namespace Reservations

• 在 top-level namespace cuda::、nv::、cooperative_groups::（及其任何巢狀 namespace）內新增定義皆為 UB。
• 唯一允許：cuda 當作子 namespace 巢狀於非保留 namespace 內。

namespace cuda {
    struct foo;        // ERROR：在 cuda namespace 內宣告 class
    void bar();        // ERROR：在 cuda namespace 內宣告 function
    namespace utils {} // ERROR：在 cuda namespace 內宣告 namespace
}
namespace utils {
    namespace cuda {   // CORRECT：cuda 巢狀於非保留 namespace utils 內
        void bar();
    }
}
using namespace utils; // ERROR：等同把 symbol 加進 global scope 的 cuda namespace

指標、記憶體位址與變數限制

Pointers and Memory Addresses

指標解參考（*p、p->m、p[0]）只允許在記憶體所屬的 execution space 內，否則為 UB（多半 segfault）。

Variables：memory space specifier 規則

const-qualified 變數

• 無 memory space 標註的 const 變數（global/namespace/class scope）視為 host 變數；device code 不可引用或取址。
• 但若同時滿足下列條件，可在 device code 直接使用其值：
  • 在使用點之前已用 constant expression 初始化
  • 型別非 volatile
  • 型別為 built-in 整數型，或 built-in 浮點型（host compiler 為 MSVC 時浮點型例外）

volatile-qualified 變數

• volatile 僅為了與 ISO C++ 相容而保留，在 GPU 上幾乎沒有適用場景。
• 讀寫 volatile 非 atomic，編譯為 volatile 指令，不保證記憶體操作順序、也不保證硬體實際存取次數等於 PTX 指令數。
• tile code 中 volatile 對記憶體存取毫無效果。

static 變數

• device function 內允許 static 區域變數。
• 每個 execution space 各持有一份獨立的 static 變數：
  • __host__ __device__ function：host 一份、device 一份
  • __tile__ __device__ function：tile 一份、SIMT 一份
• 若 function 帶 __host__ execution space，帶顯式 memory space 的 static（static __device__/__tile__/__constant__/__shared__/__managed__）僅在 __CUDA_ARCH__ 有定義時才允許。
• 不允許 dynamic initialization（如 static int v = x;、static TrivialStruct s{x};、non-trivial constructor 皆 ERROR）。

extern 變數

• 在 whole program compilation 模式下，__device__/__tile__/__shared__/__managed__/__constant__ 變數不可用 extern 定義 external linkage。
• __tile__ 變數此限制在 separate compilation 模式下也適用。
• 唯一例外：動態配置的 __shared__ 變數（dynamic shared memory）。

函式限制

__global__ Function Parameters

• 不可 variadic：C ellipsis ... 與 va_list 皆禁止（C++11 variadic template 則允許，受 __global__ Variadic Template 限制）。
• 參數經 constant memory 傳給 device，總大小上限 32,764 bytes。
• 參數不可為 std::initializer_list 型別。
• Polymorphic（virtual）class 參數為 UB。
• Lambda/closure 型別允許（受相應限制）。
• __tile_global__ 函式：參數不可為 pass-by-value 的 class/struct/union。

__global__ Function Arguments Passing

類別與模板限制

Class-type Variables / Function Members

• 帶 __device__/__tile__/__constant__/__managed__/__shared__ 的變數，其 class 型別不可有 non-empty constructor 或 non-empty destructor（empty 的判定：trivial，或已定義、無參數、空 body、無 virtual function/base、無 NSDMI 等）。
• __global__/__tile_global__ 函式不能是 struct/class/union 的 member；可出現在 friend 宣告，但不能在 friend 中定義。

Implicitly-Declared / Explicitly-Defaulted Functions

• 非 virtual 且 implicitly-declared 或 explicitly-defaulted 的函式 F，其 execution space = 所有呼叫它的函式之 execution space 的聯集；分析時 __global__ caller 視為 __device__ caller。
• 若 F 是 implicitly-declared virtual function（如 virtual destructor），則被它覆寫且非 implicitly-declared 的 virtual function 的 execution space 也會被併入 F。

Polymorphic Classes 與 Windows Class Layout

Templates

下列型別不可當作 __global__ 函式或 __device__/__constant__ 變數（C++14）的 template argument：

• 定義於 __host__ 或 __host__ __device__ function scope 內的型別
• 匿名型別（anonymous struct、lambda），除非該型別 local 於 __device__/__global__ function
• private/protected 的 class member 型別，除非該 class local 於 __device__/__global__ function
• 由上述任一型別複合而成的型別

Tile code 限制（__tile__ / __tile_global__）

__tile__ 與 __tile_global__ 函式有額外限制。下列語言構造在 tile code 不支援：

其他限制：

• __tile__/__tile_global__ 函式必須在宣告它的同一 TU 內有 function body。
• virtual function 不可標註 __tile__。
• 不可 variadic（C ellipsis ...）。
• 不可直接或間接 recursion。
• __tile_global__ 參數不可為 pass-by-value class/struct/union。
• tile code 不可做 device-side kernel launch，tile kernel 也不可從 device 端啟動。
• tile code 不可直接存取 __half、__nv_bfloat16 及相關 extended FP 型別的 __x member。

C++11/14/17/20/23 各標準版本限制

C++11 Restrictions

可直接用於 device code 的 constexpr 變數型別：

• C++ scalar 型別（排除 pointer 與 pointer-to-member）：nullptr_t、bool、整數型、浮點型
• enumerator：enum、enum class
• class 型別（class/struct/union）且具 constexpr constructor
• 上述型別的 raw array（如 int[]），僅當用於 constexpr __device__ 或 __host__ __device__ 函式內
• constexpr __managed__ 與 constexpr __shared__ 不允許

C++14 Restrictions

C++17 Restrictions

C++20 Restrictions

C++23 Restrictions

• 無已知 C++23 特有限制，超出 C++23 Language Features 表中標記為不支援/N/A（含 Defect Report 解析）的部分皆屬「缺漏/不適用」而非額外行為限制。
• Equality Operator (P2468R2)：NVCC 未完整實作，但以近似行為處理，目前未導致已知的 user code 失敗。

考試/測驗重點

Related Notes

• 05-Technical-Appendices/03-Cpp-Language-Features
• 05-Technical-Appendices/04-Cpp-Standard-Library-and-Lambdas
• 05-Technical-Appendices/06-Function-and-Variable-Annotations
• 02-Programming-GPUs/04-CUDA-Cpp-Errors-and-Specifiers
• 05-Technical-Appendices/Practice-Technical-Appendices
• 00-Dashboard/Exam-Traps