標簽： cuda常量內(nèi)存紋理內(nèi)存LinJM 2014-11-29 21:53 2143人閱讀 評論(0) 收藏 舉報 分類： 圖像處理與分析（24）

版權(quán)聲明：本文為博主原創(chuàng)文章，未經(jīng)博主允許不得轉(zhuǎn)載。

接觸CUDA的時間并不長，最開始是在cuda-convnet的代碼中接觸CUDA代碼，當時確實看的比較痛苦。最近得空，在圖書館借了本《GPU高性能編程 CUDA實戰(zhàn)》來看看，同時也整理一些博客來加強學習效果。

https://www.baidu.com/Jeremy Lin

在上一篇博文中，我們談到了如何利用共享內(nèi)存來實現(xiàn)線程協(xié)作的問題。本篇博文我們主要來談?wù)勅绾卫贸Ａ績?nèi)存和紋理內(nèi)存來提高程序性能。

常量內(nèi)存

所謂的常量內(nèi)存，從它的名字我們就可以知道，它是用來保存在核函數(shù)執(zhí)行期間不會發(fā)生變化的數(shù)據(jù)。NVIDIA硬件提供了64KB的常量內(nèi)存，并且常量內(nèi)存采用了不同于標準全局內(nèi)存的處理方式。在某些情況下，用常量內(nèi)存來替換全局內(nèi)存可以有效地減少內(nèi)存帶寬。

常量內(nèi)存的聲明方式與共享內(nèi)存是類似的。要使用常量內(nèi)存，則需要在變量前面加上 __constant__修飾符：

[cpp] view plaincopy

__constant__?int?s[10000]??

在之前的程序中，我們?yōu)樽兞糠峙鋬?nèi)存時是先聲明一個指針，然后通過cudaMalloc()來為指針分配GPU內(nèi)存。而當我們將其改為常量內(nèi)存時，則要將這個聲明修改為在常量內(nèi)存中靜態(tài)地分配空間。我們不再需要對變量指針調(diào)用cudaMalloc()或者cudaFree()，而是在編譯時為這個變量（如數(shù)組s）提交固定的大小。

另外一個值得注意的是，當從主機內(nèi)存復制到GPU上的常量內(nèi)存時，我們需要使用一個特殊版本的cudaMemcpy()，即是：

?cudaMemcpyToSymbol()

cudaMemcpyToSymbol()和參數(shù)為cudaMemcpyHostToDevice()的cudaMemcpy()之間的唯一差異在于，cudaMemcpyToSymbol()會復制到常量內(nèi)存，而cudaMemcpy()會復制到全局內(nèi)存。

note：變量修飾符 __constant__ 將變量的訪問限制為只讀。

那么常量內(nèi)存為什么能帶來性能提升～

原因:

• 對常量內(nèi)存的單次讀操作可以廣播到其他的“鄰近(nearby)”線程，這將節(jié)約15次讀取操作；
• 常量內(nèi)存的數(shù)據(jù)將緩存起來，因此對于相同地址的連續(xù)操作將不會產(chǎn)生額外的內(nèi)存通信量。

下面我們來具體講講這兩個原因。

首先，我們需要來看看到底什么是線程束(warp)，在CUDA架構(gòu)中，線程束是指一個包含32個線程的集合，這個線程集合被“編織在一起”并且以“步調(diào)一致(Lockstep)”的形式執(zhí)行。在程序中的每一行，線程束中的每個線程都將在不同的數(shù)據(jù)上執(zhí)行相同的指令。

當處理常量內(nèi)存時，NVIDIA硬件將把單次內(nèi)存讀取操作廣播到每個半線https://www.baidu.com/程束(Half-Warp)。在半線程束中包含16個線程，即線程束中線程數(shù)量的一半。如果在半線程束中的每個線程從常量內(nèi)存的相同地址上讀取數(shù)據(jù)，那么GPU只會產(chǎn)生一次讀取請求并在隨后將數(shù)據(jù)廣播到每個線程。如果從常量內(nèi)存中讀取大量數(shù)據(jù)，那么這種方式產(chǎn)生的內(nèi)存流量只是使用全局內(nèi)存時的1/16。

但在讀取常量內(nèi)存時，所節(jié)約的并不僅限于減少94%的帶寬。由于這塊內(nèi)存的內(nèi)容是不發(fā)生變化的，因此硬件將主動把這個常量數(shù)據(jù)緩存在GPU上。在第一次從常量內(nèi)存的某個地址上讀取后，當其他半線程束請求同一個地址時，那么將命中緩存，這同樣減少了額外的內(nèi)存流量。

