Cuda By Example - 11 (Texture Memory 2-D)
跟1D一样，2D的代码也没有运行过。旧的方法看看就好。
声明二维Texture
texture<float, 2> texConstSrc;
texture<float, 2> texIn;
texture<float, 2> texOut;访问二维Texture
使用2D的Texture的便利性体现在blend_kernel函数里。不再需要通过x，y去计算一维索引。二维texture使用tex2D()去读取数据。
__global__ void blend_kernel( float *dst,bool dstOut ) {// map from threadIdx/BlockIdx to pixel positionint x = threadIdx.x + blockIdx.x * blockDim.x;int y = threadIdx.y + blockIdx.y * blockDim.y;int offset = x + y * blockDim.x * gridDim.x;float t, l, c, r, b;if (dstOut) {t = tex2D(texIn,x,y-1);l = tex2D(texIn,x-1,y);c = tex2D(texIn,x,y);r = tex2D(texIn,x+1,y);b = tex2D(texIn,x,y+1);} else {t = tex2D(texOut,x,y-1);l = tex2D(texOut,x-1,y);c = tex2D(texOut,x,y);r = tex2D(texOut,x+1,y);b = tex2D(texOut,x,y+1);}dst[offset] = c + SPEED * (t + b + r + l - 4 * c);
}然后拷贝热源数据
__global__ void copy_const_kernel( float *iptr ) {// map from threadIdx/BlockIdx to pixel positionint x = threadIdx.x + blockIdx.x * blockDim.x;int y = threadIdx.y + blockIdx.y * blockDim.y;int offset = x + y * blockDim.x * gridDim.x;float c = tex2D(texConstSrc,x,y);if (c != 0)iptr[offset] = c;
}二维Texture绑定
使用２维Texture去绑定一维数组，稍微复杂一些：
cudaChannelFormatDesc desc = cudaCreateChannelDesc<float>();
HANDLE_ERROR( cudaBindTexture2D( NULL, texConstSrc, data.dev_constSrc,desc, DIM, DIM, sizeof(float)*DIM));
HANDLE_ERROR( cudaBindTexture2D( NULL, texIn, data.dev_inSrc,desc, DIM, DIM, sizeof(float)*DIM));
HANDLE_ERROR( cudaBindTexture2D( NULL, texOut, data.dev_outSrc,desc, DIM, DIM, sizeof(float)*DIM));
解除绑定
解除绑定的方式跟１D相同
cudaUnbindTexture(texIn);
cudaUnbindTexture(texOut);
cudaUnbindTexture(texConstSrc);完整代码
#include "../common/book.h"
#include "../common/cpu_anim.h"#define DIM 1024
#define PI 3.1415926535897932f
#define MAX_TEMP 1.0f
#define MIN_TEMP 0.0001f
#define SPEED 0.25f// these exist on the GPU side
texture<float,2> texConstSrc;
texture<float,2> texIn;
texture<float,2> texOut;__global__ void blend_kernel( float *dst,bool dstOut ) {// map from threadIdx/BlockIdx to pixel positionint x = threadIdx.x + blockIdx.x * blockDim.x;int y = threadIdx.y + blockIdx.y * blockDim.y;int offset = x + y * blockDim.x * gridDim.x;float t, l, c, r, b;if (dstOut) {t = tex2D(texIn,x,y-1);l = tex2D(texIn,x-1,y);c = tex2D(texIn,x,y);r = tex2D(texIn,x+1,y);b = tex2D(texIn,x,y+1);} else {t = tex2D(texOut,x,y-1);l = tex2D(texOut,x-1,y);c = tex2D(texOut,x,y);r = tex2D(texOut,x+1,y);b = tex2D(texOut,x,y+1);}dst[offset] = c + SPEED * (t + b + r + l - 4 * c);
}__global__ void copy_const_kernel( float *iptr ) {// map from threadIdx/BlockIdx to pixel positionint x = threadIdx.x + blockIdx.x * blockDim.x;int y = threadIdx.y + blockIdx.y * blockDim.y;int offset = x + y * blockDim.x * gridDim.x;float c = tex2D(texConstSrc,x,y);if (c != 0)iptr[offset] = c;
}// globals needed by the update routine
struct DataBlock {unsigned char *output_bitmap;float *dev_inSrc;float *dev_outSrc;float *dev_constSrc;CPUAnimBitmap *bitmap;cudaEvent_t start, stop;float totalTime;float frames;
};void anim_gpu( DataBlock *d, int ticks ) {HANDLE_ERROR( cudaEventRecord( d->start, 0 ) );dim3 blocks(DIM/16,DIM/16);dim3 threads(16,16);CPUAnimBitmap *bitmap = d->bitmap;// since tex is global and bound, we have to use a flag to// select which is in/out per iterationvolatile bool dstOut = true;for (int i=0; i<90; i++) {float *in, *out;if (dstOut) {in = d->dev_inSrc;out = d->dev_outSrc;} else {out = d->dev_inSrc;in = d->dev_outSrc;}copy_const_kernel<<<blocks,threads>>>( in );blend_kernel<<<blocks,threads>>>( out, dstOut );dstOut = !dstOut;}float_to_color<<<blocks,threads>>>( d->output_bitmap,d->dev_inSrc );HANDLE_ERROR( cudaMemcpy( bitmap->get_ptr(),d->output_bitmap,bitmap->image_size(),cudaMemcpyDeviceToHost ) );HANDLE_ERROR( cudaEventRecord( d->stop, 0 ) );HANDLE_ERROR( cudaEventSynchronize( d->stop ) );float elapsedTime;HANDLE_ERROR( cudaEventElapsedTime( &elapsedTime,d->start, d->stop ) );d->totalTime += elapsedTime;++d->frames;printf( "Average Time per frame: %3.1f ms\n",d->totalTime/d->frames );
}// clean up memory allocated on the GPU
void anim_exit( DataBlock *d ) {cudaUnbindTexture( texIn );cudaUnbindTexture( texOut );cudaUnbindTexture( texConstSrc );HANDLE_ERROR( cudaFree( d->dev_inSrc ) );HANDLE_ERROR( cudaFree( d->dev_outSrc ) );HANDLE_ERROR( cudaFree( d->dev_constSrc ) );HANDLE_ERROR( cudaEventDestroy( d->start ) );HANDLE_ERROR( cudaEventDestroy( d->stop ) );
}int main( void ) {DataBlock data;CPUAnimBitmap bitmap( DIM, DIM, &data );data.bitmap = &bitmap;data.totalTime = 0;data.frames = 0;HANDLE_ERROR( cudaEventCreate( &data.start ) );HANDLE_ERROR( cudaEventCreate( &data.stop ) );int imageSize = bitmap.image_size();HANDLE_ERROR( cudaMalloc( (void**)&data.output_bitmap,imageSize ) );// assume float == 4 chars in size (ie rgba)HANDLE_ERROR( cudaMalloc( (void**)&data.dev_inSrc,imageSize ) );HANDLE_ERROR( cudaMalloc( (void**)&data.dev_outSrc,imageSize ) );HANDLE_ERROR( cudaMalloc( (void**)&data.dev_constSrc,imageSize ) );cudaChannelFormatDesc desc = cudaCreateChannelDesc<float>();HANDLE_ERROR( cudaBindTexture2D( NULL, texConstSrc,data.dev_constSrc,desc, DIM, DIM,sizeof(float) * DIM ) );HANDLE_ERROR( cudaBindTexture2D( NULL, texIn,data.dev_inSrc,desc, DIM, DIM,sizeof(float) * DIM ) );HANDLE_ERROR( cudaBindTexture2D( NULL, texOut,data.dev_outSrc,desc, DIM, DIM,sizeof(float) * DIM ) );// initialize the constant datafloat *temp = (float*)malloc( imageSize );for (int i=0; i<DIM*DIM; i++) {temp[i] = 0;int x = i % DIM;int y = i / DIM;if ((x>300) && (x<600) && (y>310) && (y<601))temp[i] = MAX_TEMP;}temp[DIM*100+100] = (MAX_TEMP + MIN_TEMP)/2;temp[DIM*700+100] = MIN_TEMP;temp[DIM*300+300] = MIN_TEMP;temp[DIM*200+700] = MIN_TEMP;for (int y=800; y<900; y++) {for (int x=400; x<500; x++) {temp[x+y*DIM] = MIN_TEMP;}}HANDLE_ERROR( cudaMemcpy( data.dev_constSrc, temp,imageSize,cudaMemcpyHostToDevice ) ); // initialize the input datafor (int y=800; y<DIM; y++) {for (int x=0; x<200; x++) {temp[x+y*DIM] = MAX_TEMP;}}HANDLE_ERROR( cudaMemcpy( data.dev_inSrc, temp,imageSize,cudaMemcpyHostToDevice ) );free( temp );bitmap.anim_and_exit( (void (*)(void*,int))anim_gpu,(void (*)(void*))anim_exit );
}