前言:

本章內容：

• 了解數據在共享內存中是如何被安排的

• 掌握從二維共享內存到線性全局內存的索引轉換

• 解決不同訪問模式中存儲體中的沖突

• 在共享內存中緩存數據以減少對全局內存的訪問

• 使用共享內存避免非合并全局內存的訪問

• 理解常量緩存和只讀緩存之間的差異

• 使用線程束洗牌指令編程

在前面的章節中, 已經介紹了幾種全局內存的訪問模式. 通過安排全局內存訪問模式, 我們學會了如何實現良好的性能并且避免了浪費事務. 未對齊的內存訪問是沒有問題的, 因為現代的GPU硬件都有一級緩存, 但在跨全局內存的非合并內存訪問, 仍然會導致帶寬利用率不會達到最佳標準. 根據算法性質和相應的訪問模式, 非合并訪問可能是無法避免的. 然而, 在許多情況下, 使用共享內存來提高全局內存合并訪問是有可能的. 共享內存是許多高性能計算應用程序的關鍵驅動力.

在本章中, 你將學習如何使用共享內存進行編程、數據在共享內存中如何被存儲、數據元素是怎樣使用不同的訪問模式被映射到內存存儲體中的. 還將掌握使用共享內存提高核函數性能的方法.

5.1 CUDA共享內存概述:

GPU中有兩種類型的內存：

• 板載內存(以內存顆粒的形式貼于顯卡PCB上)

• 片上內存(集成于芯片內部)

全局內存是較大的板載內存, 具有相對較高的延遲. 共享內存是較小的片上內存, 具有相對較低的延遲, 并且共享內存可以提供比全局內存高得多的帶寬. 可以把它當作一個可編程管理的緩存. 共享內存通常的用途有：

• 塊內線程通信的通道

• 用于全局內存數據的可編程管理的緩存

• 高速暫存存儲器, 用于轉換數據以優化全局內存訪問模式

共享內存:

這里就給原文了, 之前那些奇怪的翻譯怎么就不給

共享內存（shared memory, SMEM）其特點:

• 每個SM上都有一個獨立的共享內存
  其作用更像L1 & L2緩存
• 被SM上執行的所有線程共享
  通常用于線程間的相互協作, 大大降低了核函數所需的全局內存帶寬
• 通過程序顯式的管理
  所以稱之為可編程管理的緩存
• 帶寬比全局內存塊10倍, 而延時通常低20倍以上
  物理上更接近CUDA核心

以Kepler核心的SM為例:

SM基本存儲順序:

共享內存訪問事物:

與全局內存相同, 線程通過類似的方式訪問共享內存, 這里不再贅述

但如果過個線程訪問共享內存中個同一個字, 則在一個線程讀取該字后, 將會通過多播的形式廣播給其他線程

可編程管理的緩存:

緩存(L1 & L2)對于程序而言是透明的, 編譯器才能處理所有數據的移動, 而并非程序員

而共享內存是一個可編程管理的緩存, 所以可以通過在數據布局上提供更多的細粒度控制和改善片上數據的移動, 使得對優化應用程序代碼變得更簡單

共享內存的分配:

共享內存使用__shared__修飾符進行聲明

如:

__shared__ float tile[size_y][size_x];

如果一個共享內存的大小在編譯時是未知的(相當于每個線程使用時大小不一樣), 則需要添加extern修飾
并且==此時只能聲明一維數組==

extern __shared__ int tile[];

在每個核函數被調用時, 需要動態分配共享內存 這部分操作在主機端進行

即在<<<>>>后頭多加一個參數, 注意這里是以字節為單位

kernel<<<grid, block, isize * sizeof(int)>>>(...)

共享內存存儲體和訪問模式

優化內存性能時要度量的兩個關鍵屬性是：延遲和帶寬

共享內存可以用來隱藏全局內存延遲和帶寬對性能的影響(第四章所述)

內存存儲體:

為了獲得高內存帶寬, 共享內存被分為32個同樣大小的內存模型, 它們被稱為存儲體, 它們可以被同時訪問

這里和線程束大小32相同

此造就了以下特點:

如果通過線程束發布共享內存加載或存儲操作, 且在每個存儲體上只訪問不多于一個的內存地址, 那么該操作可由一個內存事務來完成. 否則, 該操作由多個內存事務來完成, 這樣就降低了內存帶寬的利用率

存儲體沖突:

上頭剛說到的問題

