函式與變數限定詞 (Function and Variable Annotations)

重點總覽

CUDA C/C++ 在標準 C++ 之上加了一組限定詞 (annotation)，用來標示函式在哪個情境執行、變數放在哪段記憶體、以及如何啟動 kernel。核心分三類：執行空間限定詞（__global__ / __device__ / __host__）、記憶體空間限定詞（__device__ / __constant__ / __shared__ / __managed__），以及啟動相關限定詞（__launch_bounds__ / __maxnreg__ / __cluster_dims__）。再配合內建型別（向量型別、dim3）與內建變數（threadIdx 等）構成 kernel 的基本語彙。

Function Execution Space Specifiers

執行空間限定詞標示函式在 host、SIMT（傳統 device）、或 tile 情境執行與被呼叫：

__global__ / __tile_global__ 的呼叫是非同步的：在 device 完成執行前就先返回 host thread。

多重執行空間與 __CUDA_ARCH__

宣告多個執行空間（如 __host__ __device__）的函式會為每個情境各編譯一份。用 __CUDA_ARCH__ 巨集區分 host / device 程式路徑：

__host__ __device__ void func() {
#if defined(__CUDA_ARCH__)
  // Device 端程式路徑
#else
  // Host 端程式路徑
#endif
}

Variable Memory Space Specifiers

記憶體空間限定詞標示變數在 device 上的儲存位置：

• __device__、__tile__、__constant__ 變數可從 host 用 cudaGetSymbolAddress()、cudaGetSymbolSize()、cudaMemcpyToSymbol()、cudaMemcpyFromSymbol() 存取。
• __constant__ 在 device 端唯讀，只能由 host 透過 Runtime API 修改。

__device__   float device_var   = 4.0f;  // 位於 device 記憶體
__constant__ float constant_var = 4.0f;  // 位於 constant 記憶體
int main() {
  float* device_ptr;
  cudaGetSymbolAddress((void**)&device_ptr, device_var); // 取得符號位址
  size_t symbol_size;
  cudaGetSymbolSize(&symbol_size, device_var);            // 取得大小 (4 bytes)
  float host_var;
  cudaMemcpyFromSymbol(&host_var, device_var, sizeof(host_var)); // device -> host
  host_var = 3.0f;
  cudaMemcpyToSymbol(device_var, &host_var, sizeof(host_var));   // host -> device
}

__shared__ 記憶體

• 尺寸分兩種：static（編譯期決定）或 dynamic（kernel 啟動時決定）。
• 動態尺寸變數必須宣告為 external array 或 pointer。
• 靜態尺寸變數不可在宣告處初始化。

extern __shared__ char dynamic_smem[];      // 動態 shared memory
__global__ void kernel() {
  __shared__ int smem_var1[4];              // static 尺寸
  auto smem_var2 = (int*)dynamic_smem;      // dynamic 尺寸
}
int main() {
  size_t shared_memory_size = 16;
  kernel<<<1, 1, shared_memory_size>>>();   // 第三個參數指定動態 shared bytes
  cudaDeviceSynchronize();
}

__managed__ 記憶體

__tile__ 記憶體

• 與 __device__ 類似：配置於 device 全域記憶體、可由 Runtime API 存取；差別是 __tile__ 變數可由 tile code 直接以名稱存取。
• 同一變數不可同時標記 __device__ 與 __tile__。
• __device__ 變數不可在 __tile__ / __tile_global__ code 以名稱直接存取；__tile__ 變數也不可在 __device__ / __global__ code 以名稱直接存取。但兩者都位於 device global memory，可透過指標在任一情境存取。
• __tile__ 變數不可含 pointer 或 reference 型別的子物件。

Inlining、__restrict__ 與 __grid_constant__

Inlining Specifiers（限 __host__ / __device__）

__restrict__ 指標

• 解決 pointer aliasing（多個指標指向重疊記憶體）抑制優化（如重排、共同子運算式消除）的問題。
• 是程式設計者的承諾：指標生命週期內，該記憶體只透過此指標存取，讓編譯器更積極優化。
• 所有指標參數都必須加 __restrict__ 優化才有效。

__device__ void f(const float* __restrict__ a,
                  const float* __restrict__ b,
                  float* __restrict__ c);

__grid_constant__ 參數

替 __global__ 函式參數加上 __grid_constant__，可阻止編譯器建立per-thread copy，改讓 grid 內所有 thread 透過單一位址存取，提升效能。

• 具 kernel 的生命週期；對單一 kernel 私有，其他 grid（含 sub-grid）看不到。
• 所有 thread 看到相同位址。
• 唯讀：修改該物件或其子物件（含 mutable 成員）為 undefined behavior。
• 參數須為 const-qualified 非 reference 型別；在 __tile_global__ 參數上會被忽略。
• 所有函式宣告（含 function template 特化與實例化）對 __grid_constant__ 參數的標註必須與主宣告一致，否則編譯錯誤。

Annotation Summary（Table 41）

Built-in Types and Variables

Host Compiler 型別擴充

