NVCC：NVIDIA CUDA 編譯器 (The NVIDIA CUDA Compiler)

重點總覽

nvcc 是什麼

nvcc（NVIDIA CUDA Compiler）是 NVIDIA 提供、用來編譯 CUDA C/C++ 與 PTX 程式碼的 toolchain，屬於 CUDA Toolkit 的一部分。它本身是一個頂層工具，會協調整個編譯流程，依各階段呼叫合適的子工具（compiler、linker、PTX assembler、Cubin assembler）。

• nvcc 驅動的是 offline 編譯（編譯時就產生目標碼）。
• 相對地，CUDA runtime compiler nvrtc 驅動的是 online / Just-in-Time (JIT) 編譯（執行期才編譯）。

2.7.1 CUDA Source Files and Headers

用 nvcc 編譯的原始檔可同時含 host code（在 CPU 執行）與 device code（在 GPU 執行）。nvcc 依副檔名判斷內容類型。

• 含 device code 的原始檔慣例用 .cu；含 device code 的 header 慣例用 .cuh，以與 host-only header（.h 等）區分。

2.7.2 NVCC Compilation Workflow

初始階段，nvcc 會把 device code 與 host code 分離，分別交給 GPU compiler 與 host compiler。

                 example.cu
                     │ nvcc 分流
       ┌─────────────┴──────────────┐
   host code                    device code
       │                            │
  host compiler              GPU compiler (每個 virtual ISA 各跑一次)
 (gcc/clang/...)                    │  例如 compute_90
       │                          PTX assembly
       │                            │ ptxas
   host object                    Cubin (對應 SM 版本的 hardware ISA)
       │                            │
       └──────────► link ◄──────────┘
                     │
              可內嵌多個 PTX + Cubin
                的 Fatbin 容器
                     ▼
                  executable

關鍵事實：

• nvcc 編譯 host code 需要相容的 host compiler；CUDA Toolkit 對 Linux/Windows 各定義了 host compiler 支援政策。
• 只含 host code 的檔案，可用 nvcc 或 直接用 host compiler 編譯；產生的 object file 可在 link 階段與含 GPU code 的 nvcc object file 合併。
• GPU compiler 把 C/C++ device code 編成 PTX assembly，並為命令列上每個指定的 virtual ISA（如 compute_90）各跑一次。
• 個別 PTX 再交給 ptxas，產生對應 target hardware ISA 的 Cubin；hardware ISA 以其 SM version 識別。
• 可把多個 PTX 與 Cubin target 內嵌進單一 Fatbin 容器，讓一個 binary 同時支援多個 virtual 與 target hardware ISA。

範例 example.cu：

// ----- example.cu -----
#include <stdio.h>
__global__ void kernel() {
    printf("Hello from kernel\n");
}
void kernel_launcher() {
    kernel<<<1, 1>>>();
    cudaDeviceSynchronize();
}
int main() {
    kernel_launcher();
    return 0;
}

此檔同時含 device（__global__ kernel）與 host（main/kernel_launcher）程式碼，正是 nvcc 分流編譯的典型輸入。

2.7.3 NVCC Basic Usage

基本編譯指令：

nvcc <source_file>.cu -o <output_file>

nvcc 接受常見 compiler flag：

nvcc example.cu -I path_to_include/ -L path_to_library/ -lcublas -o <output_file>

2.7.3.1 PTX 與 Cubin 生成

預設下，nvcc 會為 CUDA Toolkit 支援的最早（最低 compute_XY / sm_XY）GPU 架構產生 PTX 與 Cubin，以最大化相容性。

• -arch：為特定 GPU 架構產生 PTX 與 Cubin。
• -gencode：為多個 GPU 架構產生 PTX 與 Cubin。
• --list-gpu-code：列出所有支援的 real GPU 架構；--list-gpu-arch：列出所有支援的 virtual GPU 架構。

nvcc example.cu -arch=compute_<XY>  # 例 compute_80：PTX-only，GPU 向前相容
nvcc example.cu -arch=sm_<XY>       # 例 sm_80：PTX + Cubin，GPU 向前相容
nvcc example.cu -arch=native        # 自動偵測並為當前 GPU 產生 Cubin（無 PTX，無向前相容）
nvcc example.cu -arch=all           # 為所有支援架構產生 Cubin，另含最新 PTX 供向前相容
nvcc example.cu -arch=all-major     # 為所有主要架構產生 Cubin（如 sm_80, sm_90），另含最新 PTX

進階：可個別指定 PTX 與 Cubin target。

# 為 virtual arch compute_80 產生 PTX，並編成 sm_86 的 Cubin，且保留 compute_80 PTX
nvcc example.cu -arch=compute_80 -gpu-code=sm_86,compute_80   # PTX + Cubin
# 編成 sm_86、sm_89 的 Cubin（不保留 PTX）
nvcc example.cu -arch=compute_80 -gpu-code=sm_86,sm_89
nvcc example.cu -gencode=arch=compute_80,code=sm_86,sm_89     # 同上
# 多個 gencode：compute_80→sm_86,sm_89 與 compute_90→sm_90
nvcc example.cu -gencode=arch=compute_80,code=sm_86,sm_89 \
                -gencode=arch=compute_90,code=sm_90

2.7.3.2 Host Code Compilation Notes

不含 device code 或 symbol 的 compilation unit（原始檔 + 其 header）可直接用 host compiler 編譯。

