計算能力5.x
Architecture:
多處理器由以下部分組成:
用于算術運算的128個CUDA核心(請參見算術運算吞吐量的算術指令),
32個用于單精度浮點超越函數的特殊功能單元,
4個warp排程器
當多處理器被執行時,它首先在四個排程器中進行配置設定。 然後,在每個指令釋出時間,每個排程程式為其配置設定的一個warp指令發出一條指令,如果有的話,該指令即可執行。
多處理器具有:
一個由所有功能單元共享的隻讀常量高速緩存,可加速駐留在裝置記憶體中的常量記憶體空間的讀取速度,
一個24 KB的統一L1 /紋理緩存,用于緩存全局記憶體中的讀取,
計算能力5.0或96 KB共享記憶體的裝置共享記憶體的64 KB共享記憶體用于計算能力的裝置5.2。
紋理單元也使用統一的L1 /紋理緩存,實作紋理和表面存儲器中提到的各種尋址模式和資料過濾。
還有一個由所有多處理器共享的L2緩存,用于緩存對本地或全局記憶體的通路,包括臨時寄存器溢出。 應用程式可以通過檢查l2CacheSize裝置屬性來查詢L2緩存大小。
高速緩存行為(例如,讀取是否高速緩存在統一的L1 /紋理高速緩存和L2中或L2中)可以使用對加載指令的修飾符以逐通路為基礎部配置設定置。
全局記憶體:
全局記憶體通路總是緩存在L2中,而L2中的緩存與計算功能3.x的裝置(參見全局記憶體)的行為方式相同。
隻讀核心整個生命周期的資料也可以通過使用<code>__ldg()</code>函數(請參見隻讀資料高速緩存加載函數)讀取,将其緩存在上一節所述的統一L1 /紋理高速緩存中。 當編譯器檢測到某些資料的隻讀條件滿足時,它将使用<code>__ldg()</code>來讀取它。 編譯器可能無法始終檢測到某些資料的隻讀條件滿足。 使用const和<code>__restrict__</code>限定符标記用于加載此類資料的指針會增加編譯器檢測隻讀條件的可能性.
對于核心整個生命周期不是隻讀的資料不能緩存在計算能力為5.0的裝置的統一L1 /紋理緩存中。 對于計算能力5.2的裝置,預設情況下,它不會緩存在統一的L1 /紋理緩存中,但可以使用以下機制啟用緩存:
使用PTX參考手冊中所述的相應修改器,使用内聯彙編執行讀取;
使用-Xptxas -dlcm = ca編譯标志進行編譯,在這種情況下,所有讀取都将被緩存,但使用内聯彙編使用禁用緩存的修飾符執行的讀取除外;
使用-Xptxas -fscm = ca編譯标志進行編譯,在這種情況下,所有讀取都将被緩存,包括使用内聯彙編執行的讀取,而不管所使用的修改器如何。
當使用上面列出的三種機制啟用緩存時,計算能力5.2的裝置将全局記憶體讀取緩存到統一的L1 /紋理緩存中,用于所有核心啟動,除了核心啟動之外,其内的線程塊消耗太多的多處理器資源。 這些異常由分析器報告。
共享記憶體:
共享記憶體有32個組織,這樣連續的32位字映射到連續的組。 每個bank每個時鐘周期的帶寬為32位。
對于一個warp的共享記憶體請求不會在通路同一個32位字内的任何位址的兩個線程之間産生bank沖突(即使這兩個位址落在同一個bank中):在這種情況下,對于讀取通路, 被廣播到請求線程和寫入通路,每個位址僅由其中一個線程寫入(該線程執行寫入操作是未定義的)