• 128-bit 整數 __int128：Linux 上 host 編譯器定義 __SIZEOF_INT128__ 時支援。
• 128-bit 浮點 __float128 / _Float128：在 compute capability 10.0 以上的 GPU 可用；Linux x86 上 host 編譯器定義 __SIZEOF_FLOAT128__ 或 __FLOAT128__ 時支援。__float128 常數運算式可能以較低精度處理。
• _Complex 型別只在 host 程式支援。
• 128-bit 整數與浮點型別不支援 tile code。

Built-in Variables（device-only）

• dim3 與 uint3 都是含 x/y/z 三個 unsigned 的 trivial struct；C++11 起 dim3 各分量預設為 1。
• 上述變數不支援 tile code；tile 中改用 cuda::tiles::bid() 與 cuda::tiles::num_blocks()。

Built-in Vector Types

CUDA 提供從基本整數/浮點衍生的向量型別（host 與 device 皆可用），命名規則為 <base>1～<base>4，例如 char1..char4、int1..int4、float1..float4、double1..double4 等。

• 分量透過 .x / .y / .z / .w 存取（第一到第四個）。
• 每種型別都有工廠函式 make_<type_name>()，例如 make_int4(...)。
• 非 nvcc 編譯 host 程式時，需 #include <cuda_runtime.h> 才能用這些型別與函式。

int sum(int4 v) { return v.x + v.y + v.z + v.w; }
int4 add_one(int x, int y, int z, int w) {
  return make_int4(x + 1, y + 1, z + 1, w + 1);
}

對齊重點（Table 43 摘要）：

Kernel Configuration

每次呼叫 __global__ / __tile_global__ 函式都必須提供 execution configuration，語法為在函式名與引數列之間插入 <<<grid_dim, block_dim, dynamic_smem_bytes, stream>>>：

__global__ void kernel(float* parameter);
kernel<<<grid_dim, block_dim, dynamic_smem_bytes>>>(parameter);

• execution configuration 的引數會先於實際函式引數被求值。
• 若 grid_dim / block_dim 超過 device 上限，或 dynamic_smem_bytes 超過扣除靜態配置後可用的 shared memory，呼叫會失敗。

Thread Block Cluster（__cluster_dims__）

• compute capability 9.0 以上可指定編譯期 cluster 維度，語法 __cluster_dims__([x, [y, [z]]])。

__global__ void __cluster_dims__(2, 1, 1) kernel(float* parameter); // X=2, Y=Z=1

• 無引數的 __cluster_dims__() 表示 kernel 以 grid cluster 啟動，cluster 維度延到啟動時指定；啟動時若沒指定會產生 launch-time error。
• 不可用於 __tile_global__ kernel。
• 也可在執行期用 cudaLaunchKernelEx（搭配 cudaLaunchConfig_t 與 cudaLaunchAttributeClusterDimension）指定；grid 維度應為 cluster 大小的整數倍。

Launch Bounds（__launch_bounds__）

__global__ void
__launch_bounds__(maxThreadsPerBlock, minBlocksPerMultiprocessor, maxBlocksPerCluster)
MyKernel(...) { ... }

• 編譯器據此推導暫存器上限 L，確保 minBlocksPerMultiprocessor 個 block（各 maxThreadsPerBlock threads）能常駐 SM；初始暫存器數超過 L 就壓低（可能增加 local memory / 指令數）。
• 若初始暫存器數低於 L：只指定 maxThreadsPerBlock（未指定 minBlocksPerMultiprocessor）時，編譯器用它決定 n↔n+1 個常駐 block 轉換的暫存器門檻；若同時指定 minBlocksPerMultiprocessor 與 maxThreadsPerBlock，編譯器可能反而把用量提高到 L，以減少指令數、改善單執行緒指令的延遲隱藏。
• kernel 若以超過 maxThreadsPerBlock 的 threads/block 或超過 maxBlocksPerCluster 的 blocks/cluster 啟動，會啟動失敗。

Maximum Registers per Thread（__maxnreg__）

__global__ void __maxnreg__(maxNumberRegistersPerThread) MyKernel(...) { ... }

• 指定 block 內單一 thread 可配置的最大暫存器數；編成 .maxnreg PTX directive。
• __launch_bounds__() 與 __maxnreg__() 不可同時套在同一 kernel。
• --maxrregcount <N> 可控制整個檔案所有 __global__ 函式的暫存器用量，但對含 __maxnreg__ 的 kernel 會被忽略。

考試/測驗重點

Related Notes

• 05-Technical-Appendices/05-Cpp-Language-Restrictions
• 05-Technical-Appendices/07-Synchronization-and-Atomic-Functions
• 02-Programming-GPUs/04-CUDA-Cpp-Errors-and-Specifiers
• 01-Introduction-to-CUDA/04-GPU-Memory-Hierarchy
• 05-Technical-Appendices/01-Compute-Capabilities
• 05-Technical-Appendices/09-Compiler-Hints-and-Warp-Matrix
• 05-Technical-Appendices/Practice-Technical-Appendices
• 00-Dashboard/Exam-Traps