批次記憶體傳輸與環境變數 (Batched Transfers and Environment Variables)

重點總覽

項目	重點
Batching 概念	把多個（通常較小的）task 合成單一較大的 operation，攤提（amortize）個別 dispatch 的 CPU/driver overhead；例如 cuBLAS 的 batched matrix multiplication
`cudaMemcpyBatchAsync`	批次記憶體傳輸 API（含 3D 變體 `cudaMemcpyBatch3DAsync`），用 memory copy attributes 提供排序、位置與 overlap 提示，讓 driver 能整體最佳化
`srcAccessOrder`	三種值：`...Stream` / `...DuringApiCall` / `...Any`，決定 source 的存取排序，分別對應 pinned/managed、ephemeral 指標、立即 stage
Location hints	`srcLocHint`/`dstLocHint` 用 `cudaMemLocation`（type + id）提示 source/dest 位置（device / host NUMA node），與 `cudaMemPrefetchAsync` 共用同一結構
flags（SM vs CE）	`cudaMemcpyFlagPreferOverlapWithCompute` 提示優先用 Copy Engine 並與 compute overlap；非 Tegra 平台忽略
`CUDA_DEVICE_MAX_CONNECTIONS`	增加其值可減少不同 stream 的獨立工作因 false dependency 被序列化；MPS 下預設值較低
`CUDA_MODULE_LOADING`	預設為 lazy module loading；設為 `EAGER` 可把 module 載入 overhead 移到初始化階段
設定時機	環境變數應在「啟動 app 前」設定；在 application 內設定可能無效

Batching 的動機

Batching（批次化） 泛指把數個（通常較小的）task 組成單一（通常較大的）operation。批次的各組件不必完全相同，但通常相同；典型例子是 cuBLAS 提供的 batch matrix multiplication。

與 CUDA Graphs 及 PDL（Programmatic Dependent Launch） 相同，batching 的目的都是減少個別 dispatch task 時的 overhead。
在記憶體傳輸上，launch 一次 memory transfer 會帶來一些 CPU 與 driver overhead。
一般的 cudaMemcpyAsync() 目前的形式不一定能提供足夠資訊讓 driver 最佳化傳輸（例如關於 source/destination 的 hint）。

Tip

在 Tegra 平台上可選擇用 SMs 或 Copy Engines (CEs) 執行傳輸（目前由 driver 的 heuristic 決定）。用 SM 可能傳得快，但會佔住部分算力；用 CE 可能較慢，卻能空出 SM 去做其他工作，反而提升整體 application 效能。

              傳輸速度      釋放 SM 算力
   SMs  :  快              否（佔用 compute power）
   CEs  :  慢              是（SM 可做其他工作）→ 整體 app 效能可能更高
   (僅 Tegra 可選；其他平台由 driver heuristic 決定)

cudaMemcpyBatchAsync：批次傳輸 API

cudaMemcpyBatchAsync()（以及其 3D 變體 cudaMemcpyBatch3DAsync()）讓批次記憶體傳輸可被最佳化。除了 source/destination 指標陣列與 size 陣列外，API 還用 memory copy attributes 來指定排序預期、source/dest 位置 hint，以及是否偏好把 transfer 與 compute overlap（目前僅 Tegra + CE 支援）。

最簡單的同質（homogeneous）案例：pinned host → pinned device。

std::vector<void *> srcs(batch_size);
std::vector<void *> dsts(batch_size);
std::vector<size_t> sizes(batch_size);
for (size_t i = 0; i < batch_size; i++) {
    cudaMallocHost(&srcs[i], sizes[i]);      // pinned host memory
    cudaMalloc(&dsts[i], sizes[i]);          // device memory
    cudaMemsetAsync(srcs[i], sizes[i], stream);
}
// 同一批所有 copy 共用一組 attribute
cudaMemcpyAttributes attrs = {};
attrs.srcAccessOrder = cudaMemcpySrcAccessOrderStream;
size_t attrsIdxs = 0;                        // attribute 的索引
cudaMemcpyBatchAsync(&dsts[0], &srcs[0], &sizes[0], batch_size,
                     &attrs, &attrsIdxs, 1 /*numAttrs*/, nullptr /*failIdx*/, stream);

