L2 Cache Control

重點總覽

當 kernel 反覆存取同一塊 global memory（persisting），可把一部分 L2 cache「保留」（set-aside）給這些存取優先使用，藉此提升頻寬、降低延遲；只被讀一次的資料則屬於 streaming。此能力需要 compute capability 8.0 以上。

L2 Cache Set-Aside for Persisting Accesses

可從 L2 cache 撥出一塊「保留區」（set-aside）給 persisting 存取使用。

• persisting 存取對這塊保留區有優先使用權；normal 或 streaming 存取只有在保留區未被 persisting 使用時才能借用。
• 保留大小可在上限內調整，上限為 prop.persistingL2CacheMaxSize。

cudaGetDeviceProperties(&prop, device_id);
size_t size = min(int(prop.l2CacheSize * 0.75), prop.persistingL2CacheMaxSize);
cudaDeviceSetLimit(cudaLimitPersistingL2CacheSize, size); // 保留 3/4 L2 或最大允許值

上面取「L2 的 75%」與「硬體允許上限」中較小者，作為 set-aside 大小。

L2 Policy for Persisting Accesses

access policy window 指定一段連續的 global memory 範圍與其在 L2 的持久化屬性；落在窗口內的存取依機率套用 hitProp，其餘套用 missProp。可掛在 CUDA Stream 或 CUDA Graph Kernel Node 上。

cudaStreamAttrValue stream_attribute;                       // Stream 層級屬性結構
stream_attribute.accessPolicyWindow.base_ptr  = reinterpret_cast<void*>(ptr);
stream_attribute.accessPolicyWindow.num_bytes = num_bytes;  // < accessPolicyMaxWindowSize
stream_attribute.accessPolicyWindow.hitRatio  = 0.6;        // cache 命中比例提示
stream_attribute.accessPolicyWindow.hitProp   = cudaAccessPropertyPersisting;
stream_attribute.accessPolicyWindow.missProp  = cudaAccessPropertyStreaming;
cudaStreamSetAttribute(stream, cudaStreamAttributeAccessPolicyWindow, &stream_attribute);

此後在該 stream 執行的 kernel，對 [ptr..ptr+num_bytes) 的存取比其他位置更可能持久留在 L2。Graph 版本則用 cudaKernelNodeAttrValue 搭配 cudaGraphKernelNodeSetAttribute(node, cudaKernelNodeAttributeAccessPolicyWindow, &node_attribute)，欄位完全相同。

• num_bytes 必須小於 cudaDeviceProp::accessPolicyMaxWindowSize。
• 上例 hitRatio = 0.6：窗口內約 60% 存取套用 persisting，40% 套用 streaming。
• 哪些存取被歸為 persisting 是隨機的，機率約等於 hitRatio；實際機率分布取決於硬體架構與記憶體範圍。

hitRatio 如何避免 thrashing

以 set-aside 為 16KB、num_bytes 為 32KB 的窗口為例：

set-aside L2 = 16KB        window = 32KB
hitRatio=1.0  ┌──────────────┬──────────────┐  整窗 32KB 想塞進 16KB
              │  ←──── 不斷互相驅逐(thrash) ──→ │
              └──────────────┴──────────────┘
hitRatio=0.5  ┌──────────────┐ 只標 16KB persisting，剛好放下，不抖動
              └──────────────┘

• hitRatio < 1.0 可手動控制不同 stream 的 accessPolicyWindow 能在 L2 快取多少資料。
• 例：兩個並行 kernel 各有 16KB 窗口、共用 16KB 保留區；都用 hitRatio = 1.0 會互相驅逐 cache line，改用 0.5 則較不會驅逐自己或對方的 persisting cache line。

L2 Access Properties

針對不同 global memory 存取定義三種屬性：

L2 Persistence Example

完整流程：撥出 set-aside → 設窗口 → kernel 反覆使用 → 關閉窗口並 reset L2。

cudaStreamCreate(&stream);
cudaGetDeviceProperties(&prop, device_id);
size_t size = min(int(prop.l2CacheSize * 0.75), prop.persistingL2CacheMaxSize);
cudaDeviceSetLimit(cudaLimitPersistingL2CacheSize, size);            // 保留 set-aside
size_t window_size = min(prop.accessPolicyMaxWindowSize, num_bytes); // 取較小者

