Warp 函式、CUDA 巨集與裝置函式 (Warp Functions, CUDA Macros and Device Functions)

重點總覽

Warp 函式讓「同一 warp 內的 thread」彼此通訊與協作。其中 vote / match / reduce / shuffle 等同步原語皆以 *_sync 結尾並帶 mask 參數（連同 __syncwarp 統稱 warp __sync intrinsic），mask 用來在硬體執行前確保收斂；__activemask() 為例外，不帶 mask、也非同步原語。本篇另含 CUDA 專屬巨集（__CUDA_ARCH__ 等）與裝置端工具函式（位址空間判斷/轉換、低階快取載入/儲存、__trap、__nanosleep、DPX）。

Warp Active Mask：__activemask()

unsigned __activemask();

回傳 32-bit 整數遮罩，第 N bit=1 表示呼叫當下第 N lane 為 active；inactive（含已退出程式）的 lane 為 0。

if (pred) {
  // Invalid：at_least_one 的值不確定
  at_least_one = __activemask() > 0;
}

unsigned mask = __activemask();
// 假設 mask == 0xFFFFFFFF（全部 bit 設立、所有 thread active）
int predicate = threadIdx.x % 2 == 0;        // 偶數 lane 為 1、奇數為 0
int result = __any_sync(mask, predicate);    // active 集合可能已不被保留

Warp Vote、Match 與 Reduce

Warp Vote（reduction-and-broadcast）

每個非退出 thread 輸入一個整數 predicate，與 0 比較後做歸約並把單一結果廣播回所有參與 thread。

Warp Match（broadcast-and-compare）

unsigned __match_any_sync(unsigned mask, T value);
unsigned __match_all_sync(unsigned mask, T value, int* pred);

• __match_any_sync：回傳 mask 內「與自己 bitwise 值相同」的非退出 thread 之遮罩。
• __match_all_sync：若 mask 內所有非退出 thread 的值皆相同 → 回 mask 並把 *pred 設 true；否則回 0、*pred 設 false。
• T 可為 int、unsigned、long、unsigned long、long long、unsigned long long、float、double。

Warp Reduce（CC 8.x 以上）

T __reduce_add_sync(unsigned mask, T value);   // T: 有/無號整數
T __reduce_min_sync(unsigned mask, T value);
T __reduce_max_sync(unsigned mask, T value);
unsigned __reduce_and_sync(unsigned mask, unsigned value);
unsigned __reduce_or_sync (unsigned mask, unsigned value);
unsigned __reduce_xor_sync(unsigned mask, unsigned value);

• add/min/max：對 mask 內各 thread 的 value 做算術歸約，T 為有號或無號整數。
• and/or/xor：對 value 做 bitwise 歸約。

Warp Shuffle 函式

T __shfl_sync     (unsigned mask, T value, int srcLane,  int width = warpSize);
T __shfl_up_sync  (unsigned mask, T value, unsigned delta, int width = warpSize);
T __shfl_down_sync(unsigned mask, T value, unsigned delta, int width = warpSize);
T __shfl_xor_sync (unsigned mask, T value, int laneMask, int width = warpSize);

不經 shared memory，在 warp 內非退出 thread 間交換 value。

T 可為：

• int、unsigned、long、unsigned long、long long、unsigned long long、float、double。
• 含 cuda_fp16.h 時可用 __half、__half2。
• 含 cuda_bf16.h 時可用 __nv_bfloat16、__nv_bfloat162。

典型用法範例

// 廣播：全 warp 從 lane 0 取 value
value = __shfl_sync(0xFFFFFFFF, value, 0);

// 蝴蝶式 full-warp 歸約：log2(32) = 5 步
for (int i = 1; i <= 16; i *= 2)
  value += __shfl_xor_sync(0xFFFFFFFF, value, i);

// 8-thread sub-partition 的 inclusive scan：log2(8) = 3 步
for (int delta = 1; delta <= 4; delta *= 2) {
  int tmp = __shfl_up_sync(0xFFFFFFFF, value, delta, /*width=*/8);
  if ((laneId % 8 - delta) >= 0) value += tmp;  // source_lane < 0 的不變
}

Warp __sync Intrinsic Constraints（mask 規則）

所有 __shfl_*_sync、__match_*_sync、__reduce_*_sync、__syncwarp 都靠 mask 指明哪些 lane 參與，並在硬體執行前確保收斂。mask 每個 bit 對應一個 lane ID（threadIdx.x % warpSize），intrinsic 會等到 mask 內所有非退出 thread 都抵達呼叫點。

正確執行需滿足：

