前言:

本章內(nèi)容：

• 了解數(shù)據(jù)在共享內(nèi)存中是如何被安排的

• 掌握從二維共享內(nèi)存到線性全局內(nèi)存的索引轉(zhuǎn)換

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

• 在共享內(nèi)存中緩存數(shù)據(jù)以減少對(duì)全局內(nèi)存的訪問

• 使用共享內(nèi)存避免非合并全局內(nèi)存的訪問

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

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

在前面的章節(jié)中, 已經(jīng)介紹了幾種全局內(nèi)存的訪問模式. 通過安排全局內(nèi)存訪問模式, 我們學(xué)會(huì)了如何實(shí)現(xiàn)良好的性能并且避免了浪費(fèi)事務(wù). 未對(duì)齊的內(nèi)存訪問是沒有問題的, 因?yàn)楝F(xiàn)代的GPU硬件都有一級(jí)緩存, 但在跨全局內(nèi)存的非合并內(nèi)存訪問, 仍然會(huì)導(dǎo)致帶寬利用率不會(huì)達(dá)到最佳標(biāo)準(zhǔn). 根據(jù)算法性質(zhì)和相應(yīng)的訪問模式, 非合并訪問可能是無法避免的. 然而, 在許多情況下, 使用共享內(nèi)存來提高全局內(nèi)存合并訪問是有可能的. 共享內(nèi)存是許多高性能計(jì)算應(yīng)用程序的關(guān)鍵驅(qū)動(dòng)力.

在本章中, 你將學(xué)習(xí)如何使用共享內(nèi)存進(jìn)行編程、數(shù)據(jù)在共享內(nèi)存中如何被存儲(chǔ)、數(shù)據(jù)元素是怎樣使用不同的訪問模式被映射到內(nèi)存存儲(chǔ)體中的. 還將掌握使用共享內(nèi)存提高核函數(shù)性能的方法.

5.1 CUDA共享內(nèi)存概述:

GPU中有兩種類型的內(nèi)存：

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

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

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

• 塊內(nèi)線程通信的通道

• 用于全局內(nèi)存數(shù)據(jù)的可編程管理的緩存

• 高速暫存存儲(chǔ)器, 用于轉(zhuǎn)換數(shù)據(jù)以優(yōu)化全局內(nèi)存訪問模式

共享內(nèi)存:

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

共享內(nèi)存（shared memory, SMEM）其特點(diǎn):

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

以Kepler核心的SM為例:

SM基本存儲(chǔ)順序:

共享內(nèi)存訪問事物:

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

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

可編程管理的緩存:

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

而共享內(nèi)存是一個(gè)可編程管理的緩存, 所以可以通過在數(shù)據(jù)布局上提供更多的細(xì)粒度控制和改善片上數(shù)據(jù)的移動(dòng), 使得對(duì)優(yōu)化應(yīng)用程序代碼變得更簡單

共享內(nèi)存的分配:

共享內(nèi)存使用__shared__修飾符進(jìn)行聲明

如:

__shared__ float tile[size_y][size_x];

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

extern __shared__ int tile[];

在每個(gè)核函數(shù)被調(diào)用時(shí), 需要?jiǎng)討B(tài)分配共享內(nèi)存 這部分操作在主機(jī)端進(jìn)行

即在<<<>>>后頭多加一個(gè)參數(shù), 注意這里是以字節(jié)為單位

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

共享內(nèi)存存儲(chǔ)體和訪問模式

優(yōu)化內(nèi)存性能時(shí)要度量的兩個(gè)關(guān)鍵屬性是：延遲和帶寬

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

內(nèi)存存儲(chǔ)體:

為了獲得高內(nèi)存帶寬, 共享內(nèi)存被分為32個(gè)同樣大小的內(nèi)存模型, 它們被稱為存儲(chǔ)體, 它們可以被同時(shí)訪問

這里和線程束大小32相同

此造就了以下特點(diǎn):

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

存儲(chǔ)體沖突:

上頭剛說到的問題

當(dāng)多個(gè)地址請(qǐng)求落在相同的內(nèi)存存儲(chǔ)體中時(shí), 就會(huì)發(fā)生存儲(chǔ)體沖突, 這會(huì)導(dǎo)致請(qǐng)求被重復(fù)執(zhí)行
硬件會(huì)將存儲(chǔ)體沖突的請(qǐng)求分割到盡可能多的獨(dú)立的無沖突事務(wù)中, 有效帶寬的降低是由一個(gè)等同于所需的獨(dú)立內(nèi)存事務(wù)數(shù)量的因素導(dǎo)致的

和上一章講到的相似, 當(dāng)線程束發(fā)出共享內(nèi)存請(qǐng)求時(shí), 有以下3種典型的模式：

• 并行訪問：多個(gè)地址訪問多個(gè)存儲(chǔ)體