當多個地址請求落在相同的內存存儲體中時, 就會發生存儲體沖突, 這會導致請求被重復執行
硬件會將存儲體沖突的請求分割到盡可能多的獨立的無沖突事務中, 有效帶寬的降低是由一個等同于所需的獨立內存事務數量的因素導致的

和上一章講到的相似, 當線程束發出共享內存請求時, 有以下3種典型的模式：

• 并行訪問：多個地址訪問多個存儲體

• 串行訪問：多個地址訪問同一個存儲體
  如線程束中的32個線程都訪問同一個存儲體中的不同地址, 將需要32個內存事務, 所消耗的時間也是單一請求的32倍

• 廣播訪問：單一地址讀取單一存儲體
  此僅適用多個線程訪問一個存儲體中的同一個地址, 此時不發生存儲體沖突

  此種情況雖然僅需要一個內存事務, 但是由于訪問的數據量很小, 所以帶寬的利用度很差

訪問模式:

共享內存存儲體的寬度規定了共享內存地址與共享內存存儲體的對應關系

• 計算能力2.x的設備中為4字節（32位）

• 計算能力3.x的設備中為8字節（64位）

對于Fermi設備, 存儲體的寬度是32位并且有32個存儲體. 每個存儲體在每兩個時鐘周期內都有32位的帶寬. 連續的32位字映射到連續的存儲體中

使用共享內存的字節地址計算出存儲體的索引:
$$存儲體索引=\frac{字節地址}{字節數/存儲體}\%32個存儲體$$
也就是說, 存儲體在共享內存中的分布是這樣的:

這樣的布局是為了相鄰的字被分配到不同的存儲體中, 在線程塊中的線程執行連續訪問時, 能分配到不同的存儲體中, 以最大限度的提高線程束中可能的并發訪問數量

同樣的, 同一個線程束中的多個線程對同一個地址訪問時會使用廣播, 并不會引發存儲體沖突, 但如果是寫入操作的話則需要排隊, 并且順序未知

對于Kepler架構而言:

其同樣有32個存儲體, 但是其有32位和64位兩種地址模式, 后者顯然能更好的降低存儲體沖突的概率(總是產生相同或更少的存儲體沖突)

而在32位模式下, 64位的存儲體被分割成倆:

如圖, 同時訪問Bank0 的0和32索引單元并不會引發存儲體沖突, 因為他們屬于一個存儲體中連續的64位, 在一個時鐘周期中可以同時傳送

但是, 如果訪問的不是連續的64位, 如以下兩種情況, 則會導致存儲體沖突:

內存填充:

內存填充是避免存儲體沖突的一種方法

假設有5個存儲體, 其中的數據以如下排布:

如果要訪問bank0的不同地址, 則會發生5項內存沖突
而內存填充的思想就是通過額外的字, 將原本儲存在統一存儲體中的數據分散到不同的存儲體中

如圖, 在N=5個元素之后添加一個額外的字, 其元素排布將變成如下:

內存填充的思想&優點:

• 對于行, 在進行行主序讀取時, 仍能保證不發生存儲體沖突
• 對于列, 由于打亂了原先在同一列中的元素排布, 所以對于列主序讀取, 也能做到不發生存儲體沖突

綜上, 就是無論使用行主序 或 列主序, 都不會發生存儲體沖突

內存填充的缺點:

• 添加了額外的無用數據, 將使線程塊可用的總共享內存減少

• 由于其涉及到存儲體的具體數量, 所以不同架構的顯卡中應用內存填充將使用不同的策略
  需要修改寫入和訪問的索引

  不修改會導致應用到不同架構上可能出現存儲體沖突

訪問模式配置:

之前說到Kepler架構有兩種共享內存工作模式, 默認是在4字節(32位)

使用此函數可以在運行時查看:

cudaError_t cudaDeviceGetSharedMemConfig(enum cudaSharedMemConfig *pConfig);

使用此函數進行共享內存工作模式的配置:

cudaError_t cudaDeviceSetSharedMemConfig(enum cudaSharedMemConfig config);

一個大的存儲體可能為共享內存訪問產生更高的帶寬, 但是可能會導致更多的存儲體沖突

根據情況設置

配置共享內存:

CUDA為配置一級緩存和共享內存的大小提供了兩種方法：

• 按設備進行配置

• 按核函數進行配置

設備全局配置:

使用以下函數配置一級緩存和共享內存的大小:

cudaError_t cudaDeviceSetCacheConfig(enum cudaFuncCache cacheConfig);

支持的參數如下:

一般有兩個配置策略:

• 當核函數使用較多的共享內存時, 傾向于更多的共享內存

• 當核函數使用更多的寄存器時, 傾向于更多的一級緩存

