編譯器最佳化提示、診斷與 Warp Matrix (Compiler Hints, Diagnostics and Warp Matrix)

重點總覽

本篇涵蓋三組 device code 擴充：編譯器最佳化提示（給編譯器「假設/期望」資訊以產生更好的碼）、除錯與診斷（assert、__brkpt()、診斷 pragma），以及 Warp Matrix Functions (WMMA)（用 Tensor Core 做 D=A*B+C 的 warp 級矩陣乘加）。提示類函式在 device code 永遠可用，host 端支援與否視 host 編譯器而定；WMMA 需 compute capability 7.0+、整個 warp 協作，且 fragment 是架構特定的不透明結構。

編譯器最佳化提示 (Compiler Optimization Hints)

最佳化提示用額外資訊裝飾程式碼，幫助編譯器產生更好的碼。

• 這些 built-in 函式在 device code 一律可用。
• host code 支援與否取決於 host 編譯器。

#pragma unroll

編譯器預設會展開 trip count 已知的小迴圈；此 pragma 可控制任一迴圈的展開。必須緊接在迴圈正前方，且只作用於該迴圈。後面可選擇性接一個整數常數運算式：

struct MyStruct { static constexpr int value = 4; };
inline constexpr int Count = 4;

__device__ void foo(int* p1, int* p2) {
  // 無引數，迴圈將完全展開
  #pragma unroll
  for (int i = 0; i < 12; ++i) p1[i] += p2[i] * 2;

  #pragma unroll (Count + 1)        // unroll 值 = 5
  for (int i = 0; i < 12; ++i) p1[i] += p2[i] * 4;

  #pragma unroll 1                  // unroll 值 = 1，停用展開
  for (int i = 0; i < 12; ++i) p1[i] += p2[i] * 8;

  #pragma unroll value  // unroll 值 = 4
  for (int i = 0; i < 12; ++i) p1[i] += p2[i] * 16;

  #pragma unroll -1                 // 負值，pragma 被忽略
  for (int i = 0; i < 12; ++i) p1[i] += p2[i] * 2;
}

__builtin_assume_aligned()

讓編譯器假設回傳指標至少對齊到 align bytes。

void* __builtin_assume_aligned(const void* ptr, size_t align);
void* __builtin_assume_aligned(const void* ptr, size_t align, <integral type> offset);

• 三參數版讓編譯器假設 (char*) ptr - offset 至少對齊到 align bytes。
• align 必須是 2 的次方且為整數字面值 (integer literal)。

void* res1 = __builtin_assume_aligned(ptr, 32);     // res1 至少 32-byte 對齊
void* res2 = __builtin_assume_aligned(ptr, 32, 8);  // res2 = (char*)ptr - 8 至少 32-byte 對齊

__builtin_assume() 與 __assume()

void __builtin_assume(bool predicate);
void __assume(bool predicate);   // 僅限 Microsoft Compiler

讓編譯器假設布林引數為真。

__device__ bool is_greater_than_zero(int value) { return value > 0; }
__device__ bool f(int value) {
  __builtin_assume(value > 0);
  return is_greater_than_zero(value);  // 直接回傳 true，不再實際評估條件
}

__builtin_expect()

long __builtin_expect(long input, long expected);

• 告知編譯器 input 預期等於 expected，並回傳 input 的值。
• 通常用於提供分支預測資訊，行為類似 C++20 的 [[likely]] / [[unlikely]] 屬性。

if (__builtin_expect(var, 0))  // 提示編譯器：var 很可能等於 0
  doit();

__builtin_unreachable()

void __builtin_unreachable(void);

告知編譯器控制流永不會到達此呼叫點；若執行期真的到達則為 UB。可避免產生不可達分支的碼，並關閉不可達碼的編譯警告。

switch (in) {
  case 1: return 4;
  case 2: return 10;
  default: __builtin_unreachable();  // 告知 default 永不被觸及
}

自訂 ABI 與 MMA Throughput Pragma

#pragma nv_abi

讓在 separate compilation mode 編譯的程式，藉由保留函式所用暫存器數量，達到接近 whole program compilation 的效能。

#pragma nv_abi preserve_n_data(EXPR) preserve_n_control(EXPR)

• 兩個引數皆選擇性、可任意順序，但至少要提供一個。
• preserve_n_data(EXPR)：限制保留的 data registers 數量。
• preserve_n_control(EXPR)：限制保留的 control registers 數量。