• 串行訪問：多個(gè)地址訪問同一個(gè)存儲(chǔ)體
  如線程束中的32個(gè)線程都訪問同一個(gè)存儲(chǔ)體中的不同地址, 將需要32個(gè)內(nèi)存事務(wù), 所消耗的時(shí)間也是單一請(qǐng)求的32倍

• 廣播訪問：單一地址讀取單一存儲(chǔ)體
  此僅適用多個(gè)線程訪問一個(gè)存儲(chǔ)體中的同一個(gè)地址, 此時(shí)不發(fā)生存儲(chǔ)體沖突

  此種情況雖然僅需要一個(gè)內(nèi)存事務(wù), 但是由于訪問的數(shù)據(jù)量很小, 所以帶寬的利用度很差

訪問模式:

共享內(nèi)存存儲(chǔ)體的寬度規(guī)定了共享內(nèi)存地址與共享內(nèi)存存儲(chǔ)體的對(duì)應(yīng)關(guān)系

• 計(jì)算能力2.x的設(shè)備中為4字節(jié)（32位）

• 計(jì)算能力3.x的設(shè)備中為8字節(jié)（64位）

對(duì)于Fermi設(shè)備, 存儲(chǔ)體的寬度是32位并且有32個(gè)存儲(chǔ)體. 每個(gè)存儲(chǔ)體在每兩個(gè)時(shí)鐘周期內(nèi)都有32位的帶寬. 連續(xù)的32位字映射到連續(xù)的存儲(chǔ)體中

使用共享內(nèi)存的字節(jié)地址計(jì)算出存儲(chǔ)體的索引:
$$存儲(chǔ)體索引=\frac{字節(jié)地址}{字節(jié)數(shù)/存儲(chǔ)體}\%32個(gè)存儲(chǔ)體$$
也就是說, 存儲(chǔ)體在共享內(nèi)存中的分布是這樣的:

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

同樣的, 同一個(gè)線程束中的多個(gè)線程對(duì)同一個(gè)地址訪問時(shí)會(huì)使用廣播, 并不會(huì)引發(fā)存儲(chǔ)體沖突, 但如果是寫入操作的話則需要排隊(duì), 并且順序未知

對(duì)于Kepler架構(gòu)而言:

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

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

如圖, 同時(shí)訪問Bank0 的0和32索引單元并不會(huì)引發(fā)存儲(chǔ)體沖突, 因?yàn)樗麄儗儆谝粋€(gè)存儲(chǔ)體中連續(xù)的64位, 在一個(gè)時(shí)鐘周期中可以同時(shí)傳送

但是, 如果訪問的不是連續(xù)的64位, 如以下兩種情況, 則會(huì)導(dǎo)致存儲(chǔ)體沖突:

內(nèi)存填充:

內(nèi)存填充是避免存儲(chǔ)體沖突的一種方法

假設(shè)有5個(gè)存儲(chǔ)體, 其中的數(shù)據(jù)以如下排布:

如果要訪問bank0的不同地址, 則會(huì)發(fā)生5項(xiàng)內(nèi)存沖突
而內(nèi)存填充的思想就是通過額外的字, 將原本儲(chǔ)存在統(tǒng)一存儲(chǔ)體中的數(shù)據(jù)分散到不同的存儲(chǔ)體中

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

內(nèi)存填充的思想&優(yōu)點(diǎn):

• 對(duì)于行, 在進(jìn)行行主序讀取時(shí), 仍能保證不發(fā)生存儲(chǔ)體沖突
• 對(duì)于列, 由于打亂了原先在同一列中的元素排布, 所以對(duì)于列主序讀取, 也能做到不發(fā)生存儲(chǔ)體沖突

綜上, 就是無論使用行主序 或 列主序, 都不會(huì)發(fā)生存儲(chǔ)體沖突

內(nèi)存填充的缺點(diǎn):

• 添加了額外的無用數(shù)據(jù), 將使線程塊可用的總共享內(nèi)存減少

• 由于其涉及到存儲(chǔ)體的具體數(shù)量, 所以不同架構(gòu)的顯卡中應(yīng)用內(nèi)存填充將使用不同的策略
  需要修改寫入和訪問的索引

  不修改會(huì)導(dǎo)致應(yīng)用到不同架構(gòu)上可能出現(xiàn)存儲(chǔ)體沖突

訪問模式配置:

之前說到Kepler架構(gòu)有兩種共享內(nèi)存工作模式, 默認(rèn)是在4字節(jié)(32位)

使用此函數(shù)可以在運(yùn)行時(shí)查看:

cudaError_t cudaDeviceGetSharedMemConfig(enum cudaSharedMemConfig *pConfig);

使用此函數(shù)進(jìn)行共享內(nèi)存工作模式的配置:

cudaError_t cudaDeviceSetSharedMemConfig(enum cudaSharedMemConfig config);

一個(gè)大的存儲(chǔ)體可能為共享內(nèi)存訪問產(chǎn)生更高的帶寬, 但是可能會(huì)導(dǎo)致更多的存儲(chǔ)體沖突