核函數單獨配置:

cudaError_t cudaFuncSetCacheConfig(const void *func, enum cudaFuncCache cacheConfig);

參數與上頭相同

其中func是指定配置的核函數的指針

對于每個核函數, 僅需要調用一次配置函數即可

同步:

既然是并行計算語言, 必然會有同步機制, CUDA提供幾個運行時函數來執行塊內同步:

這里又開始迷惑HAPI翻譯了, 翻譯的爛就算了, 譯者還不給原文名

• 障礙
  塊內的所有線程都到達barrier點后才會繼續執行

• 內存柵欄
  所有調用的線程必須等到全部內存修改對其余調用線程可見時才能繼續執行

后者的理解需要先了解一下CUDA的弱排序內存模型

這是什么鬼翻譯, 這里比較好的翻譯應該是弱內存順序模型或弱內存模型 Weak Memory Models

理解了準確意思即可

弱排序內存模型

GPU線程在不同內存（如共享內存、全局內存、鎖頁主機內存或對等設備的內存）中寫入數據的順序, 不一定和這些數據在源代碼中訪問的順序相同

一個線程的寫入順序對其他線程可見時, 它可能和寫操作被執行的實際順序不一致
同樣, 如果指令之間是相互獨立的, 線程從不同內存中讀取數據的順序和讀指令在程序中出現的順序不一定相同

為了顯式地強制程序以一個確切的順序執行, 必須在應用程序代碼中插入內存柵欄和障礙
這是保證與其他線程共享資源的核函數行為正確的唯一途徑

顯式障礙:

在核函數中, 通過使用以下函數來設置障礙:

void __syncthreads();

它要求塊中的線程必須等待直到所有線程都到達該點
并確保在障礙點之前, 被這些線程訪問的所有全局和共享內存對同一塊中的所有線程都可見

所以__syncthreads通常用于協調同一塊中線程間的通信, 如訪問同一地址的內存空間時可能產生的問題(寫后讀、讀后寫、寫后寫)

使用這玩意時還需要注意死鎖問題:

當線程塊中的線程走不同的程序路徑時, 在分支中使用__syncthreads()可能導致部分線程永遠無法到達同步點而形成死鎖:

內存柵欄:

這里需要簡單了解一下并發中的可見性 & 有序性:

• 緩存導致了可見性問題
• 編譯優化導致了有序性問題

可以理解可見性就是:
一個線程修改了內存數據, 其他同步范圍內的線程都能夠正確訪問到這個被修改后的數值, 而非是修改前的數值
(緩存問題會導致部分修改的數值僅在緩存中, 而并沒有同步到其他線程可見的地步, 這個在Java并發編程中有涉及)

內存柵欄的功能可確保柵欄前的任何內存寫操作對柵欄后的其他線程都是可見的

根據所需范圍, 有3種內存柵欄：塊、網格或系統, 分別對應三種柵欄函數:

void __threadfence_block(); //線程塊級別 void __threadfence(); //網格級別 void __threadfence_system(); //系統級別

其都是在不同范圍內保證所有寫操作對范圍內的所有線程可見

而一個比較特殊的是__threadfence_block()塊內內存同步, 書里是這樣講的:

內存柵欄不執行任何線程同步, 所以對于一個塊中的所有線程來說, 沒有必要實際執行這個指令

又開始謎語人了, 之前哪里有說過?

這里先放著

volatile修飾符:

C++中的volatile修飾符也能用在CUDA中, 使用后編譯器會取消對該變量的緩存優化, 每次改變都會執行內存同步( 即不進行數據緩存, 而直接寫回到內存中)

5.2 共享內存的數據布局:

為了全面了解如何有效地使用共享內存, 本節將使用共享內存研究幾個簡單的例子, 其中包括下列主題：

• 方陣與矩陣數組

• 行主序與列主序訪問

• 靜態與動態共享內存的聲明

• 文件范圍與內核范圍的共享內存

• 內存填充與無內存填充

當使用共享內存設計核函數時, 重點應放在以下兩個概念上：

• 跨內存存儲體映射數據元素

• 從線程索引到共享內存偏移的映射

當這些概念了然于心時, 就可以設計一個高效的核函數了, 它可以避免存儲體沖突, 并充分利用共享內存的優勢

方形共享內存:

方形共享內存說白了就是方形排布的共享內存:

可以直接使用一個二維線程塊來訪問, 分為行主序 & 列主序

第一種是行主序, 線程塊的行對應著內存塊的行
第二種則相反