• 每個呼叫 thread，其對應 bit 在 mask 中必須=1。
• 每個非呼叫 thread，其對應 bit 必須=0（已退出 thread 被忽略）。
• mask 內所有非退出 thread 必須以相同 mask 值執行該 intrinsic。
• 不同 thread 可用不同 mask 並發呼叫，只要這些 mask 彼此 disjoint（不重疊）——即使在 divergent 控制流中亦合法。

正確：if (threadIdx.x < 4)   __all_sync(0b1111,      pred);   // 0,1,2,3 都參與
正確：disjoint               __syncwarp(tid<16 ? 0xFFFF : 0xFFFF0000);
錯誤：active 卻沒設           __all_sync(0b0000011,   pred);   // 2,3 active 但漏設
錯誤：設了卻非 active         __all_sync(0b1111111,   pred);   // 4,5,6 非 active
錯誤：參與者 mask 重疊不一致  __all_sync(tid==0 ? 1 : 0xFFFFFFFF, pred);

CUDA-Specific Macros

__CUDA_ARCH__

代表「正在編譯的目標虛擬架構」，其值可能與實際裝置 compute capability 不同。只在 device code（__device__、__host__ __device__、__global__）中有定義，可用來分流架構特化路徑、或區分 host/device 程式碼。

• 數值關係：__CUDA_ARCH__ = <version> * 10。例如 nvcc --generate-code arch=compute_80,code=sm_90 → __CUDA_ARCH__ 為 800。

__CUDA_ARCH_SPECIFIC__ 與 __CUDA_ARCH_FAMILY_SPECIFIC__

用來識別具「架構特定」與「家族特定」特性的 GPU，同樣只在 device code 定義，分別對應 nvcc 的 compute_<version>a 與 compute_<version>f。

CUDA Feature Testing Macros

當 CUDA 前端編譯器支援某特性時定義對應巨集：

__nv_pure__ 屬性

標記 pure 函式（對參數無 side effect、可讀但不改全域變數），host/device 皆可用。編譯器將其轉成 GNU pure 屬性或 MSVC noalias 屬性。

__device__ __nv_pure__ int add(int a, int b) { return a + b; }

CUDA-Specific Functions

位址空間判斷（Address Space Predicate）

__device__ unsigned __isGlobal(const void* ptr);
__device__ unsigned __isShared(const void* ptr);
__device__ unsigned __isConstant(const void* ptr);
__device__ unsigned __isGridConstant(const void* ptr);  // __grid_constant__ kernel 參數
__device__ unsigned __isLocal(const void* ptr);

若 ptr 為指向該位址空間物件的 generic 位址則回 1，否則回 0；對 NULL 指標行為未指定。

位址空間轉換（Address Space Conversion）

當編譯器無法判定指標的位址空間（跨 translation unit、或與 PTX 互動）時使用：

低階載入/儲存（快取運算子）

皆支援所有 C++ 基本型別、CUDA 向量型別（x3 分量除外）與擴充浮點型別（__half、__half2、__nv_bfloat16、__nv_bfloat162）。

__trap()、__nanosleep()

void __trap();                                 // 中止 kernel、向 host 觸發中斷
__device__ void __nanosleep(unsigned nanoseconds);  // 約睡 ns 奈秒，上限約 1 ms

// __nanosleep 經典用例：exponential back-off mutex
__device__ void mutex_lock(unsigned* mutex) {
  unsigned ns = 8;
  while (atomicCAS(mutex, 0, 1) == 1) {
    __nanosleep(ns);
    if (ns < 256) ns *= 2;     // 退避時間倍增，上限 256ns
  }
}

DPX（Dynamic Programming eXtension）指令

對最多三個 16-bit 或 32-bit 有/無號整數做 min/max 與 fused add-min/max，並可選 ReLU（夾到 0）。

__vimax3_s32_relu(-15, 8, 5);      // max(-15, 8, 5, 0) = 8
__vimax3_s32_relu(-15, -2, -4);    // 全負 → ReLU 後 = 0
__viaddmax_s32_relu(-5, 6, -2);    // max(-5+6, -2, 0) = max(1,-2,0) = 1
__vimax3_u16x2(0x00050002, 0x00070004, 0x00020006); // 逐 16-bit → 0x00070006

考試/測驗重點

Related Notes

• 05-Technical-Appendices/07-Synchronization-and-Atomic-Functions
• 05-Technical-Appendices/09-Compiler-Hints-and-Warp-Matrix
• 01-Introduction-to-CUDA/02-Execution-Model-and-SIMT
• 03-Advanced-CUDA/04-Using-PTX-and-Hardware-Model
• 05-Technical-Appendices/11-Math-Functions-and-Intrinsics
• 05-Technical-Appendices/Practice-Technical-Appendices
• 00-Dashboard/Exam-Traps