全局記憶體:
計算能力3.x的裝置的全局記憶體通路緩存在L2中,計算能力3.5或3.7的裝置的全局記憶體通路也可緩存在上一節中描述的隻讀資料緩存中; 他們通常不會被L1緩存。 計算能力3.5的某些裝置和計算能力3.7的裝置允許通過-Xptxas -dlcm = ca選項将參與緩存L1中的全局記憶體通路緩存到nvcc。
高速緩存行是128位元組,映射到裝置記憶體中128位元組對齊的段。 在L1和L2中緩存的記憶體通路都使用128位元組的記憶體事務處理,而L2中緩存的記憶體通路僅使用32位元組的記憶體事務處理。 例如,在分散存儲器通路的情況下,L2中的緩存是以可以減少過取。
如果每個線程通路的單詞大小大于4個位元組,則warp的記憶體請求首先被分割為獨立發出的單獨的128位元組記憶體請求:
兩個記憶體請求,每個半個warp請求一個,如果大小為8個位元組,
四個記憶體請求,每個四分之一warp一個,如果大小為16個位元組。
然後将每個記憶體請求分解為獨立釋出的緩存行請求。 緩存行請求在緩存命中的情況下以L1或L2緩存的吞吐量服務,否則以裝置記憶體的吞吐量服務。
請注意,線程可以按任何順序通路任何單詞,包括相同的單詞。
如果一個warp執行的非原子指令寫入全局記憶體中warp的多個線程中的同一個位置,則隻有一個線程執行寫操作,而哪個線程沒有定義。
隻讀核心整個生命周期的資料也可以通過使用<code>__ldg()</code>函數(請參閱隻讀資料高速緩存加載函數)讀取,将其緩存在上一節中描述的隻讀資料高速緩存中。 當編譯器檢測到某些資料的隻讀條件滿足時,它将使用<code>__ldg()</code>來讀取它。 編譯器可能無法始終檢測到某些資料的隻讀條件滿足。 使用const和<code>__restrict__</code>限定符标記用于加載此類資料的指針會增加編譯器檢測隻讀條件的可能性。
圖16顯示了全局記憶體通路和相應記憶體事務的一些示例。

共享記憶體:
共享記憶體有32個bank,有兩種尋址模式,如下所述。
可以使用cudaDeviceGetSharedMemConfig()查詢尋址模式,并使用cudaDeviceSetSharedMemConfig()進行設定(有關更多詳細資訊,請參閱參考手冊)。 每個bank每個時鐘周期的帶寬為64位。
圖17顯示了一些跨接通路的例子。
圖18顯示了一些涉及廣播機制的記憶體讀取通路示例。
64位模式:
連續的64位字映射到連續的bank。
對于一個warp的共享記憶體請求不會在通路同一個64位字内任何子字的兩個線程之間産生bank沖突(即使這兩個子字的位址落在同一個bank中):在這種情況下 對于讀取通路,将64位字廣播給請求的線程,對于寫入通路,每個子字隻由其中一個線程寫入(該線程執行寫入操作是未定義的)。
32位模式:
連續的32位字映射到連續的bank。
對于一個warp的共享記憶體請求不會在通路同一32位字内的任何子字的兩個線程之間或兩個32位字中的索引i和j處于相同的64字對齊段内 (即第一個索引是64的倍數的段)并且使得j = i + 32(即使兩個子字的位址落入同一個存儲體):在這種情況下,對于讀取通路,32 位字被廣播給送出請求的線程和寫入通路,每個子字隻被其中一個線程寫入(該線程執行寫入操作是未定義的)