很容易能看到, 第一種行主序的方法能呈現出更好的性能和更少的存儲體沖突:
由于線程束是按x優先進行劃分的, 所以鄰近threadIdx.x 的線程會被劃分到同一個線程束中, 這樣訪問共享內存時, 線程束中的每個線程都能訪問到不同的存儲體

行主序訪問 & 列主序訪問:

這里就是實踐行主序 & 列主序的區別, 比較性能差異

行主序訪問:

此時沒有存儲體沖突

列主序訪問:

此時會導致大量的存儲體沖突

使用nvprof能很好的看到性能差異:

書里使用的是K40c

執行時間的差異:

存儲體沖突的差異:

在nvprof中使用以下兩個指標檢測存儲體沖突:

行主序寫 & 列主序讀:

下面的核函數實現了共享內存中按行主序寫入和按列主序讀取

所以這個例子有啥意義, 這不是猜都能猜到的么

動態共享內存:

這里使用上頭講到的動態內存

動態共享內存可以在核函數之外聲明, 其作用域將是整個文件
也可以在核函數之內聲明, 其作用域將僅限于核函數

例程:

核函數中按行主序寫入, 按列主序讀取

nvprof結果:

所以表明了使用動態共享內存也會存在相同的問題

填充動態聲明的共享內存:

這里是對動態共享內存執行內存填充

填充動態聲明的共享內存數組更加復雜

因為在以上核函數中用于存儲數據的全局內存小于填充的共享內存, 所以需要3個索引：一個索引用于按照行主序寫入共享內存, 一個索引用于按照列主序讀取共享內存, 一個索引用于未填充的全局內存的合并訪問

這些結果和填充靜態聲明的共享內存是一致的

所以這里證明的是, 無論是靜態共享內存還是動態共享內存都能被有效的填充

方形共享內存內核性能的比較:

到目前為止, 從所有執行過的內核運行時間可以看出：

• 使用填充的內核可提高性能, 因為它減少了存儲體沖突

• 帶有動態聲明共享內存的內核增加了少量的消耗

矩形共享內存:

這一部分的行文邏輯基本上和上一節相同, 討論共享內存的幾個點, 只不過吧上頭的方陣替換為了矩陣

矩形共享內存是一個更普遍的二維共享內存, 他與方形共享內存的區別就是行列數不等 ( 矩陣 & 方陣的區別)

本部分的所有核函數調用都使用以下執行配置:

行主序訪問 & 列主序訪問:

這里的結果 & 結論基本上與上頭的方陣相同
所以簡單看下就好

就是將上頭的方陣替換為了矩陣內存, 并執行內存轉置操作:

這里使用的應該是16個數據, 而并非之前方陣的32個, 所以數據不同, 但是結論是相同的

共享內存的存儲和加載請求, 由setRowReadRow核函數中的一個事務完成. 同樣的請求在setColReadCol函數中由8個事務完成. Kepler K40的存儲體寬度是8個字, 一列16個4字節的數據元素被安排到8個存儲體中, 如圖5-6所示, 因此, 該操作有一個8路沖突

行主序寫 & 列主序讀:

使用共享內存進行矩陣轉置的核函數. 通過最大化低延遲的加載和存儲來提高性能, 并合并全局內存訪問

內核有3個內存操作：

• 寫入每個線程束的共享內存行, 以避免存儲體沖突

• 讀取每個線程束中的共享內存列, 以完成矩陣轉置

• 使用合并訪問(上一章講到的)寫入每個線程束的全局內存行

該存儲操作是無沖突的, 加載操作報告了一個8路沖突

與預期相同
store時是行主序, load時是列主序

動態共享內存:

還是緊接著上頭的例子進行修改, 將其中的靜態內存改為動態內存, 繼續實現矩陣轉置

結果與使用靜態內存相同

所以結論就是:

動態分配共享內存不會影響存儲體沖突

填充靜態共享內存:

在前面的宏中若將填充數據元素的數量從2改到1, 則nvprof報告有兩個事務完成共享內存的加載操作, 即發生一個雙向存儲體沖突

所以結論是:

填充的元素個數與行列數是有關系的, 數量不當仍將導致存儲體沖突

填充動態共享內存:

大致套路和靜態共享內存相同:

結論就是:

動態內存的填充比靜態內存的仍然要復雜
其有專門的計算index 的代碼

矩形共享內存內核性能的比較:

在一般情況下, 和上一節說到的一樣:

• 核函數使用共享內存填充消除存儲體沖突以提高性能
• 使用動態共享內存的核函數會顯示有少量的消耗

5.3 減少全局內存訪問:

