CUDA学习日志：常量内存和纹理内存
版權聲明：本文為博主原創文章，未經博主允許不得轉載。
接觸CUDA的時間并不長，最開始是在cuda-convnet的代碼中接觸CUDA代碼，當時確實看的比較痛苦。最近得空，在圖書館借了本《GPU高性能編程 CUDA實戰》來看看，同時也整理一些博客來加強學習效果。
https://www.baidu.com/Jeremy Lin
在上一篇博文中，我們談到了如何利用共享內存來實現線程協作的問題。本篇博文我們主要來談談如何利用常量內存和紋理內存來提高程序性能。
常量內存
所謂的常量內存，從它的名字我們就可以知道，它是用來保存在核函數執行期間不會發生變化的數據。NVIDIA硬件提供了64KB的常量內存，并且常量內存采用了不同于標準全局內存的處理方式。在某些情況下，用常量內存來替換全局內存可以有效地減少內存帶寬。
常量內存的聲明方式與共享內存是類似的。要使用常量內存，則需要在變量前面加上 __constant__修飾符：
[cpp] view plaincopy在之前的程序中，我們為變量分配內存時是先聲明一個指針，然后通過cudaMalloc()來為指針分配GPU內存。而當我們將其改為常量內存時，則要將這個聲明修改為在常量內存中靜態地分配空間。我們不再需要對變量指針調用cudaMalloc()或者cudaFree()，而是在編譯時為這個變量（如數組s）提交固定的大小。
另外一個值得注意的是，當從主機內存復制到GPU上的常量內存時，我們需要使用一個特殊版本的cudaMemcpy()，即是：
?cudaMemcpyToSymbol()
cudaMemcpyToSymbol()和參數為cudaMemcpyHostToDevice()的cudaMemcpy()之間的唯一差異在于，cudaMemcpyToSymbol()會復制到常量內存，而cudaMemcpy()會復制到全局內存。
note：變量修飾符 __constant__ 將變量的訪問限制為只讀。
那么常量內存為什么能帶來性能提升～
原因:
- 對常量內存的單次讀操作可以廣播到其他的“鄰近(nearby)”線程，這將節約15次讀取操作；
- 常量內存的數據將緩存起來，因此對于相同地址的連續操作將不會產生額外的內存通信量。
下面我們來具體講講這兩個原因。
首先，我們需要來看看到底什么是線程束(warp)，在CUDA架構中，線程束是指一個包含32個線程的集合，這個線程集合被“編織在一起”并且以“步調一致(Lockstep)”的形式執行。在程序中的每一行，線程束中的每個線程都將在不同的數據上執行相同的指令。
當處理常量內存時，NVIDIA硬件將把單次內存讀取操作廣播到每個半線https://www.baidu.com/程束(Half-Warp)。在半線程束中包含16個線程，即線程束中線程數量的一半。如果在半線程束中的每個線程從常量內存的相同地址上讀取數據，那么GPU只會產生一次讀取請求并在隨后將數據廣播到每個線程。如果從常量內存中讀取大量數據，那么這種方式產生的內存流量只是使用全局內存時的1/16。
但在讀取常量內存時，所節約的并不僅限于減少94%的帶寬。由于這塊內存的內容是不發生變化的，因此硬件將主動把這個常量數據緩存在GPU上。在第一次從常量內存的某個地址上讀取后，當其他半線程束請求同一個地址時，那么將命中緩存，這同樣減少了額外的內存流量。
不過正如上一篇博文講的__syncthread()不能亂用一樣，常量內存也不能亂用。它可能會對性能產生負面的影響。半線程束廣播功能實際上是一把雙刃劍。雖然當所有16個線程都讀取相同地址時，這個功能可以極大提升性能，但當所有16個線程分別讀取不同的地址時，它實際上會降低性能。因為這16次不同的讀取操作會被串行化，從而需要16倍的時間來發出請求。但如果從全局內存中讀取，那么這些請求會同時發出。
紋理內存
和常量內存一樣，紋理內存是另一種類型的只讀內存，在特定的訪問模式中，紋理內存同樣能夠提升性能并減少內存流量。
紋理內存緩存在芯片上，因此在某些情況中，它能夠減少對內存的請求并提供更高效的內存帶寬。紋理緩存是專門為那些在內存訪問模式中存在大量空間局部性(Spatial Locality)的圖形應用程序而設計的。在某個計算應用程序中，這意味著一個線程讀取的位置可能與鄰近線程的讀取位置“非常接近”，如下圖所示。
從數學的角度，上圖中的4個地址并非連續的，在一般的CPU緩存中，這些地址將不會緩存。但由于GPU紋理緩存是專門為了加速這種訪問模式而設計的，因此如果在這種情況中使用紋理內存而不是全局內存，那么將會獲得性能的提升。
下面，我們來看看如何使用紋理內存。
首先，需要將輸入的數據聲明為texture類型的引用。比如：
[cpp] view plaincopy
然后，就是在為這三個緩沖區分配GPU內存后，需要通過cudaBindTexture()將這些變量綁定到內存緩存區。這相當于告訴CUDA運行時兩件事：
- 我們希望將指定的緩沖區作為紋理來使用；
- 我們希望將紋理引用作為紋理的“名字”。
當用cudaBindTexture綁定后，紋理變量就設置好了，現在可以啟動核函數。
然而，當讀取核函數中的紋理時，需要通過特殊的函數來告訴GPU將讀取請求轉發到紋理內存而不是標準的全局內存。因此，當讀取內存時，需要使用特殊的方式：從線性內存中讀取（拾取），使用的函數是tex1Dfetch()：
[cpp] view plaincopy
這些函數用紋理坐標x拾取綁定到紋理參考texRef 的線性內存的區域。不支持紋理過濾和尋址模式。對于整數型，這些函數將會將整數型轉化為單精度浮點型。除了這些函數，還支持2-和4-分量向量的拾取。
最后，當應用程序運行結束后，還要清除紋理的綁定。
[cpp] view plaincopy
二維紋理內存的聲明如下：
[cpp] view plaincopy綁定函數：
cudaBindTexture2D()
紋理拾取函數：
tex2D()
本文地址：http://blog.csdn.net/linj_m/article/details/41522573
更多資源請 關注博客：LinJM-機器視覺 ?微博：林建民-機器視覺
總結
以上是生活随笔為你收集整理的CUDA学习日志：常量内存和纹理内存的全部內容，希望文章能夠幫你解決所遇到的問題。
- 上一篇： wpf 文件上传到服务器_07-文件上传
- 下一篇： cuda profiler使用