I am implementing bilateral filter on GPU using brute force method. I am doing the following things in mu code:
(i) copying input image to texture memory
(ii) i have calculated spatial gaussian kernel in host code itself and passing it as input to kernel code.
(iii) passing output image as input to kernel code
(iv) calculating output pixel value by multiplying with gaussian kernel and range filter values. Output image size has been assigned the same size as input image.
But i am getting segmentation error while assigning the output value image. Also i m getting 0 as output image value, when i am copying using cudaMemcpy function. Please help me out as I have my course presentation next week, need to show the working code [:(]… I am pasting my kernel code here:
//Kernel
texture<float, 2, cudaReadModeElementType> tex;
global void BilateralFilter(float* outputImage, float* d_kernel, int kernelRad, float sigmar, int width, int height) {
int i = blockIdx.x*blockDim.x + threadIdx.x;
int j = blockIdx.y*blockDim.y + threadIdx.y;
float in = (float) i/width;
float jn = (float) j/height;
float ival0 = tex2D(tex, in, jn);
float ival = 0.0;
float rangFilt = 0.0;
float weight = 0.0;
float out = 0.0;
float temp = 0.0;
for (int ti = -kernelRad; ti < kernelRad; ti++) {
for (int tj = -kernelRad; tj < kernelRad; tj++) {
in = (float) (i + ti)/width;
in = fmax(0.0, in);
in = fmin(in, 1.0);
jn = (float) (j + tj)/height;
jn = fmax(0.0, jn);
jn = fmin(1.0, jn); //for restricting the values outside prescribed range
ival = tex2D(tex, in, jn);
rangFilt = expf(-(ival - ival0)*(ival - ival0)/(sigmar*sigmar));
temp = rangFilt*d_kernel[(ti + kernelRad) + (2*kernelRad + 1)*(tj + kernelRad)];
out += temp*ival;
weight += temp;
}
}
outputImage[i + j*width] = (float) out/weight;
}