根據(jù)情況設(shè)置

配置共享內(nèi)存:

CUDA為配置一級(jí)緩存和共享內(nèi)存的大小提供了兩種方法：

• 按設(shè)備進(jìn)行配置

• 按核函數(shù)進(jìn)行配置

設(shè)備全局配置:

使用以下函數(shù)配置一級(jí)緩存和共享內(nèi)存的大小:

cudaError_t cudaDeviceSetCacheConfig(enum cudaFuncCache cacheConfig);

支持的參數(shù)如下:

一般有兩個(gè)配置策略:

• 當(dāng)核函數(shù)使用較多的共享內(nèi)存時(shí), 傾向于更多的共享內(nèi)存

• 當(dāng)核函數(shù)使用更多的寄存器時(shí), 傾向于更多的一級(jí)緩存

核函數(shù)單獨(dú)配置:

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

參數(shù)與上頭相同

其中func是指定配置的核函數(shù)的指針

對(duì)于每個(gè)核函數(shù), 僅需要調(diào)用一次配置函數(shù)即可

同步:

既然是并行計(jì)算語言, 必然會(huì)有同步機(jī)制, CUDA提供幾個(gè)運(yùn)行時(shí)函數(shù)來執(zhí)行塊內(nèi)同步:

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

• 障礙
  塊內(nèi)的所有線程都到達(dá)barrier點(diǎn)后才會(huì)繼續(xù)執(zhí)行

• 內(nèi)存柵欄
  所有調(diào)用的線程必須等到全部內(nèi)存修改對(duì)其余調(diào)用線程可見時(shí)才能繼續(xù)執(zhí)行

后者的理解需要先了解一下CUDA的弱排序內(nèi)存模型

這是什么鬼翻譯, 這里比較好的翻譯應(yīng)該是弱內(nèi)存順序模型或弱內(nèi)存模型 Weak Memory Models

理解了準(zhǔn)確意思即可

弱排序內(nèi)存模型

GPU線程在不同內(nèi)存（如共享內(nèi)存、全局內(nèi)存、鎖頁主機(jī)內(nèi)存或?qū)Φ仍O(shè)備的內(nèi)存）中寫入數(shù)據(jù)的順序, 不一定和這些數(shù)據(jù)在源代碼中訪問的順序相同

一個(gè)線程的寫入順序?qū)ζ渌€程可見時(shí), 它可能和寫操作被執(zhí)行的實(shí)際順序不一致
同樣, 如果指令之間是相互獨(dú)立的, 線程從不同內(nèi)存中讀取數(shù)據(jù)的順序和讀指令在程序中出現(xiàn)的順序不一定相同

為了顯式地強(qiáng)制程序以一個(gè)確切的順序執(zhí)行, 必須在應(yīng)用程序代碼中插入內(nèi)存柵欄和障礙
這是保證與其他線程共享資源的核函數(shù)行為正確的唯一途徑

顯式障礙:

在核函數(shù)中, 通過使用以下函數(shù)來設(shè)置障礙:

void __syncthreads();

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

所以__syncthreads通常用于協(xié)調(diào)同一塊中線程間的通信, 如訪問同一地址的內(nèi)存空間時(shí)可能產(chǎn)生的問題(寫后讀、讀后寫、寫后寫)

使用這玩意時(shí)還需要注意死鎖問題:

當(dāng)線程塊中的線程走不同的程序路徑時(shí), 在分支中使用__syncthreads()可能導(dǎo)致部分線程永遠(yuǎn)無法到達(dá)同步點(diǎn)而形成死鎖:

內(nèi)存柵欄:

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

• 緩存導(dǎo)致了可見性問題
• 編譯優(yōu)化導(dǎo)致了有序性問題

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

內(nèi)存柵欄的功能可確保柵欄前的任何內(nèi)存寫操作對(duì)柵欄后的其他線程都是可見的

根據(jù)所需范圍, 有3種內(nèi)存柵欄：塊、網(wǎng)格或系統(tǒng), 分別對(duì)應(yīng)三種柵欄函數(shù):

void __threadfence_block(); //線程塊級(jí)別 void __threadfence(); //網(wǎng)格級(jí)別 void __threadfence_system(); //系統(tǒng)級(jí)別

其都是在不同范圍內(nèi)保證所有寫操作對(duì)范圍內(nèi)的所有線程可見

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

內(nèi)存柵欄不執(zhí)行任何線程同步, 所以對(duì)于一個(gè)塊中的所有線程來說, 沒有必要實(shí)際執(zhí)行這個(gè)指令

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

這里先放著

volatile修飾符:

C++中的volatile修飾符也能用在CUDA中, 使用后編譯器會(huì)取消對(duì)該變量的緩存優(yōu)化, 每次改變都會(huì)執(zhí)行內(nèi)存同步( 即不進(jìn)行數(shù)據(jù)緩存, 而直接寫回到內(nèi)存中)

