CUDA纹理使用方法.docVIP

一、例子1：fluidsGL中： （一）host端 1）声明实际纹理数据 float2 *hvfield = NULL; float2 *dvfield = NULL; 2) 对host数据和device数据进行初始化 hvfield = (cData*)malloc(sizeof(cData) * DS); memset(hvfield, 0, sizeof(cData) * DS); // Allocate and initialize device data cudaMallocPitch((void**)dvfield, tPitch, sizeof(cData)*DIM, DIM); cudaMemcpy(dvfield, hvfield, sizeof(cData) * DS, cudaMemcpyHostToDevice); 调用device端函数，建立并绑定纹理数据 setupTexture(DIM, DIM); bindTexture(); //注意,其中将array和纹理标识texref绑定 （二）device端 1)声明纹理对象标识，以及纹理实际数据数组 texturefloat2, 2 texref; static cudaArray *array = NULL; 2）实现纹理建立、绑定函数 void setupTexture(int x, int y) { // Wrap mode appears to be the new default texref.filterMode = cudaFilterModeLinear; cudaChannelFormatDesc desc = cudaCreateChannelDescfloat2(); cudaMallocArray(array, desc, y, x); CUT_CHECK_ERROR(cudaMalloc failed); } void bindTexture(void) { cudaBindTextureToArray(texref, array); CUT_CHECK_ERROR(cudaBindTexture failed); } void unbindTexture(void) { cudaUnbindTexture(texref); CUT_CHECK_ERROR(cudaUnbindTexture failed); } //传入数据 void updateTexture(cData *data, size_t wib, size_t h, size_t pitch) { cudaMemcpy2DToArray(array, 0, 0, data, pitch, wib, h, cudaMemcpyDeviceToDevice); CUT_CHECK_ERROR(cudaMemcpy failed); } void deleteTexture(void) { cudaFreeArray(array); CUT_CHECK_ERROR(cudaFreeArray failed); } 3）访问纹理数据 __global__ void advectVelocity_k(cData *v, float *vx, float *vy, int dx, int pdx, int dy, float dt, int lb) { int gtidx = blockIdx.x * blockDim.x + threadIdx.x; int gtidy = blockIdx.y * (lb * blockDim.y) + threadIdx.y * lb; int p; cData vterm, ploc; float vxterm, vyterm; // gtidx is the domain location in x for this thread if (gtidx dx) { for (p = 0; p lb; p++) { // fi is the domain location in y for this thread int fi = gtidy + p; if (fi dy) { int fj = fi * pdx + gtidx;