CUDA __constant__遵從全局內存。哪個緩存？

Question

我使用一個__constant__變量來代替將很多參數傳遞給內核。這個變量是一個結構數組，其中包含許多指向全局數據的指針（這些指針將是參數列表）;一個用於多個不同數據集的數組，用於調用內核。然後內核訪問這個數組並取消引用全局相應的數據。我的問題是，這個數據是通過L2還是常量緩存來緩存？而且，如果後者和如果通過__ldg()加載，它是否經過L1或仍然是常量高速緩存？

更具體地說，數據本身位於全局中，但是內核將取消引用__constant__變量。這會對緩存產生不利影響嗎？

Answer 1

通過立即數常量（操作碼中的常量）或索引常量（通過ldc指令訪問）訪問的常量變量由（bank，offset）對訪問，而不是通過地址訪問。這些讀取通過直接常量和索引常量高速緩存。在一些芯片上，這些是相同的緩存。不斷訪問的示例如下：

// immediate constant 
ADD r0, r1, c[bank][offset] 

// r1 has packed version of bank, offset 
LDC r0, r1 

通過cc2.0及以上版本的參數，您將看到直接不斷的訪問。

常量訪問經歷了恆定的內存層次結構，最終導致全局地址可以在系統內存或設備內存中。

如果將常量變量設置爲全局指針，則將通過數據層次結構讀取數據。

如果您定義了一個const變量，編譯器可以選擇將只讀數據放入存儲區/偏移量或地址中。

如果您查看SASS（nvdisasm或工具），您將看到LD說明。根據芯片的不同，這些數據可能會緩存在L1/Tex緩存中，然後是L2緩存。

SHARED 
LDS/STS/ATOMS    -> shared memory 

GENERIC 
LD/ST (generic to shared) -> shared memory 
LD/ST (generic to global) -> L1/TEX -> L2 
LD/ST (generic to local) -> L1/TEX -> L2 

LOCAL 
LDL/STL (local)   -> L1/TEX -> L2 

GLOBAL 
LDG/STG (global)   -> TEX -> L2 

INDEXED CONSTANT 
LDC -> indexed constant cache -> ...-> L2 

L2未命中可以進入設備內存或固定系統內存。

如果您提到常量變量很可能會通過立即常量（假設合理常量大小的最佳性能）訪問，並且解除引用的指針將導致全局內存訪問。

在GK110上LDG指令被緩存在紋理緩存中。

On Maxwell LDG.CI指令緩存在紋理緩存中。將LDG.CA操作緩存在紋理緩存（GM20x）中。所有其他LDG訪問都會經過紋理緩存，但不會超出warp指令的生命週期。