呼叫 Kernel、Compilation 與本章總結 (Calling Kernels, Compilation, and Summary)

重點總覽 (Overview)

呼叫 Kernel:執行配置參數 (Calling Kernels: Execution Configuration)

實作完 kernel 後,最後一步是從 host code 呼叫 kernel 以 launch grid。CUDA C 在傳統 C 的函式呼叫語法上,加入夾在 <<< 與 >>> 之間的 execution configuration parameters。

// Launch ceil(n/256) blocks of 256 threads each
vecAddKernel<<<ceil(n/256.0), 256>>>(A_d, B_d, C_d, n);

• execution configuration parameters 只在 launch kernel 時使用,普通函式呼叫沒有這組參數。
• 同一個 kernel 可在 host code 不同位置用不同 thread 數呼叫。

Ceiling Division:為何要無條件進位

要確保 grid 內 thread 數足夠覆蓋所有 n 個元素,block 數必須取「期望 thread 數 / block size」的無條件進位 (ceiling):

ceil(n/256.0)   // 正確:用 256.0 (浮點) 確保產生浮點商,ceil 才能正確進位
// n/256        // 錯誤:整數除法會無條件捨去,漏掉最後一個不滿的 block

Grid 幾何示意 (n=1000, blockSize=256)

gridDim = ceil(1000/256.0) = 4 blocks      blockDim = 256 threads/block

 Block 0          Block 1          Block 2          Block 3
[t0 ... t255]    [t0 ... t255]    [t0 ... t255]    [t0 ........ t255]
 i = 0..255       i = 256..511     i = 512..767     i = 768..1023
   全部工作          全部工作          全部工作       i<1000 工作 │ i>=1000 閒置
                                                    (768..999)  (1000..1023 被 if 擋掉)

   i = blockIdx.x * blockDim.x + threadIdx.x   <-- 每個 thread 的全域索引

透明擴展性 (Transparent Scalability)

<<<>>> 中的 block 數量隨資料量 n 決定,而非隨硬體決定:

關鍵性質:所有 block 操作不同資料區段,彼此獨立,可以任意順序執行。

同一份 binary,不同硬體上自動擴展:

  Grid (例如 16 個 blocks)
  ┌──┬──┬──┬──┬──┬──┬──┬──┬──┬──┬──┬──┬──┬──┬──┬──┐
  │B0│B1│B2│B3│B4│B5│B6│B7│B8│B9│..│..│..│..│..│15│
  └──┴──┴──┴──┴──┴──┴──┴──┴──┴──┴──┴──┴──┴──┴──┴──┘
        │                                  │
   小 GPU:一次跑 1~2 個 block         大 GPU:一次跑 64~128 個 block
   (序列化較多 → 較慢)                  (平行度高 → 較快)

編譯流程 (Compilation: NVCC and PTX)

kernel 用到的 CUDA 擴充(如 __global__、<<<>>>、built-in 變數)不被傳統 C 編譯器接受,必須用能辨識這些擴充的編譯器 —— NVCC (NVIDIA C Compiler)。

                    CUDA C program (.cu)
                  host code + device code
                           │
                  ┌────────▼────────┐
                  │      NVCC        │  依 CUDA 關鍵字分離 host / device code
                  └───┬─────────┬────┘
        Host Code     │         │      Device Code (__global__/__device__)
     (straight ANSI C)│         │
            ┌─────────▼──┐   ┌──▼───────────────┐
            │ Host C/C++ │   │ NVCC → PTX files  │  (virtual binary 虛擬二進位)
            │ Compiler   │   └──┬───────────────┘
            └─────┬──────┘      │  runtime component of NVCC
                  │          ┌──▼───────────────┐
        傳統 CPU process     │ real object files │
                             └──┬───────────────┘
                                ▼
                        執行於 CUDA-capable GPU

• PTX 是 virtual binary (虛擬指令集),不是最終機器碼。
• PTX 由 NVCC 的 runtime component 進一步編譯成特定 GPU 的真正 object code(這也是 transparent scalability / 跨硬體可攜性的基礎之一)。

本章總結:CUDA C 擴充 (Chapter Summary: CUDA C Extensions)

本章介紹的 CUDA C 對 C 的核心擴充,可歸納為四類:

// ① 三種 function qualifier 的呼叫關係
__global__ void k(...);   // 在 device 執行;host 可呼叫(launch grid)
__device__ float d(...);  // 在 device 執行;只能被 kernel/device function 呼叫
__host__   void h(...);   // 在 host   執行;只能被 host function 呼叫(預設)
__host__ __device__ float f(...); // 編譯器產生 host + device 兩版,單一份原始碼

考試/面試重點 (Exam / Test Patterns)

Related Notes

• 02-Heterogeneous-Data-Parallel-Computing/01-Data-Parallelism-and-CUDA-Program-Structure
• 02-Heterogeneous-Data-Parallel-Computing/02-Vector-Addition-Host-Code-Device-Memory-and-Data-Transfer
• 02-Heterogeneous-Data-Parallel-Computing/03-Kernel-Functions-and-Threading
• 03-Multidimensional-Grids-And-Data/01-Multidimensional-Grid-Organization
• 04-Compute-Architecture-And-Scheduling/01-GPU-Architecture-and-Block-Scheduling
• 06-Performance-Considerations/03-Thread-Coarsening