5.2 共享內(nèi)存的數(shù)據(jù)布局:

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

• 方陣與矩陣數(shù)組

• 行主序與列主序訪問

• 靜態(tài)與動(dòng)態(tài)共享內(nèi)存的聲明

• 文件范圍與內(nèi)核范圍的共享內(nèi)存

• 內(nèi)存填充與無內(nèi)存填充

當(dāng)使用共享內(nèi)存設(shè)計(jì)核函數(shù)時(shí), 重點(diǎn)應(yīng)放在以下兩個(gè)概念上：

• 跨內(nèi)存存儲(chǔ)體映射數(shù)據(jù)元素

• 從線程索引到共享內(nèi)存偏移的映射

當(dāng)這些概念了然于心時(shí), 就可以設(shè)計(jì)一個(gè)高效的核函數(shù)了, 它可以避免存儲(chǔ)體沖突, 并充分利用共享內(nèi)存的優(yōu)勢(shì)

方形共享內(nèi)存:

方形共享內(nèi)存說白了就是方形排布的共享內(nèi)存:

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

第一種是行主序, 線程塊的行對(duì)應(yīng)著內(nèi)存塊的行
第二種則相反

很容易能看到, 第一種行主序的方法能呈現(xiàn)出更好的性能和更少的存儲(chǔ)體沖突:
由于線程束是按x優(yōu)先進(jìn)行劃分的, 所以鄰近threadIdx.x 的線程會(huì)被劃分到同一個(gè)線程束中, 這樣訪問共享內(nèi)存時(shí), 線程束中的每個(gè)線程都能訪問到不同的存儲(chǔ)體

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

這里就是實(shí)踐行主序 & 列主序的區(qū)別, 比較性能差異

行主序訪問:

此時(shí)沒有存儲(chǔ)體沖突

列主序訪問:

此時(shí)會(huì)導(dǎo)致大量的存儲(chǔ)體沖突

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

書里使用的是K40c

執(zhí)行時(shí)間的差異:

存儲(chǔ)體沖突的差異:

在nvprof中使用以下兩個(gè)指標(biāo)檢測存儲(chǔ)體沖突:

行主序?qū)?& 列主序讀:

下面的核函數(shù)實(shí)現(xiàn)了共享內(nèi)存中按行主序?qū)懭牒桶戳兄餍蜃x取

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

動(dòng)態(tài)共享內(nèi)存:

這里使用上頭講到的動(dòng)態(tài)內(nèi)存

動(dòng)態(tài)共享內(nèi)存可以在核函數(shù)之外聲明, 其作用域?qū)⑹钦麄€(gè)文件
也可以在核函數(shù)之內(nèi)聲明, 其作用域?qū)H限于核函數(shù)

例程:

核函數(shù)中按行主序?qū)懭? 按列主序讀取

nvprof結(jié)果:

所以表明了使用動(dòng)態(tài)共享內(nèi)存也會(huì)存在相同的問題

填充動(dòng)態(tài)聲明的共享內(nèi)存:

這里是對(duì)動(dòng)態(tài)共享內(nèi)存執(zhí)行內(nèi)存填充

填充動(dòng)態(tài)聲明的共享內(nèi)存數(shù)組更加復(fù)雜

因?yàn)樵谝陨虾撕瘮?shù)中用于存儲(chǔ)數(shù)據(jù)的全局內(nèi)存小于填充的共享內(nèi)存, 所以需要3個(gè)索引：一個(gè)索引用于按照行主序?qū)懭牍蚕韮?nèi)存, 一個(gè)索引用于按照列主序讀取共享內(nèi)存, 一個(gè)索引用于未填充的全局內(nèi)存的合并訪問

這些結(jié)果和填充靜態(tài)聲明的共享內(nèi)存是一致的

所以這里證明的是, 無論是靜態(tài)共享內(nèi)存還是動(dòng)態(tài)共享內(nèi)存都能被有效的填充

方形共享內(nèi)存內(nèi)核性能的比較:

到目前為止, 從所有執(zhí)行過的內(nèi)核運(yùn)行時(shí)間可以看出：

• 使用填充的內(nèi)核可提高性能, 因?yàn)樗鼫p少了存儲(chǔ)體沖突

• 帶有動(dòng)態(tài)聲明共享內(nèi)存的內(nèi)核增加了少量的消耗

矩形共享內(nèi)存:

這一部分的行文邏輯基本上和上一節(jié)相同, 討論共享內(nèi)存的幾個(gè)點(diǎn), 只不過吧上頭的方陣替換為了矩陣

矩形共享內(nèi)存是一個(gè)更普遍的二維共享內(nèi)存, 他與方形共享內(nèi)存的區(qū)別就是行列數(shù)不等 ( 矩陣 & 方陣的區(qū)別)

本部分的所有核函數(shù)調(diào)用都使用以下執(zhí)行配置:

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

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

