文章目錄

• 前言
• 1. 紋理內存的理解
• 2. C++ CUDA實現常量內存使用
• 3. 執行結果
• 總結
• 學習資料

VS2017 CUDA編程學習1：CUDA編程兩變量加法運算
VS2017 CUDA編程學習2：在GPU上執行線程
VS2017 CUDA編程學習3：CUDA獲取設備上屬性信息
VS2017 CUDA編程學習4：CUDA并行處理初探 - 向量加法實現
VS2017 CUDA編程學習5：CUDA并行執行-線程
VS2017 CUDA編程學習6: GPU存儲器架構
VS2017 CUDA編程學習7：線程同步-共享內存
VS2017 CUDA編程學習8：線程同步-原子操作
VS2017 CUDA編程學習9：常量內存

---

前言

這里繼續CUDA編程的學習，今天學習了設備紋理內存的使用，這里分享給大家！

---

1. 紋理內存的理解

紋理內存，是設備中一種只讀存儲器，在設備內存中也有對應的緩沖器Cache來加速IO操作。紋理內存對于程序中有空間鄰近性訪問的數據很高效。這里的空間鄰近性是指每個線程的讀取位置和其他線程的讀取位置鄰近，這在圖像處理中非常有用，因為圖像處理中經常需要進行鄰域比較，比如4鄰域和8鄰域。

當然全局內存也可以進行這種有空間鄰近性特性數據的存儲與訪問，但是速度上要慢很多，需要進行大量的顯存讀取操作。但是紋理內存則只需要從顯存讀取一次就可以（不理解這種說法，書中這么說的），紋理內存支持2維和3維紋理讀取操作。

---

2. C++ CUDA實現常量內存使用

這里以一個簡單的例子：從紋理內存中讀取數據并賦值。

紋理內存需要和“紋理引用”與CUDA數組配套使用來實現。
紋理引用通過$texture<param1,param2,param3>$類型變量來定義，param1表示紋理元素的類型，param2表示紋理引用的類型（1-表示1D, 2-表示2D, 3-表示3D），param3則表示讀取模式，這是一個可選參數，說明是否要執行自動類型轉換。

注意，紋理引用變量要確保被定義為全局靜態變量（個人理解，這樣才能在內核函數中使用），同時還要確保這個變量不能作為參數傳遞給任何其他函數。

下面的例子中通過紋理引用讀取線程ID作為索引位置的數據，然后復制到d_out指針指向全局內存中。

在main函數中則使用CUDA數組$cudaArray$并通過$cudaMemcpyToArray(param1,param2,param3,param4,param5,param6)$對CUDA數組進行賦值，param1表示目標CUDA數組變量， param2和param3表示將主機數據賦值到CUDA數組橫向和縱向的偏移量，（0， 0）表示CUDA數組的左上角開始賦值。

使用$cudaBindTextureToArray()$函數將紋理引用和CUDA數組綁定，這樣就可以通過訪問紋理引用來獲取紋理內存中存儲的CUDA數組數據。這里使用texture相應的函數進行訪問。當使用完紋理內存后，要執行解除紋理引用與CUDA數組的綁定的代碼，即調用$cudaUnbindTexture()$。

詳細代碼如下所示：

#include <stdio.h> #include <iostream> #include <cuda.h> #include <cuda_runtime.h> #include <cuda_runtime_api.h> #include <device_launch_parameters.h>#ifndef __CUDACC__ #define __CUDACC__ #endif #include <cuda_texture_types.h>//texture<> depend head file #include <texture_fetch_functions.h>//tex1D() depend head file#define NUM_THREADS 10 #define N 10//定義1維的紋理引用 texture<float, 1, cudaReadModeElementType> textureRef; //textureReference textureRef;//定義內核函數：從紋理內存中獲取數據并賦值給設備內存 __global__ void gpu_texture_memory(int n, float* d_out) {int idx = blockIdx.x *blockDim.x + threadIdx.x;if (idx < n){float temp = tex1D(textureRef, float(idx));d_out[idx] = temp;} }int main() {int num_blocks = N / NUM_THREADS + ((N % NUM_THREADS) ? 1 : 0);float* d_out;cudaMalloc(&d_out, sizeof(float) * N);float* h_out = (float*)malloc(sizeof(float) * N);float h_in[N];for (int i = 0; i < N; i++){h_in[i] = float(i);}//定義CUDA數組cudaArray *cu_Array;cudaMallocArray(&cu_Array, &textureRef.channelDesc, N, 1);cudaMemcpyToArray(cu_Array, 0, 0, h_in, sizeof(float)*N, cudaMemcpyHostToDevice);//綁定CUDA數組到紋理引用變量cudaBindTextureToArray(&textureRef, cu_Array, &textureRef.channelDesc);//調用內核函數gpu_texture_memory << <num_blocks, NUM_THREADS >> > (N, d_out);//拷貝結果到主機cudaMemcpy(h_out, d_out, sizeof(float)*N, cudaMemcpyDeviceToHost);printf("在GPU上使用紋理內存：\n");//打印結果for (int i = 0; i < N; i++){printf("%f\n", h_out[i]);}//回收分配的資源free(h_out);cudaFree(d_out);cudaFreeArray(cu_Array);cudaUnbindTexture(&textureRef);system("pause");return 0; }

---

3. 執行結果

---

總結

說實話，這里還是無法體會到紋理內存的提速，這里只是簡單介紹了紋理內存的用法，由于這里只是記錄個人的學習筆記，所以如果有誤，還請原諒，如果能夠指正，萬分感謝！

學習資料

《基于GPU加速的計算機視覺編程》

總結

以上是生活随笔為你收集整理的VS2017 CUDA编程学习10：纹理内存的全部內容，希望文章能夠幫你解決所遇到的問題。

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