开发者

Downsampling image using Texture Memory

开发者 https://www.devze.com 2023-02-14 06:54 出处:网络
I need to downsample an image.Some reading suggested me that if I use texture memory then this function comes on free and is faster(I am looking for a bilinear interpolation).Could some one tell me ho

I need to downsample an image.Some reading suggested me that if I use texture memory then this function comes on free and is faster(I am looking for a bilinear interpolation).Could some one tell me how exactly to write the kernel for this ?Here is what I have currently : (I use (1,1) thread blocks)

__global__ void texturekernel( int * final_red){
     int f = (blockIdx.x * blockDim.x) + threadIdx.x;
     int c = (blockIdx.y * blockDim.y) + threadIdx.y;
     int id=blockIdx.x+256*blockIdx.y;//256 is the width of downsampled image ..original was 512
     final_red[id]=tex2D( refTexture,c+0.5f,f+0.5f);//This is just for the red channel
     //where reftexture is defined as texture <float, 2, cudaReadModeElementType> refTexture;

  };

This version currently gives me all 0's in the output.

Edited (In this version I am trying to downsample 2 2000*512 size images into 2 1000*256):

  texture <float, 2, cudaReadModeElementType> refTexture; // global variable !
    cudaArray* myArray; 
    cudaChannelFormatDesc description = cudaCreateChannelDesc<float>();
    cudaError rs=cudaMallocArray (   &myArray,&description, 512,2000*2);//



//This line below is part of loop where input image is read row by row ..rowchecker keeps track of the row    

cudaMemcpyToArray(myArray,0,rowchecker++,array_temp_开发者_开发知识库red,sizeof(int)*test_columns,cudaMemcpHostToDevice);

     refTexture.normalized=false;
     refTexture.addressMode[0]=cudaAddressModeClamp;
     refTexture.addressMode[1]=cudaAddressModeClamp;
     refTexture.filterMode=cudaFilterModePoint;

     cudaBindTextureToArray( refTexture,myArray);

              dim3 blockSize(1,1);
              int n_blocks_x=256;
              int n_blocks_y=1000*2;
              dim3 gridSize(n_blocks_x,n_blocks_y);
              cudaMalloc((void**)&finalarray,(2000)*(512)*2/4*sizeof(int));
              texturekernel<<<gridSize,blockSize>>>(finalarray );


int id=blockIdx.x+256*blockIdx.y;

This statement crosses the limits of your final_red.

Try this:

__global__ void texturekernel( int * final_red){

     int f = blockIdx.x * blockDim.x + threadIdx.x;
     int c = blockIdx.y * blockDim.y + threadIdx.y;

     int id =c/2 * 256 + f/2;

     final_red[id] = tex2D( refTexture,c+0.5f,f+0.5f);

  }
0

精彩评论

暂无评论...
验证码 换一张
取 消