就是將上頭的方陣替換為了矩陣內(nèi)存, 并執(zhí)行內(nèi)存轉(zhuǎn)置操作:

這里使用的應(yīng)該是16個(gè)數(shù)據(jù), 而并非之前方陣的32個(gè), 所以數(shù)據(jù)不同, 但是結(jié)論是相同的

共享內(nèi)存的存儲(chǔ)和加載請(qǐng)求, 由setRowReadRow核函數(shù)中的一個(gè)事務(wù)完成. 同樣的請(qǐng)求在setColReadCol函數(shù)中由8個(gè)事務(wù)完成. Kepler K40的存儲(chǔ)體寬度是8個(gè)字, 一列16個(gè)4字節(jié)的數(shù)據(jù)元素被安排到8個(gè)存儲(chǔ)體中, 如圖5-6所示, 因此, 該操作有一個(gè)8路沖突

行主序?qū)?& 列主序讀:

使用共享內(nèi)存進(jìn)行矩陣轉(zhuǎn)置的核函數(shù). 通過最大化低延遲的加載和存儲(chǔ)來提高性能, 并合并全局內(nèi)存訪問

內(nèi)核有3個(gè)內(nèi)存操作：

• 寫入每個(gè)線程束的共享內(nèi)存行, 以避免存儲(chǔ)體沖突

• 讀取每個(gè)線程束中的共享內(nèi)存列, 以完成矩陣轉(zhuǎn)置

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

該存儲(chǔ)操作是無沖突的, 加載操作報(bào)告了一個(gè)8路沖突

與預(yù)期相同
store時(shí)是行主序, load時(shí)是列主序

動(dòng)態(tài)共享內(nèi)存:

還是緊接著上頭的例子進(jìn)行修改, 將其中的靜態(tài)內(nèi)存改為動(dòng)態(tài)內(nèi)存, 繼續(xù)實(shí)現(xiàn)矩陣轉(zhuǎn)置

結(jié)果與使用靜態(tài)內(nèi)存相同

所以結(jié)論就是:

動(dòng)態(tài)分配共享內(nèi)存不會(huì)影響存儲(chǔ)體沖突

填充靜態(tài)共享內(nèi)存:

在前面的宏中若將填充數(shù)據(jù)元素的數(shù)量從2改到1, 則nvprof報(bào)告有兩個(gè)事務(wù)完成共享內(nèi)存的加載操作, 即發(fā)生一個(gè)雙向存儲(chǔ)體沖突

所以結(jié)論是:

填充的元素個(gè)數(shù)與行列數(shù)是有關(guān)系的, 數(shù)量不當(dāng)仍將導(dǎo)致存儲(chǔ)體沖突

填充動(dòng)態(tài)共享內(nèi)存:

大致套路和靜態(tài)共享內(nèi)存相同:

結(jié)論就是:

動(dòng)態(tài)內(nèi)存的填充比靜態(tài)內(nèi)存的仍然要復(fù)雜
其有專門的計(jì)算index 的代碼

矩形共享內(nèi)存內(nèi)核性能的比較:

在一般情況下, 和上一節(jié)說到的一樣:

• 核函數(shù)使用共享內(nèi)存填充消除存儲(chǔ)體沖突以提高性能
• 使用動(dòng)態(tài)共享內(nèi)存的核函數(shù)會(huì)顯示有少量的消耗

5.3 減少全局內(nèi)存訪問:

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

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

使用共享內(nèi)存的并行歸約:

首先是一個(gè)僅使用全局內(nèi)存的歸約核函數(shù), 作為所有核函數(shù)的起點(diǎn)與性能的基點(diǎn):

而后是帶有共享內(nèi)存的全局內(nèi)存操作的歸約函數(shù)

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

二者對(duì)比如下:

使用共享內(nèi)存的核函數(shù)比只使用全局內(nèi)存的核函數(shù)快了1.84倍

使用nvprof的倆參數(shù)查看全局內(nèi)存加載&存儲(chǔ)事務(wù):

使用展開的并行歸約

這里就是在上一節(jié)的例子中加上之前的循環(huán)展開方法:

以下內(nèi)核展開了4個(gè)線程塊, 即每個(gè)線程處理來自于4個(gè)數(shù)據(jù)塊的數(shù)據(jù)元素
可預(yù)期的效果是:

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

• 全局內(nèi)存存儲(chǔ)事務(wù)減少了1/4

• 整體內(nèi)核性能的提升

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

使用動(dòng)態(tài)共享內(nèi)存的并行歸約

這里一筆帶過, 直接上結(jié)論;

用動(dòng)態(tài)分配共享內(nèi)存實(shí)現(xiàn)的核函數(shù)和用靜態(tài)分配共享內(nèi)存實(shí)現(xiàn)的核函數(shù)之間沒有顯著的差異

有效帶寬:

由于歸約核函數(shù)是受內(nèi)存帶寬約束的, 所以評(píng)估它們時(shí)所使用的適當(dāng)?shù)男阅苤笜?biāo)是有效帶寬