使用共享內存的主要原因之一是要緩存片上的數據, 從而減少核函數中全局內存訪問的次數

在本節中, 將重新使用第三章中的并行歸約核函數, 但是這里使用共享內存作為可編程管理緩存以減少全局內存的訪問

使用共享內存的并行歸約:

首先是一個僅使用全局內存的歸約核函數, 作為所有核函數的起點與性能的基點:

而后是帶有共享內存的全局內存操作的歸約函數

此核函數就是利用共享內存將全局內存中的數據進行了緩存, 而后的歸約都只在共享內存中進行(替代了直接讀寫全局內存的操作)

二者對比如下:

使用共享內存的核函數比只使用全局內存的核函數快了1.84倍

使用nvprof的倆參數查看全局內存加載&存儲事務:

使用展開的并行歸約

這里就是在上一節的例子中加上之前的循環展開方法:

以下內核展開了4個線程塊, 即每個線程處理來自于4個數據塊的數據元素
可預期的效果是:

• 通過在每個線程中提供更多的并行I/O, 增加全局內存的吞吐量

• 全局內存存儲事務減少了1/4

• 整體內核性能的提升

qs, 加載量保持不變, 但是由于是4展開, 所以存儲量下降(原先需要存儲多次的過程被壓縮到了一個線程中進行)

使用動態共享內存的并行歸約

這里一筆帶過, 直接上結論;

用動態分配共享內存實現的核函數和用靜態分配共享內存實現的核函數之間沒有顯著的差異

有效帶寬:

由于歸約核函數是受內存帶寬約束的, 所以評估它們時所使用的適當的性能指標是有效帶寬

有效帶寬是在核函數的完整執行時間內I/O的數量（以字節為單位）

對于內存約束的應用程序, 有效帶寬是一個估算實際帶寬利用率的很好的指標

計算公式:
$$有效帶寬(GB/s)=\frac{(讀字節數+寫字節數)}{運行時間?10^9}$$

以下是前頭的4個函數的有效帶寬:

顯然, 可以通過展開塊來獲得有效帶寬的顯著改進
每個線程運行中同時有多個請求, 會導致內存總線高飽和

5.4 合并的全局內存訪問:

使用共享內存也能幫助避免產生未合并的全局內存訪問

之前的矩陣轉置核函數中, 讀操作是合并的, 但寫操作是交叉訪問的

在使用共享內存之后, 可以將共享內存作為緩存, 先在共享內存中進行交叉訪問, 利用共享內存的低延時&高帶寬降低時間損耗, 完成后在整塊寫回到全局內存中, 以實現合并寫入

在本章前面的部分, 測試了一個矩陣轉置核函數, 該核函數使用單個線程塊對共享內存中的矩陣行進行寫入, 并讀取共享內存中的矩陣列

在本節中, 將擴展該核函數, 具體方法是使用多個線程塊對基于交叉的全局內存訪問重新排序到合并訪問

基準轉置核函數:

和上一節的行文邏輯相同, 先確定一個性能比較的基準

下面的核函數是一個僅使用全局內存的矩陣轉置的樸素實現
其中, 全局內存讀操作在線程束內是被合并的, 而全局內存寫操作在相鄰線程間是交叉訪問的

而后這個核函數將作為優化的性能上限
其中讀寫操作都將被合并, 仍執行相同數量的IO

后頭測試用的矩陣大小將使用2¹² * 2¹², 線程塊大小為32*16

基準核函數的運行結果:

副本內核比樸素內核快了將近3倍

由于樸素內核寫入全局內存, 使其帶有了4096個元素的跨度, 所以一個單一線程束的存儲內存操作是由32個全局內存事務完成的. 可以使用以下nvprof指標來確認這一點

使用共享內存的矩陣轉置:

為了避免交叉全局內存訪問, 可以使用二維共享內存來緩存原始矩陣的數據

實現的核函數:
可以看做是上一節中的setRowReadCol的擴展, 前者使用的單一線程塊, 而后者將其擴展為了使用多個線程塊和數據塊

核函數的程序步驟:

kerneltransposeSmem函數可被分解為以下幾個步驟：

線程束執行合并讀取一行, 該行存儲在全局內存中的原始矩陣塊中.

然后, 該線程束按行主序將該數據寫入共享內存中, 因此, 這個寫操作沒有存儲體沖突.

因為線程塊的讀/寫操作是同步的, 所以會有一個填滿全局內存數據的二維共享內存數組.

該線程束從二維共享內存數組中讀取一列. 由于共享內存沒有被填充, 所以會發生存儲體沖突.

