異質資料平行運算 練習題 (Practice - Data Parallelism and CUDA C Program Structure)

Data parallelism：對資料集不同部分要做的計算彼此獨立 (independent)，因此可同時平行進行。
亮度公式：L = 0.21*r + 0.72*g + 0.07*b。每個輸出 O[i] 只依賴對應的輸入 pixel I[i]，不依賴任何其他 pixel，故所有 pixel 可一對一指派給 thread 同時計算。

Data parallelism 依資料 (data decomposition) 拆解；task parallelism 依任務 (task decomposition) 拆解（例：vector add 與 matrix-vector mult 為兩個獨立 task，或 I/O、資料傳輸）。
Data parallelism 是主要來源：資料集越大可平行度越大，足以餵飽大規模平行硬體，效能能隨每一代擁有更多執行資源的硬體一起成長。task 數量有限，無法提供這種隨硬體放大的擴展。

CUDA 屬 SPMD (Single-Program Multiple-Data)：所有 thread 跑同一份 kernel 程式碼、處理不同資料。
差別：SIMD 要求所有處理單元在任一瞬間都執行同一條指令；SPMD 不要求各單元同時跑到同一條指令。（GPU 硬體在 warp 層級才呈現 SIMD 特性。）

• __global__：device 執行、host 呼叫（dynamic parallelism 時也可 device 呼叫）、呼叫即啟動一個新 grid，回傳須為 void。
• __device__：device 執行、只能被 kernel 或另一 device function 呼叫、不啟動新 thread。
• __host__：host 執行、只能被 host function 呼叫。
  無任何限定詞 → 預設為 __host__ (host function)。

cudaMalloc((void**)&A_d, n*sizeof(float));

• 第一參數 = 指標變數的位址 (void**)&A_d：因為 cudaMalloc 要把配置好的位址寫回你的指標變數；它是泛型函式故轉型為 generic pointer。（return value 留給回報錯誤，所以才是兩參數格式。）
• 第二參數 = bytes 數 n*sizeof(float) = n*4（不是元素數）。
• cudaFree(A_d) 傳指標值即可，不需位址。

參數順序為 cudaMemcpy(dst, src, size, kind)——目的地在前（與英文「copy A to B」直覺相反）。

• 搬入：cudaMemcpy(A_d, A_h, size, cudaMemcpyHostToDevice);
• 搬回：cudaMemcpy(C_h, C_d, size, cudaMemcpyDeviceToHost);
  同一函式靠調換 dst/src 與選對 kind 即可雙向傳輸；寫反方向是經典 bug。

i = blockIdx.x * blockDim.x + threadIdx.x（每個 thread 一對一對應一個資料元素 → loop parallelism）。
需要 if (i < n) 是因為 grid 總 thread 數 = ceil(n/blockDim.x) * blockDim.x ≥ n，幾乎總是大於 n。若不檢查，尾端多出的 thread 會越界讀寫 device global memory，造成錯誤結果或 runtime error。此檢查讓同一 kernel 能處理任意長度 n。

NVCC 依 CUDA 關鍵字分離 host / device code：

• Host code（純 ANSI C）→ 交給 host 標準 C/C++ compiler → 以傳統 CPU process 執行。
• Device code（__global__/__device__ 等）→ NVCC 編成 PTX（virtual binary 虛擬二進位），再由 NVCC 的 runtime component 進一步編成特定 GPU 的真正 object code 後在 CUDA-capable GPU 上執行。

block 數 = ceil(8000 / 1024) = ceil(7.8125) = 8 個 block。
grid threads = 8 × 1024 = 8192（其中 8000 個做事，尾端 192 個被 if (i < n) 擋掉）。答案為 8192。

(a) 每 block = 128。
(b)(c) block 數 = (200000+127)/128 = 1563（整數除法達成 ceiling division）；grid threads = 1563 × 128 = 200064，故 (c) = 1563、(b) = 200064。
(d) 第 02 行（算 i）在 if 之前，所有 thread 都執行 → 200064。
(e) 第 04 行在 if (i<N) 內，只有通過邊界檢查的 thread 執行 → 200000。

(a) 相鄰兩元素：每個 thread 拿一段連續的 2 格 → i = (blockIdx.x*blockDim.x + threadIdx.x) * 2（Ex.2 答案 C）。
(b) 兩個 section：每個 block 覆蓋 2*blockDim.x 連續元素，thread 在每段內只取一格 → i = blockIdx.x*blockDim.x*2 + threadIdx.x，第二格為 i + blockDim.x（Ex.3 答案 D）。
關鍵差異：(a) 同一 thread 取的兩格相鄰；(b) 同一 thread 取的兩格相隔 blockDim.x（利於後面章節討論的 memory coalescing）。

對每組運算，需搬 2 個 float 輸入 + 1 個 float 輸出（共 12 bytes）卻只做 1 次加法：
Arithmetic Intensity ≈ 1 FLOP / 12 bytes ≈ 0.083 FLOP/byte，屬重度 memory-bound。再加上 cudaMalloc 配置、H→D 搬入、D→H 搬回、cudaFree 等 overhead，整體可能比 CPU 版慢。
划算的做法：選計算量相對資料量大很多的 kernel（高 arithmetic intensity），並讓大型資料常駐 device memory 跨多次 kernel 呼叫重複使用，攤平搬運 overhead。

因為各 block 操作不同資料區段、彼此無順序相依，runtime 才能依硬體可用資源自由排程。

• 小 GPU：一次只能跑 1~2 個 block，其餘序列化 → 同一份程式跑得較慢。
• 大 GPU：一次可跑 64~128 個 block，平行度高 → 較快。
  兩者不需重新編譯或改寫，這就是 transparent scalability。<<<>>> 的 block 數由資料量 n 決定（ceil(n/256.0)），而非由硬體決定；硬體只決定「同時跑幾個 block」。若程式對 block 順序做假設，就會破壞這種跨硬體可攜性。