有效帶寬是在核函數(shù)的完整執(zhí)行時(shí)間內(nèi)I/O的數(shù)量（以字節(jié)為單位）

對(duì)于內(nèi)存約束的應(yīng)用程序, 有效帶寬是一個(gè)估算實(shí)際帶寬利用率的很好的指標(biāo)

計(jì)算公式:
$$有效帶寬(GB/s)=\frac{(讀字節(jié)數(shù)+寫字節(jié)數(shù))}{運(yùn)行時(shí)間?10^9}$$

以下是前頭的4個(gè)函數(shù)的有效帶寬:

顯然, 可以通過展開塊來獲得有效帶寬的顯著改進(jìn)
每個(gè)線程運(yùn)行中同時(shí)有多個(gè)請(qǐng)求, 會(huì)導(dǎo)致內(nèi)存總線高飽和

5.4 合并的全局內(nèi)存訪問:

使用共享內(nèi)存也能幫助避免產(chǎn)生未合并的全局內(nèi)存訪問

之前的矩陣轉(zhuǎn)置核函數(shù)中, 讀操作是合并的, 但寫操作是交叉訪問的

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

在本章前面的部分, 測試了一個(gè)矩陣轉(zhuǎn)置核函數(shù), 該核函數(shù)使用單個(gè)線程塊對(duì)共享內(nèi)存中的矩陣行進(jìn)行寫入, 并讀取共享內(nèi)存中的矩陣列

在本節(jié)中, 將擴(kuò)展該核函數(shù), 具體方法是使用多個(gè)線程塊對(duì)基于交叉的全局內(nèi)存訪問重新排序到合并訪問

基準(zhǔn)轉(zhuǎn)置核函數(shù):

和上一節(jié)的行文邏輯相同, 先確定一個(gè)性能比較的基準(zhǔn)

下面的核函數(shù)是一個(gè)僅使用全局內(nèi)存的矩陣轉(zhuǎn)置的樸素實(shí)現(xiàn)
其中, 全局內(nèi)存讀操作在線程束內(nèi)是被合并的, 而全局內(nèi)存寫操作在相鄰線程間是交叉訪問的

而后這個(gè)核函數(shù)將作為優(yōu)化的性能上限
其中讀寫操作都將被合并, 仍執(zhí)行相同數(shù)量的IO

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

基準(zhǔn)核函數(shù)的運(yùn)行結(jié)果:

副本內(nèi)核比樸素內(nèi)核快了將近3倍

由于樸素內(nèi)核寫入全局內(nèi)存, 使其帶有了4096個(gè)元素的跨度, 所以一個(gè)單一線程束的存儲(chǔ)內(nèi)存操作是由32個(gè)全局內(nèi)存事務(wù)完成的. 可以使用以下nvprof指標(biāo)來確認(rèn)這一點(diǎn)

使用共享內(nèi)存的矩陣轉(zhuǎn)置:

為了避免交叉全局內(nèi)存訪問, 可以使用二維共享內(nèi)存來緩存原始矩陣的數(shù)據(jù)

實(shí)現(xiàn)的核函數(shù):
可以看做是上一節(jié)中的setRowReadCol的擴(kuò)展, 前者使用的單一線程塊, 而后者將其擴(kuò)展為了使用多個(gè)線程塊和數(shù)據(jù)塊

核函數(shù)的程序步驟:

kerneltransposeSmem函數(shù)可被分解為以下幾個(gè)步驟：

線程束執(zhí)行合并讀取一行, 該行存儲(chǔ)在全局內(nèi)存中的原始矩陣塊中.

然后, 該線程束按行主序?qū)⒃摂?shù)據(jù)寫入共享內(nèi)存中, 因此, 這個(gè)寫操作沒有存儲(chǔ)體沖突.

因?yàn)榫€程塊的讀/寫操作是同步的, 所以會(huì)有一個(gè)填滿全局內(nèi)存數(shù)據(jù)的二維共享內(nèi)存數(shù)組.

該線程束從二維共享內(nèi)存數(shù)組中讀取一列. 由于共享內(nèi)存沒有被填充, 所以會(huì)發(fā)生存儲(chǔ)體沖突.

然后該線程束執(zhí)行數(shù)據(jù)的合并寫入操作, 將其寫入到全局內(nèi)存的轉(zhuǎn)置矩陣中的某行

核函數(shù)具體的實(shí)現(xiàn)就暫且略過了(詳見書里), 這里來看其實(shí)現(xiàn)的特點(diǎn):

全局內(nèi)存的讀取是合并的

共享內(nèi)存的寫入沒有發(fā)生存儲(chǔ)體沖突

共享內(nèi)存的讀取發(fā)生存儲(chǔ)體沖突
是按列讀取

全局內(nèi)存的寫入是合并的

性能對(duì)比:

全局內(nèi)存存儲(chǔ)的重復(fù)數(shù)量從32減少到2