然后該線程束執行數據的合并寫入操作, 將其寫入到全局內存的轉置矩陣中的某行

核函數具體的實現就暫且略過了(詳見書里), 這里來看其實現的特點:

全局內存的讀取是合并的

共享內存的寫入沒有發生存儲體沖突

共享內存的讀取發生存儲體沖突
是按列讀取

全局內存的寫入是合并的

性能對比:

全局內存存儲的重復數量從32減少到2

這是由于轉置塊中的塊寬為16, 所以線程束前半部分的寫操作和線程束后半部分的寫操作間隔了4080
因此線程束的寫入請求是有兩個事務完成的

將線程塊大小更改到32×32會把重復次數減少到1, 但是前者(32*16)將顯現出更多的并行性

顯然, 讀取二維共享內存數組中的一列會產生存儲體沖突

使用填充共享內存的矩陣轉置:

這里就是應用之前的填充

通過給二維共享內存數組tile中的每一行添加列填充, 可以將原矩陣相同列中的數據元素均勻地劃分到共享內存存儲體中

需要填充的列數取決于設備的計算能力和線程塊的大小
對于一個大小為32×16的線程塊被測試內核來說, 在Tesla K40中必須增加兩列填充, 在Tesla M2090中必須增加一列填充

修改之前的共享內存聲明如下:

使用展開的矩陣轉置:

就是在添加一個循環展開

下面的核函數展開兩個數據塊的同時處理：每個線程現在轉置了被一個數據塊跨越的兩個數據元素
這種轉化的目標是通過創造更多的同時加載和存儲以提高設備內存帶寬利用率

核函數的其他詳細實現直接去看書, 其特點都在上頭的這個圖里

增大并行性:

這里是通過調整線程塊的維度來提升性能

塊大小為16×16時展示出了最好的性能, 因為它有更多的并發線程塊, 從而有最好的設備并行性

表5-7總結了在Tesla K40中從transposeSmemUnrollPadDyn函數上獲得全局內存吞吐量和共享內存存儲體沖突的nvprof結果. 雖然線程塊配置為32×16時最大程度地減少了存儲體沖突, 但線程塊配置為16×16時最大程度地增加了全局內存吞吐量

由此, 可以得出結論, 與共享內存吞吐量相比, 內核受到全局內存吞吐量的約束更多

5.5 常量內存:

常量內存是一種專用的內存
其對內核代碼而言是只讀的，但它對主機而言既是可讀又是可寫的

常量內存位于設備的DRAM上（和全局內存一樣），并且有一個專用的片上緩存
每個SM常量內存緩存大小的限制為64KB

與其他類型的內存不同, 常量內存有一個最優訪問模式:

• 當線程束中的so哦有線程都訪問相同的位置, 此時訪問模式是最優的
• 如果線程束訪問不同的地址, 則需要串行訪問

所以常量內存的讀取成本與線程束中讀取的地址數量息息相關

使用__constant__聲明一個常量變量

由于常量內存在設備上只讀, 所以必須在主機上進行初始化:

cudaError_t cudaMemcpyToSymbol(const void *symbol, const void *src, size_t count, size_t offset __dv(0), enum cudaMemcpyKind kind __dv(cudaMemcpyHostToDevice));

cudaMemcpyToSymbol函數將src指向的數據復制到設備上由symbol指定的常量內存中。枚舉變量kind指定了傳輸方向，默認情況下，kind是cudaMemcpyHostToDevice。

常量內存的幾個特點

• 生存期與應用程序相同
• 對網格內的所有線程可見
• 主機也可以直接訪問

使用常量內存實現一維模板:

又開始了, 神奇的翻譯

這里介紹了一個莫名其妙的九點模板(搜都搜不到, 什么HAPI翻譯 )

這里重點不是理解這個九點模板, 而是考慮到他的使用場景:

9個x作為輸入, 一個輸出

公式中有c0, c1, c2, c3 四個常數, 并且每個線程都需要
所以可以用廣播式的訪問模式, 線程束中的每個線程同時引用相同的常量內存地址

所實現的核函數

代碼實現具體看書, 這里重點關注他的常量內存的使用:

與只讀緩存的比較:

這里要講到Kepler架構中添加的獨立只讀數據緩存:

這里需要注意, 好像僅僅是Kepler架構中有這玩意, 在后續的架構中并沒有這玩意:

可以看到, Kepler的SM中僅有48KB的只讀緩存
所以, 制度緩存在分散讀取方面比一級緩存更好, 當線程束中的線程都讀取相同地址時, 不應使用只讀緩存

