纹理绑定有两种,一个是绑定到线性内存就是用cudaMalloc();cudaMemcpy();开辟的内存空间,另一种是绑定到cudaMallocArray, cudaMemcpyToArray开辟到的二维数组或者三维数组。
先说比较简单的就是绑定到cudamalloc开辟到的内存空间。
首先是纹理声明:关于纹理的声明和绑定都要在cu文件进行,在其他文件进行会比较麻烦。
首先是开辟显存和赋值给显存。
首先是纹理声明:
1、texture<float, 1, cudaReadModeElementType> texRef;
2、为显存开辟空间,并初始化
cudaMalloc((void **)&rain_table,256*4*3);
cudaMemcpy(rain_table ,table, 256*4*3,cudaMemcpyHostToDevice);
3、纹理绑定:
cudaBindTexture(0,texRef,rain_table );
4、_在global_ kernel函数使用纹理:
vterm = tex1Dfetch(texRef, 3*index);
5、解除纹理:
cudaUnbindTexture(texRef);
cudaFree(table)
以上是一维纹理的使用,接着介绍二维纹理数组的使用:
纹理声明为二维纹理:
注:在1.x和3.x的卡使用纹理回比全局变量的要好,但2.x刚好相反,因为2.X有更大的L1缓存,所以2.x的卡建议不使用纹理,除非测出的速度要比使用全局变量的要快。
1、texture<float, 2, cudaReadModeElementType> texRef;
2、开辟显存:
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0,cudaChannelFormatKindFloat);
cudaArray* cuArray;
cudaMallocArray(&cuArray, &channelDesc, 3, 256); //为cuArray开辟空间
cudaMemcpyToArray(cuArray, 0, 0, table,256*3*4, cudaMemcpyHostToDevice);
3、纹理绑定
cudaBindTextureToArray(texRef, cuArray, channelDesc);
4、使用
vterm = tex2D(texRef, 0,index);
5、解除绑定
cudaBindTextureToArray(texRef,cuArray,channelDesc);