放置位置與效果：

#pragma nv_mma_throughput

針對矩陣乘加 (MMA) 啟用專門調校的編譯器最佳化。

• 僅在 entry-function scope（即 __global__ 函式）上有效。
• 效果可能沿呼叫路徑傳播到被該 entry function 呼叫的 __device__ 函式。

除錯與診斷 (Debugging and Diagnostics)

Assertion：assert()

#define assert(expression) /* unspecified */

當 expression 等於 0 時中止 kernel 執行：

• 若在 debugger 內執行，觸發 breakpoint 供檢視 device 狀態。
• 否則，每個 expression == 0 的執行緒，在透過 cudaDeviceSynchronize()、cudaStreamSynchronize() 或 cudaEventSynchronize() 與 host 同步後，向 stderr 印出訊息：

<filename>:<line number>:<function>:
block: [blockIdx.x,blockIdx.y,blockIdx.z],
thread: [threadIdx.x,threadIdx.y,threadIdx.z]
Assertion `<expression>` failed.

• assert 僅供除錯用途，會影響效能，production code 建議停用。
• 停用方式：在 include assert.h / <cassert> 前定義 NDEBUG 巨集，或用編譯旗標 -DNDEBUG。
• expression 不應有副作用，否則停用 assert 會改變程式功能。
• assert() 在 __device__ 與 __tile__ code 中皆可用。

Breakpoint：__brkpt()

void __brkpt();

從任一 device thread 呼叫即可暫停 kernel 函式的執行。

Diagnostic Pragmas

調整特定診斷訊息被觸發時的嚴重度，用警告訊息中顯示的 error number 指定對象：

#pragma nv_diag_xxx <error_number1>, <error_number2> ...

#pragma nv_diag_suppress 177   // "declared but never referenced"
void foo() { int i = 0; }
#pragma nv_diag_default 177
void bar() { int i = 0; }

也可用 push/pop 儲存與還原當前診斷 pragma 狀態：

#pragma nv_diagnostic push
#pragma nv_diag_suppress 177
void foo() { int i = 0; }
#pragma nv_diagnostic pop
void bar() { int i = 0; }  // 此處會再次發出警告

Warp Matrix Functions (WMMA) 與 Tensor Core

C++ warp matrix 運算利用 Tensor Cores 加速 D=A*B+C 形式的矩陣問題。

• 支援 compute capability 7.0 以上裝置的混合精度浮點資料。
• 需要 warp 內所有執行緒協作。
• 只有在條件對整個 warp 求值結果相同時，才允許用於條件式碼，否則執行極可能 hang。
• 所有相關函式與型別定義於 namespace nvcuda::wmma；sub-byte 運算屬 preview，定義於 nvcuda::wmma::experimental，API/資料結構可能變動。

核心 API

template<typename Use, int m, int n, int k, typename T, typename Layout=void>
class fragment;

void load_matrix_sync(fragment<...>& a, const T* mptr, unsigned ldm);
void load_matrix_sync(fragment<...>& a, const T* mptr, unsigned ldm, layout_t layout);
void store_matrix_sync(T* mptr, const fragment<...>& a, unsigned ldm, layout_t layout);
void fill_fragment(fragment<...>& a, const T& v);
void mma_sync(fragment<...>& d, const fragment<...>& a,
              const fragment<...>& b, const fragment<...>& c, bool satf=false);

fragment 模板參數

第一個模板參數 Use 指定該 fragment 在運算中的角色，決定 tile 維度：

• 乘數的 T 可為 double、float、__half、__nv_bfloat16、char、unsigned char；累加器的 T 可為 double、float、int、__half（僅特定組合受支援）。
• matrix_a / matrix_b 必須指定 Layout：row_major 或 col_major（指該方向元素在記憶體連續）。
• accumulator 的 Layout 應保留預設 void；只有在 load/store 時才另外指定 row/col layout。

對齊、步幅與同步規則

mma_sync 額外規則：satf 與各 fragment 模板參數須全 warp 一致，且 A、B、C、D 的 m、n、k 必須相符。若 satf（saturate to finite value）為 true：

直接存取 fragment 元素

因元素映射未指定，個別元素一般須在 store_matrix_sync 後從記憶體存取。但若全 warp 對所有 fragment 元素一致地做 element-wise 運算，可直接用成員：

