本範例示範 texture memory 的操作:
#include<stdio.h> #include<time.h> #include<stdlib.h> texture<float, 1, cudaReadModeElementType> texture_array; __global__ static void texExp01(float *vec){ vec[threadIdx.x] = tex1Dfetch(texture_array, threadIdx.x); } __global__ static void texExp02(float *vec){ vec[threadIdx.x] = tex1D(texture_array, threadIdx.x / (20.0/4.0)); } int main(int argc,char **argv) { float result[20], *gpu_result; float a[5] = {1, 2, 3, 4, 5}; int i; cudaMalloc((void**) &gpu_result, sizeof(float) * 20); float *gpu_a1; cudaMalloc((void**) &gpu_a1, sizeof(float) * 5); cudaMemcpy(gpu_a1, a, sizeof(float) * 5, cudaMemcpyHostToDevice); cudaBindTexture(0, texture_array, gpu_a1); texExp01<<<1, 5>>>(gpu_result); cudaThreadSynchronize(); cudaMemcpy(&result, gpu_result, sizeof(int) * 5, cudaMemcpyDeviceToHost); printf("Exp 1:\n"); for(i=0;i<5;i++){ printf("\t%.2f\n", result[i]); } cudaArray *gpu_a2; cudaMallocArray(&gpu_a2, &texture_array.channelDesc, 5, 1); cudaBindTextureToArray(texture_array, gpu_a2); texture_array.filterMode = cudaFilterModeLinear; cudaMemcpyToArray(gpu_a2, 0, 0, a, sizeof(float) * 5, cudaMemcpyHostToDevice); texExp02<<<1, 20>>>(gpu_result); cudaThreadSynchronize(); cudaMemcpy(&result, gpu_result, sizeof(float) * 20, cudaMemcpyDeviceToHost); printf("Exp 2:\n"); for(i=0;i<20;i++){ printf("\t%2d %.2f\n", i, result[i]); } return 0; }操作步驟與簡單解說如下:
- 宣告出 texture 的參考
- 語法:「texture<資料類型, 維度, 讀取模式> 名稱;」
- 資料類型:可使用一般的型別,例如 int、float 等等
- 維度:1 或 2
- 讀取模式: "cudaReadModeElementType"為原始方式讀取, "cudaReadModeNormalizedFloat"於資料是整數時,會正規化到 [0, 1] 或 [-1, 1]
- Bind texture:將 gpu 記憶體和 texture 做連結
- 使用 linear memory 時
- cudaBindTexture(offset, texture 的參考, GPU 記憶體)
- offset 一般給 0
- 使用 CUDA array 時
- cudaBindTextureToArray(texture 的參考, GPU 記憶體)
- 支援內插。 材質參考的 filterMode 為 cudaFilterModeLinear 時,是線性內插; 材質參考的 filterMode 為 cudaFilterModePoint 時,是向左尋找
- Kernel 使用
- 使用 linear memory 時:tex1Dfetch
- 使用 CUDA array 時:tex1D
- unbind texture (範例中省略)