CUDA學習筆記(一) 詳情 - 筆記,cuda,緩存,Memory,API,Python,後端開發數據分析大師博客

__ldg

功能：從 global memory 通過只讀緩存（read-only cache）加載數據。

template <typename T>
__device__ __forceinline__ T __ldg(const T* ptr);

注意這並不是函數原型，但是你可以這麼理解，T並不是所有類型都能適配的。

在老架構（Maxwell/Kepler）上，它會使用 LDG 指令，利用只讀緩存（texture cache / L1 read-only）。
在現代架構（Pascal+）仍然有效，但新架構 L2 的自動緩存更智能。
應用場景：當你的 global memory 數據只讀且多線程共享，使用 __ldg 可以減少對 L1 的壓力，提高帶寬利用。

__global__ void kernel(const float* data, float* out) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    float val = __ldg(&data[idx]);
    out[idx] = val * 2.0f;
}

架構代號	Compute Capability	Read-Only Cache 狀態	`__ldg()` 行為
Kepler (GK110)	3.5	✅ 獨立的只讀緩存	有效，大幅優化
Maxwell / Pascal	5.x / 6.x	✅ 仍保留，只讀緩存共享與 L1	有效，但作用減弱
Volta / Turing / Ampere	7.x / 8.x	⚠️ 與 L1 Cache 合併	`__ldg()` 存在，但與普通 `*ptr` 效果幾乎一樣
Hopper (H100)	9.0	❌ 不再單獨實現	`__ldg()` 仍編譯，但僅作普通 load

在 Ampere 及以後：

Read-only cache 不再是獨立單元；
所有 load 都經過一個統一的 L1 Cache（Unified Data Cache）；
__ldg() 仍然存在，編譯器不會報錯；
但 PTX 層會退化為普通的 LDG.E 或 LDG.G 指令，與 *ptr 無區別。

__prefetch_global

功能：提前將 global memory 數據加載到 cache，隱藏訪問延遲。

template <typename T>
__device__ void __prefetch_global(const T* ptr);

説明：

Ampere 架構開始支持。
只是 hint，不保證立即訪問到數據，只是告訴硬件“接下來可能用到”。
對延遲敏感的循環訪問場景非常有用。

__global__ void kernel(float* data) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    __prefetch_global(&data[idx]);
    float val = data[idx];  // 實際訪問時可能已經在 L1/L2
}

但是普通的遊戲顯卡，比如3090，4090可能不支持

架構	GPU類型	SM版本	是否支持 `__prefetch_global()`	説明
A100	數據中心	sm_80	✅ 支持	pipeline copy / prefetch 全支持
H100	數據中心	sm_90	✅ 支持	支持更多 pipeline 指令
RTX 3090 / 3080	消費級	sm_86	❌ 不支持	指令未開放，prefetch 會報錯
RTX 4090 (Ada)	消費級	sm_89	❌ 不支持	同樣未公開 prefetch API

ldcs() / ldcg() / __ldca()

功能：控制加載時的緩存策略。

API	含義
`__ldcs()`	load via streaming cache（只 L2）
`__ldcg()`	load global memory (cached)
`__ldca()`	load via L1 cache

template <typename T>
__device__ T __ldcs(const T* ptr);
template <typename T>
__device__ T __ldcg(const T* ptr);
template <typename T>
__device__ T __ldca(const T* ptr);

説明：

Ampere/Hopper 架構提供，允許程序員控制 L1/L2 的命中策略。
對優化多線程數據共享、減少 L1 衝突有幫助。
對大多數通用算法不是必需，更多是高性能調優手段。

stcs() / stcg() / __stwb()

功能：控制 store（寫）時緩存策略。

API	含義
`__stcs()`	streaming store (L2 only)
`__stcg()`	cached store
`__stwb()`	write-back store

説明：

可以降低 L1 負擔或者控制寫合併策略。
對性能優化高併發寫場景有時有效。

__pipeline_memcpy_async

功能：異步拷貝數據到 shared memory / register / L2 pipeline。

原型（示意）：

template <typename T>
__device__ void __pipeline_memcpy_async(T* dst, const T* src, size_t n);

説明：

Ampere+ GPU 支持。
類似 cp.async 指令，但 API 級別，更安全。
允許 overlap memory copy 與 compute，隱藏 global memory 訪問延遲。
必須配合 __pipeline_commit() 和 __pipeline_wait_prior() 控制階段。

`__pipeline_commit()` / `__pipeline_wait_prior()`

功能：控制 pipeline 的異步拷貝階段。

用法：

__pipeline_commit()：提交之前的 async copy 指令到 pipeline。
__pipeline_wait_prior()：等待 pipeline 中之前提交的 copy 完成。

説明：

通常配合循環或雙緩衝使用，實現 compute 與 copy 的 overlap。
相當於 Ampere+ 的“軟件可控 cp.async”。

for(int i=0;i<N;i+=tile){
    __pipeline_memcpy_async(sm_tile, &gmem[i], tile);
    __pipeline_commit();
    __pipeline_wait_prior();  // 等待上一次拷貝完成
    compute(sm_tile);
}

GPU緩存結構

CUDA GPU 的存儲層級（從慢到快）大致是：

Global Memory (DRAM)
        │
        ▼
      L2 Cache  ← 所有 SM 共享
        │
 ┌──────┴──────────┐
 │                 │
 ▼                 ▼
L1 Data Cache   Read-Only Cache (texture)
 │
 ▼
Registers / Shared Memory