不過正如上一篇博文講的__syncthread()不能亂用一樣，常量內(nèi)存也不能亂用。它可能會對性能產(chǎn)生負面的影響。半線程束廣播功能實際上是一把雙刃劍。雖然當所有16個線程都讀取相同地址時，這個功能可以極大提升性能，但當所有16個線程分別讀取不同的地址時，它實際上會降低性能。因為這16次不同的讀取操作會被串行化，從而需要16倍的時間來發(fā)出請求。但如果從全局內(nèi)存中讀取，那么這些請求會同時發(fā)出。

紋理內(nèi)存

和常量內(nèi)存一樣，紋理內(nèi)存是另一種類型的只讀內(nèi)存，在特定的訪問模式中，紋理內(nèi)存同樣能夠提升性能并減少內(nèi)存流量。

紋理內(nèi)存緩存在芯片上，因此在某些情況中，它能夠減少對內(nèi)存的請求并提供更高效的內(nèi)存帶寬。紋理緩存是專門為那些在內(nèi)存訪問模式中存在大量空間局部性(Spatial Locality)的圖形應用程序而設(shè)計的。在某個計算應用程序中，這意味著一個線程讀取的位置可能與鄰近線程的讀取位置“非常接近”，如下圖所示。

從數(shù)學的角度，上圖中的4個地址并非連續(xù)的，在一般的CPU緩存中，這些地址將不會緩存。但由于GPU紋理緩存是專門為了加速這種訪問模式而設(shè)計的，因此如果在這種情況中使用紋理內(nèi)存而不是全局內(nèi)存，那么將會獲得性能的提升。

下面，我們來看看如何使用紋理內(nèi)存。

首先，需要將輸入的數(shù)據(jù)聲明為texture類型的引用。比如：

[cpp] view plaincopy

texture<float>??texConst;??

然后，就是在為這三個緩沖區(qū)分配GPU內(nèi)存后，需要通過cudaBindTexture()將這些變量綁定到內(nèi)存緩存區(qū)。這相當于告訴CUDA運行時兩件事：

• 我們希望將指定的緩沖區(qū)作為紋理來使用；
• 我們希望將紋理引用作為紋理的“名字”。

當用cudaBindTexture綁定后，紋理變量就設(shè)置好了，現(xiàn)在可以啟動核函數(shù)。

然而，當讀取核函數(shù)中的紋理時，需要通過特殊的函數(shù)來告訴GPU將讀取請求轉(zhuǎn)發(fā)到紋理內(nèi)存而不是標準的全局內(nèi)存。因此，當讀取內(nèi)存時，需要使用特殊的方式：
從線性內(nèi)存中讀取（拾取），使用的函數(shù)是tex1Dfetch()：
[cpp] view plaincopy

template<class?Type>???

Type?tex1Dfetch(???

?texture<Type,?1,?cudaReadModeElementType>?texRef,???

?int?x);???

float?tex1Dfetch(???

?texture<unsigned?char,?1,?cudaReadModeNormalizedFloat>?texRef,???

?int?x);???

float?tex1Dfetch(???

?texture<signed?char,?1,?cudaReadModeNormalizedFloat>?texRef,???

?int?x);???

float?tex1Dfetch(???

?texture<unsigned?short,?1,?cudaReadModeNormalizedFloat>?texRef,???

?int?x);???

float?tex1Dfetch(???

?texture<signed?short,?1,?cudaReadModeNormalizedFloat>?texRef,???

?int?x);??

這些函數(shù)用紋理坐標x拾取綁定到紋理參考texRef 的線性內(nèi)存的區(qū)域。不支持紋理過濾和尋址模式。對于整數(shù)型，這些函數(shù)將會將整數(shù)型轉(zhuǎn)化為單精度浮點型。除了這些函數(shù)，還支持2-和4-分量向量的拾取。

最后，當應用程序運行結(jié)束后，還要清除紋理的綁定。

[cpp] view plaincopy

cudaUnbindTexture(texConst);??

上面談的都是一維紋理內(nèi)存，實際上也可以使用二維紋理內(nèi)存。

二維紋理內(nèi)存的聲明如下：

[cpp] view plaincopy

texture<float,2>?texConst;??

綁定函數(shù)：

cudaBindTexture2D()

紋理拾取函數(shù)：

tex2D()

本文地址：http://blog.csdn.net/linj_m/article/details/41522573

更多資源請 關(guān)注博客：LinJM-機器視覺 ?微博：林建民-機器視覺

總結(jié)

以上是生活随笔為你收集整理的CUDA学习日志：常量内存和纹理内存的全部內(nèi)容，希望文章能夠幫你解決所遇到的問題。

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