爲什麼CUDA中的恆定內存大小受到限制？

根據"CUDA C Programming Guide"，只有在多處理器常量高速緩存命中的情況下，常量內存訪問纔會有所收益（第5.3.2.4節）。否則，與整合全局存儲器讀取的情況相比，可能會有更多的半內存請求。那麼爲什麼常量內存大小限制爲64 KB？爲什麼CUDA中的恆定內存大小受到限制？

還有一個問題，以免問兩次。據我瞭解，在費米架構中，紋理緩存與L2緩存相結合。紋理的使用仍然有意義或全局內存讀取以相同的方式緩存？

常量內存（節5.3.2.4）

恆定存儲器空間駐留在設備存儲器，並且在第F.3.1和F.提到的常數的高速緩存被緩存4.1。

對於計算能力1.x的設備，對於一個warp的恆定內存請求首先被分成兩個請求，每個half-warp一個，分別發出。

然後，請求被分割爲多個獨立的請求，因爲初始請求中存在不同的內存地址，吞吐量減少了一個等於單獨請求數量的因子。

然後，以高速緩存命中情況下的常量高速緩存的吞吐量或其他情況下的設備內存吞吐量來處理結果請求。

來源

2012-04-21 AdelNick

計算能力1.0-3.0設備的常量內存大小爲64 KB。緩存工作集只有8KB（請參見CUDA編程指南v4.2表F-2）。

驅動程序，編譯器和聲明爲__device__ __constant__的變量使用常量內存。驅動程序使用常量內存來傳遞參數，紋理綁定等。編譯器在許多指令中使用常量（請參見反彙編）。

使用主機運行時功能cudaMemcpyToSymbol()和cudaMemcpyFromSymbol()（參見CUDA編程指南v4.2第B.2.2節）可以讀取和寫入置於恆定內存中的變量。常量內存位於設備內存中，但通過常量緩存進行訪問。

關於費米紋理，常量，L1和I-Cache都是每個SM中或其周圍的1級高速緩存。所有1級緩存都通過L2緩存訪問設備內存。

64 KB的常量限制是CUmodule，它是一個CUDA編譯單元。 CUmodule的概念隱藏在CUDA運行時環境下，但可通過CUDA驅動程序API訪問。

來源

2012-04-21 20:59:25

格雷格，我很抱歉可能不夠清楚。我知道如何使用常量內存。在這個問題中，我想知道爲什麼只有緩存8 KB才能提供比全局內存更好的性能，爲什麼常量內存的大小限制爲64 KB。以同樣的方式可以通過L1緩存訪問全局內存。因此，從編程人員的角度來看，使用全局內存或常量內存之間的區別是什麼，因爲它們都以相同的方式被緩存？ – AdelNick 2012-04-22 16:14:33

常量緩存的大小適合圖形和計算着色器。 CUDA API公開了常量緩存，以便開發人員可以利用額外的緩存。當所有線程訪問相同的地址時，常量緩存具有最佳性能。點擊速度非常快。錯過可以有一個L1錯過的相同表現。常量緩存可用於減少L1的抖動。計算能力1.x設備沒有L1或L2緩存，因此常量緩存用於優化某些訪問。 – 2012-04-22 19:57:02

爲什麼CUDA中的恆定內存大小受到限制？

回答

相關問題