GPU Pro 5

5. Non-separable 2D, 3D, and 4D Filtering with CUDA 479

#d e f i n e HALO 8

__global__ void Convolution_2D_Shared ( float Filter_Response ,

float Image , int DATA_W , int DATA_H)

{

int x = blockIdx . x VALID_RESPONSES_X + threadIdx . x ;

int y = blockIdx . y VALID_RESPONSES_Y + threadIdx . y ;

__shared__ float s_Image [64][96]; // y , x

// Reset shared memory

s_Image [ threadIdx . y ][ threadIdx . x ]=0.0f ;

s_Image [ threadIdx . y ][ threadIdx . x + 32] = 0.0f ;

s_Image [ threadIdx . y ][ threadIdx . x + 64] = 0.0f ;

s_Image [ threadIdx . y +32][threadIdx . x ]=0.0f ;

s_Image [ threadIdx . y +32][threadIdx . x + 32] = 0.0f ;

s_Image [ threadIdx . y +32