非同步資料複製與 L1/Shared 配置 (Async Data Copies and L1/Shared Config)

重點總覽

非同步資料複製 (Asynchronous Data Copies)

GPU 本就靠並行隱藏記憶體延遲：SM 在某個 warp 等待記憶體時切換去執行其他 warp。但即使如此，記憶體延遲仍可能成為頻寬利用率與運算資源效率的瓶頸。非同步資料複製把「發起搬移」與「等待完成」解耦 (decouple)，讓 threads 在資料於背景搬移時繼續做有用的工作。

copy-and-compute pattern 與同步版本

CUDA 常見的 copy-and-compute 模式：從 global memory 取資料 → 寫入 shared memory → 在 shared memory 上計算 (可能再寫回 global)。其 copy 階段 shared[local_idx] = global[global_idx] 會被編譯器展開為「先讀 global 進 register，再從 register 寫入 shared」。在迭代演算法中，每次都需要兩次同步。

shared[local_idx] = global_in[global_idx];
block.sync();                          // 等所有 copy 完成，compute 才能開始
compute(global_out + block_batch_idx, shared);
block.sync();                          // 等 compute 完成，才能覆寫 shared

要點：同步路徑經過 register file 中轉，且 thread block 必須阻塞等待整批 copy 完成後才能進入 compute。

非同步版本：memcpy_async + wait

#include <cooperative_groups/memcpy_async.h>
// 整個 thread group 協同把整批 batch 複製到 shared memory
cooperative_groups::memcpy_async(block, shared, global_in + block_batch_idx, block.size());
// 等待期間可對其他資料計算
cooperative_groups::wait(block);       // 等所有 copy 完成
compute(global_out + block_batch_idx, shared);

cooperative_groups::memcpy_async 把 block.size() 個元素從 global 複製到 shared。此操作彷彿由另一條 thread 執行，並在完成後與目前 thread 對 cooperative_groups::wait 的呼叫同步。

所有非同步複製的核心概念都一致：將搬移的發起與完成解耦。CUDA 提供多層 API 存取此能力：Cooperative Groups 與 libcu++ 的 memcpy_async、以及較低階的 cuda::ptx 與 primitives API。它們語意相近——彷彿由另一條 thread 執行複製，完成時可用不同的 completion 機制同步。

同步 copy：  [── load global ──][── store shared ──][block.sync 等待][── compute ──]
                         (threads 在搬移期間 idle)

非同步 copy：[memcpy_async 發起]┄┄┄(背景搬移)┄┄┄[wait]
             └──────── 同時可對其他資料 compute ────────┘   (隱藏延遲)

硬體機制與資料路徑

現代 GPU 架構提供多種非同步搬移硬體機制，支援不同資料路徑、傳輸大小與對齊需求：

配置 L1/Shared Memory 平衡 (Configuring L1/Shared Memory Balance)

同一個 SM 上的 L1 與 shared memory 使用相同的物理資源，稱為 unified data cache。在多數架構上，若 kernel 幾乎不用 shared memory，unified data cache 可被配置為提供該架構允許的最大 L1 cache。

            ┌──────────────── Unified Data Cache (固定總量) ────────────────┐
MaxL1   →   │ L1 cache                                      │ shared (最小) │
Default →   │ L1 cache              │ shared memory                         │
MaxShared → │ L1 (最小) │ shared memory                                     │
            └────────────────────────────────────────────────────────────┘
                        carveout 調整這條分界線 (per-kernel)

設定 carveout

reserved 給 shared memory 的比例可在 per-kernel 基礎上配置。在 kernel launch 前用 cudaFuncSetAttribute 設定 carveout (偏好的 shared memory 容量)：

cudaFuncSetAttribute(kernel_name,
                     cudaFuncAttributePreferredSharedMemoryCarveout, carveout);

carveout 可設為「該架構最大支援 shared memory 容量的整數百分比」，也可用三個便利 enum：

• cudaSharedmemCarveoutDefault
• cudaSharedmemCarveoutMaxL1
• cudaSharedmemCarveoutMaxShared

最大 shared memory 與支援的 carveout 大小因架構而異。

要點：

• 傳入 cudaFuncSetAttribute 的函式必須以 __global__ 宣告。
• carveout 被 driver 解讀為 hint；若執行 kernel 需要，driver 可選擇不同的 carveout 大小。

cudaFuncSetAttribute vs cudaFuncSetCacheConfig

超過 48 KB shared memory 的 opt-in

依賴每 block 超過 48 KB shared memory 的 kernel 是架構特定的，必須使用 dynamic shared memory (而非靜態大小陣列)，並以 cudaFuncSetAttribute 明確 opt-in：

// Device code
__global__ void MyKernel(...) {
    extern __shared__ float buffer[];
    // ...
}

// Host code
int maxbytes = 98304; // 96 KB
cudaFuncSetAttribute(MyKernel, cudaFuncAttributeMaxDynamicSharedMemorySize, maxbytes);
MyKernel<<<gridDim, blockDim, maxbytes>>>(...);

要點：靜態 __shared__ 陣列無法超過 48 KB；超過需 extern __shared__ + cudaFuncAttributeMaxDynamicSharedMemorySize，並在 launch 的 <<<...>>> 第三參數傳入動態 shared 大小。

考試/測驗重點

Related Notes

• 03-Advanced-CUDA/05-Thread-Scopes-and-Scoped-Atomics
• 03-Advanced-CUDA/06-Asynchronous-Barriers-and-Pipelines
• 03-Advanced-CUDA/01-Advanced-Launch-and-Clusters
• 03-Advanced-CUDA/04-Using-PTX-and-Hardware-Model
• 02-Programming-GPUs/14-Async-Streams-and-Events
• 02-Programming-GPUs/16-Unified-and-System-Memory
• 03-Advanced-CUDA/Practice-Advanced-CUDA
• 00-Dashboard/Exam-Traps