前幾個參數很直覺：source 指標陣列、destination 指標陣列、size 陣列，每個陣列都要有 batch_size 個元素。
新資訊來自 attributes：需要一個 attribute 陣列指標，與對應的 attribute index 陣列。
failIdx 參數可傳一個 size_t 陣列來記錄失敗 transfer 的 index；傳 nullptr 是安全的，只是不會記錄失敗 index。
attrsIdxs 中第 i 個元素 = 「第 i 個 attribute 套用的第一個 transfer index」。此例 attrsIdxs = 0 視為單元素陣列，代表 attrs[0] 套用到 index 0 起算的所有 transfer。

Important

srcAccessOrder = cudaMemcpySrcAccessOrderStream 表示 source 資料以正常 stream order 存取：memcpy 會 block 直到處理這些 source/destination 指標資料的先前 kernel 完成。

attribute → transfer 對應（heterogeneous 例：numAttrs = 2）
transfers:  0  1  2 ... K-1 | K  K+1 ... batch_size-1
            └──── attrs[0] ──┘ └──── attrs[1] ─────────┘
attrsIdxs:  [0]                 [K = batch_size - 10]

srcAccessOrder：三種存取排序

異質（heterogeneous）批次：一部分來自 pinned host memory，另一部分來自只存在於當前 scope 的 stack buffer（其位址為 ephemeral pointer，API call 完成後可能失效，因為呼叫是 async 的）。

cudaMemcpyAttributes attrs[2] = {};
attrs[0].srcAccessOrder = cudaMemcpySrcAccessOrderStream;        // pinned host memory
attrs[1].srcAccessOrder = cudaMemcpySrcAccessOrderDuringApiCall; // ephemeral stack buffer
size_t attrsIdxs[2];
attrsIdxs[0] = 0;
attrsIdxs[1] = batch_size - 10;
cudaMemcpyBatchAsync(&dsts[0], &srcs[0], &sizes[0], batch_size,
                     &attrs, &attrsIdxs, 2 /*numAttrs*/, nullptr /*failIdx*/, stream);

`srcAccessOrder` 值	適用情境
`cudaMemcpySrcAccessOrderStream`	source 以正常 stream order 存取；memcpy 會 block 直到先前 kernel 完成。適用 pinned / managed memory
`cudaMemcpySrcAccessOrderDuringApiCall`	source 僅在 API call 期間被存取；用於 ephemeral pointer（如 stack buffer，API 回傳後可能失效）
`cudaMemcpyAccessOrderAny`	無排序限制；當 host buffer 非 ephemeral（如 `malloc` heap）且系統無 hardware managed memory／無 coherent GPU access via address translation 時，立即 stage 傳輸最合理

Warning

對 stack 上的 ephemeral pointer，必須用 cudaMemcpySrcAccessOrderDuringApiCall，否則指標可能在 async copy 真正執行前就失效。若改用 malloc 從 heap 配置，資料就不再 ephemeral（直到顯式 free 都有效），此時依系統是否有 hardware managed memory / coherent access 決定要用 stream ordering 或 cudaMemcpyAccessOrderAny。

Location hints 與 SM/CE flags

cudaMemcpyBatchAsync 還能提供 source/destination 位置 hint，透過 attribute 的 srcLocation／dstLocation 欄位（程式碼中為 srcLocHint／dstLocHint）設定。其型別為 cudaMemLocation（含 location 的 type 與 id），與 cudaMemPrefetchAsync() 的 prefetch hint 使用同一個結構。

cudaMemLocation srcLoc = {cudaMemLocationTypeDevice,   dev_id};   // 某 device
cudaMemLocation dstLoc = {cudaMemLocationTypeHostNuma, numa_id};  // host 的某 NUMA node
// ... cudaMallocManaged(...) + cudaMemPrefetchAsync(src/dst, ..., srcLoc/dstLoc, 0, stream) ...
cudaMemcpyAttributes attrs = {};
attrs.srcAccessOrder = cudaMemcpySrcAccessOrderStream; // managed memory → stream order 合適
attrs.srcLocHint = srcLoc;
attrs.dstLocHint = dstLoc;

最後是選擇用 SM 或 CE 的 flag，欄位為 cudaMemcpyAttributes::flags：

flag	意義
`cudaMemcpyFlagDefault`	預設行為
`cudaMemcpyFlagPreferOverlapWithCompute`	提示系統偏好用 CE 並讓 transfer 與 computation overlap；非 Tegra 平台會被忽略

cudaMemcpyBatchAsync 重點摘要

可一次指定一批 memory transfer，攤提傳輸的 setup overhead。
除 source/dest 指標與 size 外，可帶一或多個 attribute，提供：記憶體種類與 source 的 stream ordering 行為、source/dest 位置 hint、以及是否偏好 overlap（或用 SM 傳輸）。
有了這些資訊，runtime 就能盡可能最佳化傳輸。