enum fragment<Use, m, n, k, T, Layout>::num_elements;
T    fragment<Use, m, n, k, T, Layout>::x[num_elements];

// 將 accumulator tile 整體乘以 0.5
wmma::fragment<wmma::accumulator, 16, 16, 16, float> frag;
float alpha = 0.5f;  // 全 warp 相同值
for (int t = 0; t < frag.num_elements; t++) frag.x[t] *= alpha;

WMMA 資料型別、矩陣尺寸與限制

支援的型別與 m-n-k 尺寸

Alternate Floating Point（CC 8.0+）

• __nv_bfloat16：替代 fp16 格式，與 f32 同範圍但精度降低（7 bits），型別在 cuda_bf16.h。其 fragment 必須搭配 float 累加器；shape 與運算同 __half。
• tf32：Tensor Core 特殊浮點格式，與 f32 同範圍、精度 >=10 bits，內部 layout 為實作定義。輸入矩陣須手動轉成 tf32：用 intrinsic __float_to_tf32（輸入輸出皆 float 型別，但輸出在數值上是 tf32）。需以 precision::tf32 精度 + float 資料型別搭配 load_matrix_sync，兩個累加器皆須 float，唯一支援尺寸 16x16x8。映射為 precision::tf32 -> float。

Double Precision（CC 8.0+）

使用 double 型別 fragment；mma_sync 以 **.rn（round to nearest even）**捨入修飾子執行。

Sub-byte Operations（preview / experimental）

存取 Tensor Core 低精度能力，位於 nvcuda::wmma::experimental，資料型別有 u4（4-bit unsigned）、s4（4-bit signed）、b1（1-bit）。元素被打包，故 num_storage_elements < num_elements：

precision::u4 -> unsigned (8 elements / 1 storage element)
precision::s4 -> int      (8 elements / 1 storage element)
precision::b1 -> unsigned (32 elements / 1 storage element)
T -> T  // 其餘型別

• sub-byte 的 layout 固定：matrix_a 須 row_major、matrix_b 須 col_major。
• ldm：u4/s4 須為 32 的倍數、b1 須為 128 的倍數（皆等於 16 bytes 倍數）。
• bmma_sync 做位元矩陣乘加 D = (A op B) + C：op 為 bmmaBitOp（bmmaBitOpXOR 需 compute_75+；bmmaBitOpAND 需 CC 8.0+）後接累加 bmmaAccumulateOpPOPC（計算 set bit 數）。

Restrictions（fragment 是架構特定的）

• tensor core 所需特殊格式會隨 major/minor 架構不同；fragment 是不透明、架構特定的 ABI 結構，開發者不可假設參數如何映射到暫存器。
• 危險：若兩函式為不同 link-compatible 架構編譯（如 sm_70 與 sm_75，其 fragment layout 不同）卻連結在一起，互傳 fragment 會導致錯誤結果或損毀，且編譯期與執行期工具都可能偵測不到。

完整範例（16x16x16 單 warp 矩陣乘法）

#include <mma.h>
using namespace nvcuda;

__global__ void wmma_ker(half* a, half* b, float* c) {
  // 宣告 fragment
  wmma::fragment<wmma::matrix_a, 16, 16, 16, half, wmma::col_major> a_frag;
  wmma::fragment<wmma::matrix_b, 16, 16, 16, half, wmma::row_major> b_frag;
  wmma::fragment<wmma::accumulator, 16, 16, 16, float> c_frag;

  wmma::fill_fragment(c_frag, 0.0f);     // 輸出初始化為零
  wmma::load_matrix_sync(a_frag, a, 16); // 載入輸入
  wmma::load_matrix_sync(b_frag, b, 16);
  wmma::mma_sync(c_frag, a_frag, b_frag, c_frag);          // 矩陣乘加
  wmma::store_matrix_sync(c, c_frag, 16, wmma::mem_row_major); // 存回輸出
}

考試/測驗重點

Related Notes

• 05-Technical-Appendices/08-Warp-Functions-and-Macros
• 05-Technical-Appendices/06-Function-and-Variable-Annotations
• 02-Programming-GPUs/08-SIMT-Memory-Performance
• 02-Programming-GPUs/17-NVCC-Compiler
• 05-Technical-Appendices/Practice-Technical-Appendices
• 00-Dashboard/Exam-Traps