這是由于轉(zhuǎn)置塊中的塊寬為16, 所以線程束前半部分的寫操作和線程束后半部分的寫操作間隔了4080
因此線程束的寫入請(qǐng)求是有兩個(gè)事務(wù)完成的

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

顯然, 讀取二維共享內(nèi)存數(shù)組中的一列會(huì)產(chǎn)生存儲(chǔ)體沖突

使用填充共享內(nèi)存的矩陣轉(zhuǎn)置:

這里就是應(yīng)用之前的填充

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

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

修改之前的共享內(nèi)存聲明如下:

使用展開的矩陣轉(zhuǎn)置:

就是在添加一個(gè)循環(huán)展開

下面的核函數(shù)展開兩個(gè)數(shù)據(jù)塊的同時(shí)處理：每個(gè)線程現(xiàn)在轉(zhuǎn)置了被一個(gè)數(shù)據(jù)塊跨越的兩個(gè)數(shù)據(jù)元素
這種轉(zhuǎn)化的目標(biāo)是通過創(chuàng)造更多的同時(shí)加載和存儲(chǔ)以提高設(shè)備內(nèi)存帶寬利用率

核函數(shù)的其他詳細(xì)實(shí)現(xiàn)直接去看書, 其特點(diǎn)都在上頭的這個(gè)圖里

增大并行性:

這里是通過調(diào)整線程塊的維度來提升性能

塊大小為16×16時(shí)展示出了最好的性能, 因?yàn)樗懈嗟牟l(fā)線程塊, 從而有最好的設(shè)備并行性

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

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

5.5 常量內(nèi)存:

常量內(nèi)存是一種專用的內(nèi)存
其對(duì)內(nèi)核代碼而言是只讀的，但它對(duì)主機(jī)而言既是可讀又是可寫的

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

與其他類型的內(nèi)存不同, 常量內(nèi)存有一個(gè)最優(yōu)訪問模式:

• 當(dāng)線程束中的so哦有線程都訪問相同的位置, 此時(shí)訪問模式是最優(yōu)的
• 如果線程束訪問不同的地址, 則需要串行訪問

所以常量內(nèi)存的讀取成本與線程束中讀取的地址數(shù)量息息相關(guān)

使用__constant__聲明一個(gè)常量變量

由于常量內(nèi)存在設(shè)備上只讀, 所以必須在主機(jī)上進(jìn)行初始化:

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

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

常量內(nèi)存的幾個(gè)特點(diǎn)

• 生存期與應(yīng)用程序相同
• 對(duì)網(wǎng)格內(nèi)的所有線程可見
• 主機(jī)也可以直接訪問

使用常量內(nèi)存實(shí)現(xiàn)一維模板:

又開始了, 神奇的翻譯

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

這里重點(diǎn)不是理解這個(gè)九點(diǎn)模板, 而是考慮到他的使用場景:

9個(gè)x作為輸入, 一個(gè)輸出

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

所實(shí)現(xiàn)的核函數(shù)

代碼實(shí)現(xiàn)具體看書, 這里重點(diǎn)關(guān)注他的常量內(nèi)存的使用:

與只讀緩存的比較:

這里要講到Kepler架構(gòu)中添加的獨(dú)立只讀數(shù)據(jù)緩存:

這里需要注意, 好像僅僅是Kepler架構(gòu)中有這玩意, 在后續(xù)的架構(gòu)中并沒有這玩意:

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

只讀緩存的使用:

當(dāng)通過只讀緩存訪問全局內(nèi)存時(shí)，需要向編譯器指出在內(nèi)核的持續(xù)時(shí)間里數(shù)據(jù)是只讀的

• 使用內(nèi)部函數(shù)__ldg

• 全局內(nèi)存的限定指針

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

與常量內(nèi)存的對(duì)比:

• 常量緩存加載的數(shù)據(jù)必須是少量的, 并且需要訪問的一致性才能獲得較好的性能
• 制度緩存加載的數(shù)據(jù)可以是比較大的, 而且能在一個(gè)非統(tǒng)一的模式下進(jìn)行訪問

所以可以得出以下結(jié)論:

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

核函數(shù)實(shí)現(xiàn):
此核函數(shù)和上頭的唯一區(qū)別就是函數(shù)聲明部分

在Tesla K40上，使用nvprof測試得出的以下結(jié)果表明，對(duì)此應(yīng)用程序使用只讀內(nèi)存時(shí)其性能實(shí)際上會(huì)降低。這是由于coef數(shù)組使用了廣播訪問模式，相比于只讀緩存，該模式更適合于常量內(nèi)存：

5.6 線程束洗牌指令:

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

洗牌指令比共享內(nèi)存有更低的延遲，并且該指令在執(zhí)行數(shù)據(jù)交換時(shí)不消耗額外的內(nèi)存

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

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

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

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

廣播:

在線程束內(nèi)交換整型變量，其基本函數(shù)標(biāo)記如下：

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