個別 cudaMemcpyAsync × N：
 CPU: [setup][setup][setup]...      ← 每次都付 CPU/driver overhead
        cp0    cp1    cp2

批次 cudaMemcpyBatchAsync × 1：
 CPU: [setup once]                  ← 攤提 setup overhead
        { cp0, cp1, cp2, ... }      ← driver 可整體最佳化

Environment Variables

CUDA 提供多種環境變數（見規格 Section 5.2）會影響執行與效能。若未明確設定，CUDA 會用合理的預設值；但在 debug 或追求效能時可能需要逐案調整。

CUDA_DEVICE_MAX_CONNECTIONS

增加其值可降低「不同 CUDA stream 的獨立工作因 false dependency 被序列化」的可能性。
false dependency 是在多個 stream 共用同一底層資源時可能被引入的。
建議先用預設值，僅在出現效能問題（例如無法歸因於 SM 資源不足等因素的「跨 stream 獨立工作被意外序列化」）時才探索其影響。
在 MPS 情況下，此變數有不同（較低）的預設值。

CUDA_MODULE_LOADING

設為 EAGER 對 latency-sensitive 應用可能較佳：把所有 module loading overhead 移到 application 初始化階段、移出 critical phase。
目前預設模式為 lazy module loading。
在 lazy 預設模式下，可在初始化階段對各 kernel 加上 "warm-up" 呼叫，強制 module 提早載入，達到類似 eager 的效果。

Important

建議在啟動 application 前就把環境變數設成新值；嘗試在 application 內設定可能完全無效。

補充（教科書級，非本節原文）

其他常見的 CUDA 環境變數還有 CUDA_VISIBLE_DEVICES（限制/重排 process 可見的 device）與 CUDA_LAUNCH_BLOCKING（設為 1 時讓 kernel launch 變同步，便於 debug）。本章正文聚焦於上述 CUDA_DEVICE_MAX_CONNECTIONS 與 CUDA_MODULE_LOADING。

module loading 模式：
  lazy (預設)  : 用到 kernel 才載入 module  → 啟動快、首次呼叫可能多 overhead
  EAGER        : 初始化時全部載入          → critical path 無載入 overhead
  lazy + warm-up calls ≈ EAGER 的效果

考試/測驗重點

情境/關鍵字	答案
batching 的目的	攤提（分攤）個別 task dispatch 的 CPU/driver overhead，與 CUDA Graphs、PDL 同理
批次傳輸 API / 3D 變體	`cudaMemcpyBatchAsync` / `cudaMemcpyBatch3DAsync`
`attrsIdxs[i]` 的意義	第 i 個 attribute 套用的「第一個 transfer index」
`failIdx` 傳 `nullptr`	安全；只是不記錄失敗 transfer 的 index
ephemeral pointer（stack buffer）	用 `cudaMemcpySrcAccessOrderDuringApiCall`
pinned / managed memory	用 `cudaMemcpySrcAccessOrderStream`（會 block 到先前 kernel 完成）
heap malloc 且無 hardware managed/coherent access	立即 stage，用 `cudaMemcpyAccessOrderAny`
位置 hint 用什麼結構	`cudaMemLocation`（type+id），與 `cudaMemPrefetchAsync` 共用
`cudaMemcpyFlagPreferOverlapWithCompute` 在非 Tegra	被忽略
哪個平台可選 SM vs CE	Tegra（由 driver heuristic 決定）
SM 傳輸特性	較快但佔用 compute power
CE 傳輸特性	較慢但釋放 SM，整體 app 效能可能更高
減少 stream false dependency 序列化	增加 `CUDA_DEVICE_MAX_CONNECTIONS`（MPS 下預設較低）
`CUDA_MODULE_LOADING` 預設 / 改善 latency	預設 lazy；設 `EAGER` 把載入移到初始化階段
lazy 模式下近似 eager 的方法	初始化時加 "warm-up" kernel 呼叫
何時設環境變數	啟動 app 前；app 內設定可能無效

重點總覽

Batching 的動機

cudaMemcpyBatchAsync：批次傳輸 API

srcAccessOrder：三種存取排序

Location hints 與 SM/CE flags

Environment Variables

CUDA_DEVICE_MAX_CONNECTIONS

CUDA_MODULE_LOADING

考試/測驗重點

Related Notes