只讀緩存的使用:

當通過只讀緩存訪問全局內存時，需要向編譯器指出在內核的持續時間里數據是只讀的

• 使用內部函數__ldg

• 全局內存的限定指針

通常選用第一種__ldg方法
尤其是在只讀緩存機制需要更多顯式控制的情況下，或者在代碼非常復雜以至于編譯器無法檢測到只讀緩存的使用是否是安全的情況下

與常量內存的對比:

• 常量緩存加載的數據必須是少量的, 并且需要訪問的一致性才能獲得較好的性能
• 制度緩存加載的數據可以是比較大的, 而且能在一個非統一的模式下進行訪問

所以可以得出以下結論:

• 常量緩存在讀取同一地址的數據中可以更好的性能
• 只讀緩存更適合于分散讀取

核函數實現:
此核函數和上頭的唯一區別就是函數聲明部分

在Tesla K40上，使用nvprof測試得出的以下結果表明，對此應用程序使用只讀內存時其性能實際上會降低。這是由于coef數組使用了廣播訪問模式，相比于只讀緩存，該模式更適合于常量內存：

5.6 線程束洗牌指令:

從用Kepler系列的GPU（計算能力為3.0或更高）開始，洗牌指令（shuffle instruction）作為一種機制被加入其中，只要兩個線程在相同的線程束中，那么就允許這兩個線程直接讀取另一個線程的寄存器

洗牌指令比共享內存有更低的延遲，并且該指令在執行數據交換時不消耗額外的內存

首先介紹一下束內線程（lane）的概念
簡單來說, 一個束內線程指的是線程束內的單一線程, 每個束內線程都有唯一的束內線程索引, 為[0,31], 但沒有單獨存儲束內線程索引的變量, 而是通過塊內線程索引threadIdx.x計算得到:

線程束洗牌指令的不同形式:

有兩組洗牌指令：一組用于整型變量，另一組用于浮點型變量。每組有4種形式的洗牌指令

這里僅介紹整型變量的4中洗牌指令, 對于單精度浮點的洗牌則與整型的完全相同

廣播:

在線程束內交換整型變量，其基本函數標記如下：

__CUDA_FP16_DECL__ __DEPRECATED__(__WSB_DEPRECATION_MESSAGE(__shfl))__half2 __shfl(const __half2 var, const int delta, const int width = warpSize)

書里的是這個形式:

這個函數能使線程束中的每個線程都可以直接從一個特定的線程中獲取某個值

線程束內所有活躍的線程都同時產生此操作，這將導致每個線程中有4字節數據的移動

參數解釋:

• 返回值:
  其他線程從root線程獲得到的值

• var
  root線程共享出來的值

• srcLane:
  Lane代表的是束內線程, 所以可知這玩意是用來指定束內線程的

• width:
  洗牌分段
  默認=warpSize=32 , 此時洗牌操作的作用范圍是整個線程束
  但是通過手動設置值可以調的更細, 使每段包含有width個線程, 并且每段上指定獨立的洗牌操作

  此時srcLane使用的線程ID與束內線程ID不同, 其使用如下公式計算:

  那么線程0～15將從線程3接收x的值，線程16～31將從線程19接收x的值（在線程束的前16個線程中其偏移量為3）

所以可知, 這個操作有點類似于MPI中的廣播

但是這里是吧__shlf中的參數寫死了的情況

如果使用動態參數, 可以得到下一節中的循環交換的效果:

復制:

__CUDA_FP16_DECL__ __DEPRECATED__(__WSB_DEPRECATION_MESSAGE(__shfl_up))__half2 __shfl_up(const __half2 var, const unsigned int delta, const int width = warpSize);__CUDA_FP16_DECL__ __DEPRECATED__(__WSB_DEPRECATION_MESSAGE(__shfl_down))__half2 __shfl_down(const __half2 var, const unsigned int delta, const int width = warpSize);

參數介紹:

• delta
  線程束偏移量

其他參數都和上頭的廣播相似

偏移量這個就是下圖所展現的

而這兩個函數的區別就是方向不同:

• up向高index方向復制
• down向低index方向復制

并且從圖中也可以看到, 需要注意的是:
頭尾部分的線程束的值保持不變, 并沒有首尾相接的交換

交叉交換:

__CUDA_FP16_DECL__ __DEPRECATED__(__WSB_DEPRECATION_MESSAGE(__shfl_xor))__half2 __shfl_xor(const __half2 var, const int delta, const int width = warpSize)

不具體闡述了

線程束內的共享數據