Global Memory (DRAM)：顯存，全局可訪問，但延遲高（幾百個 cycles）。
L2 Cache：芯片級共享緩存，所有 SM 都能訪問，延遲較低（幾十個 cycles）。
L1 Cache：每個 SM 獨立的一級緩存，延遲更低（10-20 cycles）。
Read-Only Cache：專門優化只讀訪問（只在老架構中單獨存在，現在常與 L1 合併或共享）。
Registers / Shared Memory：線程或線程塊級的高速存儲，延遲極低。

我們來對照看下不同指令的緩存行為。

指令 / API	訪問路徑	緩存層級	特點	典型用途
`__ldg`	Global → Read-only cache → Register	L2 + Read-only(L1T)	只讀緩存，不污染普通 L1	共享常量、查表數據
`__ldcg`	Global → L2 → L1 → Reg	L2 + L1	默認路徑（全緩存）	一般數據加載
`__ldca`	Global → L1 only → Reg	L1 only（不走 L2）	強制優先用 L1，L2可不命中	臨時數據，多次局部訪問
`__ldcs`	Global → L2 only → Reg	L2 only（跳過 L1）	不污染 L1 cache	大流量讀取（streaming read）

但實際上，這些全是“建議式控制”，不是顯式管理。

L1/L2/只讀緩存的對比總結

緩存層	作用範圍	容量	一致性	可寫	典型用途
L1 Data Cache	每個 SM 獨立	小（128KB~192KB）	不全局一致	可寫	局部數據緩存
Read-only Cache	每個 SM 獨立	小（48KB~128KB）	無需一致性	只讀	常量、查表
L2 Cache	全 SM 共享	大（幾 MB）	全局一致	可寫	跨 SM 通信、共享數據

實際調優經驗

場景	推薦策略	理由
常量查表 / 只讀共享	`__ldg`	走 read-only cache，不污染 L1
大數據流式讀	`__ldcs`	避免 L1 被大量數據污染
局部複用強（tile/block 內）	`__ldca`	強制走 L1，本地複用效率高
一般數據訪問	默認 / `__ldcg`	默認行為即可
Ampere+ 異步拷貝	`__pipeline_memcpy_async`	取代手動 prefetch，性能更高

SM 與緩存的關係

SM（Streaming Multiprocessor） 是 GPU 的計算核心單元，就像 CPU 裏的“核（core）”。
一張 GPU（比如 A100、RTX 4090）內部通常有 幾十到上百個 SM。

CUDA 的層級關係：

Grid → Block → Warp → Thread

而 SM 是執行的硬件實體：

概念	含義	對應的硬件
Grid	一次 kernel 啓動的所有線程集合	GPU 整體
Block	一組線程（可通信/同步）	分配給一個 SM 執行
Warp	32 個線程組成的調度單元	SM 內的調度器發射
Thread	單個執行流	SM 內的計算核心（ALU）

每個 block 在運行時都會被分配到某個 SM 上執行，
block 內的所有線程都在那個 SM 的共享資源（如 shared memory、L1 cache）中運行。

每個 SM 有自己獨立的 L1 Cache 和 Shared Memory，所以：

一個 SM 裏的線程塊（block）可以在 shared memory 中高速共享數據；
不同 SM 之間的數據共享必須通過 L2 Cache 或 Global Memory；
L1 cache 不同步（每個 SM 自己的 L1 可能不一致）；
L2 cache 在所有 SM 間是 一致的（coherent）。

這也就是為什麼有下面這種訪問策略：

緩存類型	是否 SM 共享	一致性	示例 API
L1 Cache	❌ 獨立	不一致	`__ldca()`
Read-only Cache	❌ 獨立	不一致	`__ldg()`
L2 Cache	✅ 所有 SM 共享	一致	`__ldcs()`
Global Memory	✅ 所有 SM 共享	一致	—

SM的小知識

可以用下面的代碼查看自己 GPU 的 SM 數量：

cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, 0);
std::cout << "SM數量: " << prop.multiProcessorCount << std::endl;

比如3090輸出82

“核函數一次最多啓動的 block 數不能超過 SM 數”

這是 不對的。
block 數可以遠遠多於 SM 數，只不過 同時在 GPU 上運行的 block 數受到 SM 數量和資源限制。

正確理解

核函數啓動時你寫的：

myKernel<<<numBlocks, threadsPerBlock>>>(...);

裏指定的 numBlocks 是整個 Grid 的大小。
GPU 會把這些 block 分批調度到 SM 上執行。

舉個具體例子（比如 A100）：

每個 SM 最多同時運行 2048 線程
每個 block 有 256 線程
那麼每個 SM 同時最多能運行 8 個 block
若總共有 108 個 SM，則同時最多有 108 × 8 = 864 個 block 在執行
但如果你啓動了 10,000 個 block，其他的會排隊等前面的執行完再上。

層次	存儲位置	是否共享	説明
寄存器 (Registers)	每個線程私有	❌	每個線程自己的最快速存儲空間（延遲 < 10 cycles）。編譯器自動分配。
共享內存 (Shared Memory)	每個 Block 獨佔	✅（同 Block 內共享）	位於 SM 內部的片上 SRAM。延遲 ~100 cycles，比 global memory 快很多。
L1 Cache	每個 SM 獨佔	✅（同 SM 上的所有 Block 共享）	用於緩存 global memory 加載的結果，可配置與 Shared Memory 共用物理空間。
L2 Cache	GPU 所有 SM 共享	✅（全局共享）	位於芯片片上（on-chip），負責不同 SM 之間的數據共享與一致性。
Global Memory (DRAM)	芯片外顯存	✅（全局共享）	訪問延遲幾百到上千 cycles。

本文章為轉載內容，我們尊重原作者對文章享有的著作權。如有內容錯誤或侵權問題，歡迎原作者聯繫我們進行內容更正或刪除文章。

數據分析大師博客

數據分析大師博客

博客 / 詳情