• 若任一 compilation unit 使用 CUDA runtime API，最終必須連結 CUDA runtime library。
• runtime 同時提供靜態與動態版：libcudart_static（靜態）與 libcudart（動態）。
• 預設連結靜態 runtime；要改用動態版，在 compile 或 link 指令加 --cudart=shared。
• -ccbin <compiler>：指定 host function 使用的 host compiler；環境變數 NVCC_CCBIN 亦可。
• -Xcompiler：把參數透傳給 host compiler。

nvcc example.cu -ccbin=clang++
export NVCC_CCBIN='gcc'
nvcc example.cu -Xcompiler=-O3   # -O3 被透傳給 host compiler

2.7.3.3 Separate Compilation of GPU Code

nvcc 預設為 whole-program compilation：要求用到的所有 GPU code 與 symbol 都在同一 compilation unit 內。

要讓 device function 呼叫其他 compilation unit 定義的 device function 或存取其 device variable，必須加 -rdc=true（或別名 -dc），啟用跨單元的 device code linking，這稱為 separate compilation。

separate compilation 的條件：

• 在某單元定義的 non-const __device__ 變數，於其他單元須以 extern 引用。
• 所有 const __device__ 變數都必須以 extern 定義並引用。
• 所有 .cu 原始檔都須用 -dc 或 -rdc=true 編譯。
• host 與 device function 預設即為 external linkage，不需 extern。

// ----- definition.cu -----
extern __device__ int device_variable = 5;
__device__
int device_function() { return 10; }

// ----- example.cu -----
extern __device__ int device_variable;   // 以 extern 引用
__device__
int device_function();                    // 宣告，定義在另一單元
__global__ void kernel(int* ptr) {
    device_variable = 0;
    *ptr = device_function();
}

nvcc -dc definition.cu -o definition.o
nvcc -dc example.cu    -o example.o
nvcc definition.o example.o -o program   # device link + 最終連結

每個 .cu 都以 -dc 各自編成 object，最後再由 nvcc 連結；這就是 separate compilation 的標準流程。

2.7.4 Common Compiler Options

2.7.4.1 Language Features

nvcc 支援從 C++03 到 C++23 的核心語言特性。

• -std={c++03|c++11|c++14|c++17|c++20|c++23}：指定語言標準。
• -restrict：斷言所有 kernel pointer 參數皆為 restrict pointer。
• -extended-lambda：允許在 lambda 宣告中加 __host__、__device__ 標註。
• -expt-relaxed-constexpr（實驗性）：允許 host code 呼叫 __device__ constexpr 函式、device code 呼叫 __host__ constexpr 函式。

2.7.4.2 Debugging Options

2.7.4.3 Optimization Options

• -Xptxas：把參數傳給 PTX assembler ptxas。例如 -Xptxas=-maxrregcount=N 設定每 thread 最大 register 數。
• -extra-device-vectorization：啟用更積極的 device code 向量化。
• --apply-controls=/path/to/file：傳入 advanced controls file (ACF) 給 nvcc 與 ptxas，針對特定 workload 改變預設編譯行為。

可從編譯器取得、利於進階最佳化的輸出旗標：

2.7.4.4 Link-Time Optimization (LTO)

separate compilation 因跨檔最佳化機會有限，效能可能低於 whole-program。LTO 在 link 時跨各別編譯的檔案做最佳化，代價是編譯時間增加，但可挽回大部分 whole-program 的效能，同時保留 separate compilation 的彈性。

啟用 LTO 需 -dlto 旗標或 lto_<SM version> link-time 最佳化 target：

# 方式一：-dlto
nvcc -dc -dlto -arch=sm_100 definition.cu -o definition.o
nvcc -dc -dlto -arch=sm_100 example.cu    -o example.o
nvcc -dlto definition.o example.o -o program

# 方式二：lto_<SM> target
nvcc -dc -arch=lto_100 definition.cu -o definition.o
nvcc -dc -arch=lto_100 example.cu    -o example.o
nvcc -dlto definition.o example.o -o program

注意：device 編譯（-dc）與最終 link 兩步驟都要帶上 LTO 旗標。

2.7.4.5 Profiling Options

用 Nsight Compute 與 Nsight Systems 可直接 profile，編譯時不必加額外旗標。但下列旗標能讓 profiling 工具把原始碼與生成碼對應起來：

• -lineinfo：產生 device code 行號資訊，讓 profiling 工具顯示原始碼（工具需原始碼仍在編譯時的相同位置）。
• -src-in-ptx：把原始碼保留在 PTX 中，避開上述 -lineinfo 對檔案位置的限制；需搭配 -lineinfo。

2.7.4.6 Fatbin Compression

nvcc 預設會壓縮 binary 內的 fatbin。

• -no-compress：停用 fatbin 壓縮。
• --compress-mode={default|size|speed|balance|none}：設定壓縮模式。

2.7.4.7 Compiler Performance Controls

加速與分析「編譯過程本身」的選項：

考試/測驗重點

Related Notes

• 02-Programming-GPUs/01-CUDA-Cpp-Kernels-and-Launch
• 02-Programming-GPUs/04-CUDA-Cpp-Errors-and-Specifiers
• 02-Programming-GPUs/07-SIMT-Device-Memory-Spaces
• 02-Programming-GPUs/08-SIMT-Memory-Performance
• 02-Programming-GPUs/09-SIMT-Atomics-Cooperative-Occupancy
• 01-Introduction-to-CUDA/05-CUDA-Platform
• 02-Programming-GPUs/Practice-Programming-GPUs
• 00-Dashboard/Exam-Traps