在本節中，會介紹幾個有關線程束洗牌指令的例子，并說明線程束洗牌指令的優點

洗牌指令將被應用到以下3種整數變量類型中：

• 標量變量

• 數組

• 向量型變量

下頭就全是例子, 就中間的循環移動需要看看以外, 其他的都一筆帶過

線程束內的值廣播:

這里就是對上一節講到的洗牌指令進行應用:

執行結果:

線程束內上移:

這里還是應用

線程束內下移:

線程束內環繞移動:

這里實現的就是上頭所沒有實現的環繞式移動, 即首尾相接的__shfl_up 或 __shfl_down

• 正偏移量為左移
• 負偏移量為右移

蝴蝶(交叉)交換:

交換數組值:

這個其實就是__shfl_xor()的花式應用

在下面的內核中，每個線程都有一個寄存器數組value，其大小是SEGM。每個線程從全局內存d_in中讀取數據塊到value中，使用由掩碼確定的相鄰線程交換該塊，然后將接收到的數據寫回到全局內存數組d_out中

使用數組索引交換數值:

這里實現的是在兩個線程各自的數組中以不同的偏移量交換它們之間的元素，需要有基于洗牌指令的交換函數

本部分先放著

布爾變量pred被用于識別第一個調用的線程，它是交換數據的一對線程。要交換的數據元素是由第一個線程的firstIdx和第二個線程的secondIdx偏移標識的。第一個調用線程通過交換firstIdx和secondIdx中的元素開始，但此操作僅限于本地數組。然后在兩線程間的secondIdx位置執行蝴蝶交換。最后，第一個線程交換接收自secondIdx返回到firstIdx的元素

使用線程束洗牌指令的并行歸約

這里就是將前頭的洗牌指令應用到之前的歸約例子中

基本思路非常簡單，它包括3個層面的歸約：

• 線程束級歸約

• 線程塊級歸約

• 網格級歸約

詳細的解釋可以看書:

這里直接看結果:

用洗牌指令實現線程束級并行歸約獲得了1.42倍的加速

5.7 總結:

為了獲得最大的應用性能，需要有一個能顯式管理的內存層次結構。在C語言中，沒有直接控制數據移動的方式。在本章中，介紹了不同CUDA內存層次結構類型，如共享內存、常量內存和只讀緩存。介紹了當從共享內存中引入或刪除數據時如何顯式控制以顯著提高其性能。還介紹了常量內存和只讀緩存的行為，以及如何最有效地使用它們。

共享內存可以被聲明為一維或二維數組，它能為每個程序提供一個簡單的邏輯視圖。物理上，共享內存是一維的，并能通過32個存儲體進行訪問。避免存儲體沖突是在共享內存應用優化過程中一個重要的因素。共享內存被分配在所有常駐線程塊中，因此，它是一個關鍵資源，可能會限制內核占用率。

在內核中使用共享內存有兩個主要原因：一個是用于緩存片上數據并且減少全局內存訪問量；另一個是傳輸共享內存中數據的安排方式，避免非合并的全局內存訪問。

常量內存對只讀數據進行了優化，這些數據每次都將數據廣播到許多線程中。常量內存也使用自己的SM緩存，防止常量內存的讀操作通過一級緩存干擾全局內存的訪問。因此，對合適的數據使用常量內存，不僅可優化特定項目的訪問，還可能提高整體全局內存吞吐量。

只讀紋理緩存提供了常量內存的替代方案，該方案優化了數據的分散讀取。只讀緩存訪問全局內存中的數據，但它使用一個獨立的內存訪問流水線和獨立的緩存，以使SM可以訪問數據。因此，只讀緩存共享了常量內存的許多好處，同時對不同的訪問模式也進行了優化。

洗牌指令是線程束級的內部功能，能使線程束中的線程彼此之間快速直接地共享數據。洗牌指令具有比共享內存更低的延遲，并且不需要分配額外的資源。使用洗牌指令可以減少內核中線程束同步優化的數目。然而，在許多情況下，洗牌指令不是共享內存的替代品，因為共享內存在整個線程塊中都可見。

本章對一些有特殊用途的內存類型進行了深度了解。雖然這些內存類型比全局內存使用得少，但是適當地使用它們可以提高帶寬利用率，降低整體的內存延遲。如果你正在研究優化的因素，那么牢記共享內存、常量內存、只讀緩存和洗牌指令都是非常重要的。

總結

以上是生活随笔為你收集整理的CUDA编程第五章: 共享内存常量内存的全部內容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網站內容還不錯，歡迎將生活随笔推薦給好友。