書里的是這個(gè)形式:

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

線程束內(nèi)所有活躍的線程都同時(shí)產(chǎn)生此操作，這將導(dǎo)致每個(gè)線程中有4字節(jié)數(shù)據(jù)的移動(dòng)

參數(shù)解釋:

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

• var
  root線程共享出來的值

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

• width:
  洗牌分段
  默認(rèn)=warpSize=32 , 此時(shí)洗牌操作的作用范圍是整個(gè)線程束
  但是通過手動(dòng)設(shè)置值可以調(diào)的更細(xì), 使每段包含有width個(gè)線程, 并且每段上指定獨(dú)立的洗牌操作

  此時(shí)srcLane使用的線程ID與束內(nèi)線程ID不同, 其使用如下公式計(jì)算:

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

所以可知, 這個(gè)操作有點(diǎn)類似于MPI中的廣播

但是這里是吧__shlf中的參數(shù)寫死了的情況

如果使用動(dòng)態(tài)參數(shù), 可以得到下一節(jié)中的循環(huán)交換的效果:

復(fù)制:

__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);

參數(shù)介紹:

• delta
  線程束偏移量

其他參數(shù)都和上頭的廣播相似

偏移量這個(gè)就是下圖所展現(xiàn)的

而這兩個(gè)函數(shù)的區(qū)別就是方向不同:

• up向高index方向復(fù)制
• down向低index方向復(fù)制

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

交叉交換:

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

不具體闡述了

線程束內(nèi)的共享數(shù)據(jù)

在本節(jié)中，會(huì)介紹幾個(gè)有關(guān)線程束洗牌指令的例子，并說明線程束洗牌指令的優(yōu)點(diǎn)

洗牌指令將被應(yīng)用到以下3種整數(shù)變量類型中：

• 標(biāo)量變量

• 數(shù)組

• 向量型變量

下頭就全是例子, 就中間的循環(huán)移動(dòng)需要看看以外, 其他的都一筆帶過

線程束內(nèi)的值廣播:

這里就是對(duì)上一節(jié)講到的洗牌指令進(jìn)行應(yīng)用:

執(zhí)行結(jié)果:

線程束內(nèi)上移:

這里還是應(yīng)用

線程束內(nèi)下移:

線程束內(nèi)環(huán)繞移動(dòng):

這里實(shí)現(xiàn)的就是上頭所沒有實(shí)現(xiàn)的環(huán)繞式移動(dòng), 即首尾相接的__shfl_up 或 __shfl_down

• 正偏移量為左移
• 負(fù)偏移量為右移

蝴蝶(交叉)交換:

交換數(shù)組值:

這個(gè)其實(shí)就是__shfl_xor()的花式應(yīng)用

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

使用數(shù)組索引交換數(shù)值:

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

本部分先放著

布爾變量pred被用于識(shí)別第一個(gè)調(diào)用的線程，它是交換數(shù)據(jù)的一對(duì)線程。要交換的數(shù)據(jù)元素是由第一個(gè)線程的firstIdx和第二個(gè)線程的secondIdx偏移標(biāo)識(shí)的。第一個(gè)調(diào)用線程通過交換firstIdx和secondIdx中的元素開始，但此操作僅限于本地?cái)?shù)組。然后在兩線程間的secondIdx位置執(zhí)行蝴蝶交換。最后，第一個(gè)線程交換接收自secondIdx返回到firstIdx的元素

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

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

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

• 線程束級(jí)歸約

• 線程塊級(jí)歸約

• 網(wǎng)格級(jí)歸約

詳細(xì)的解釋可以看書:

這里直接看結(jié)果:

用洗牌指令實(shí)現(xiàn)線程束級(jí)并行歸約獲得了1.42倍的加速

5.7 總結(jié):

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

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

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

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

只讀紋理緩存提供了常量內(nèi)存的替代方案，該方案優(yōu)化了數(shù)據(jù)的分散讀取。只讀緩存訪問全局內(nèi)存中的數(shù)據(jù)，但它使用一個(gè)獨(dú)立的內(nèi)存訪問流水線和獨(dú)立的緩存，以使SM可以訪問數(shù)據(jù)。因此，只讀緩存共享了常量內(nèi)存的許多好處，同時(shí)對(duì)不同的訪問模式也進(jìn)行了優(yōu)化。

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

本章對(duì)一些有特殊用途的內(nèi)存類型進(jìn)行了深度了解。雖然這些內(nèi)存類型比全局內(nèi)存使用得少，但是適當(dāng)?shù)厥褂盟鼈兛梢蕴岣邘捓寐?#xff0c;降低整體的內(nèi)存延遲。如果你正在研究優(yōu)化的因素，那么牢記共享內(nèi)存、常量內(nèi)存、只讀緩存和洗牌指令都是非常重要的。

總結(jié)

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

如果覺得生活随笔網(wǎng)站內(nèi)容還不錯(cuò)，歡迎將生活随笔推薦給好友。