stream_attribute.accessPolicyWindow.base_ptr  = reinterpret_cast<void*>(data1);
stream_attribute.accessPolicyWindow.num_bytes = window_size;
stream_attribute.accessPolicyWindow.hitRatio  = 0.6;
stream_attribute.accessPolicyWindow.hitProp   = cudaAccessPropertyPersisting;
stream_attribute.accessPolicyWindow.missProp  = cudaAccessPropertyStreaming;
cudaStreamSetAttribute(stream, cudaStreamAttributeAccessPolicyWindow, &stream_attribute);

for (int i = 0; i < 10; i++)
    cuda_kernelA<<<grid_size, block_size, 0, stream>>>(data1); // data1 被多次使用 → 受惠
cuda_kernelB<<<grid_size, block_size, 0, stream>>>(data1);     // 同 stream 另一 kernel 也受惠

stream_attribute.accessPolicyWindow.num_bytes = 0;             // 窗口大小設 0 即停用
cudaStreamSetAttribute(stream, cudaStreamAttributeAccessPolicyWindow, &stream_attribute);
cudaCtxResetPersistingL2Cache();                               // 清除 L2 中的 persistent line
cuda_kernelC<<<grid_size, block_size, 0, stream>>>(data2);     // data2 可用完整 L2(normal 模式)

重點：用完 persisting 後務必先把 num_bytes 設 0 關閉窗口，再呼叫 cudaCtxResetPersistingL2Cache()，後續 kernel 才能以 normal 優先權完整使用 L2。

set-aside ──► set window(data1) ──► kernelA ×10 ──► kernelB
                                                       │
                          num_bytes=0(關窗) ◄──────────┘
                                  │
                  cudaCtxResetPersistingL2Cache() ──► kernelC(data2, 完整 L2)

Reset L2 Access to Normal

persisting cache line 可能在用完後仍長期留在 L2，因此「重設回 normal」對 streaming/normal 存取以正常優先權使用 L2 很重要。三種重設方式：

1. 以 cudaAccessPropertyNormal 屬性重設先前的 persisting 記憶體區。
2. 呼叫 cudaCtxResetPersistingL2Cache() 把所有 persisting cache line 重設為 normal。
3. 未被觸碰的 line 最終會自動重設為 normal。

Manage Utilization of L2 Set-Aside Cache

多個並行（不同 stream）的 kernel 可各自指定 access policy window，但保留區是被所有並行 kernel 共用的：總使用量是各 kernel 個別用量的加總。當 persisting 存取總量超過保留區容量時，持久化的好處會遞減。

管理保留區利用率須一併考量：

• L2 set-aside 的大小。
• 可能並行執行的 CUDA kernel。
• 這些並行 kernel 各自的 access policy window。
• 何時、如何 reset L2，讓 normal/streaming 存取能以相同優先權使用先前保留的 L2。

Query L2 Cache Properties

L2 相關屬性屬於 cudaDeviceProp，用 cudaGetDeviceProperties 查詢：

Control L2 Cache Set-Aside Size for Persisting Memory Access

set-aside 大小以 cudaLimit 形式管理：用 cudaDeviceGetLimit 查詢、用 cudaDeviceSetLimit 設定，limit 名稱為 cudaLimitPersistingL2CacheSize，最大值為 cudaDeviceProp::persistingL2CacheMaxSize。

enum cudaLimit {
    /* other fields not shown */
    cudaLimitPersistingL2CacheSize
};

cudaLimitPersistingL2CacheSize 即是控制 L2 保留大小的旋鈕，搭配 set/get Limit 一對即可調整與讀回。

考試/測驗重點

Related Notes

• 04-CUDA-Features/03-CUDA-Graphs-Structure-and-Capture
• 04-CUDA-Features/18-Memory-Synchronization-Domains
• 04-CUDA-Features/13-Async-Copies-LDGSTS
• 03-Advanced-CUDA/07-Async-Data-Copies-and-L1-Config
• 03-Advanced-CUDA/10-Tour-of-CUDA-Features
• 04-CUDA-Features/14-Async-Copies-TMA
• 04-CUDA-Features/20-Virtual-Memory-Management
• 04-CUDA-Features/Practice-CUDA-Features
• 00-Dashboard/Exam-Traps