Issues with expanding Intel FPGA Sobel Filter example to include Median Filter.

[Altera_Forum]

Hi,  

 I am currently developing on the FPGA Cyclone V board. I have successfully ran the 3x3 sobel filter kernel example, and have expanded this to include a 5x5 example successfully also. I decided I would like to include a median filter to remove noise prior to doing the 5x5 sobel filter. I have also successfully got an output from the median filter. However when I try to implement the two, I get strange results. The code is below.    # include "../host/inc/imageDims.h" __kernel void laplacian33(global unsigned int * restrict frame_in, global unsigned char * restrict frame_out, const int totalpixels) { //Define convolution matrix int h_n_x = {{-1,-2,-1},{0,0,0},{1,2,1}}; int h_n_y = {{-1,0,1},{-2,0,2},{-1,0,1}}; char set_white = 0xff; char set_black = 0x00; // Pixel buffer of 2 columns and 3 extra pixels int pix_buff; // Initial iterations initialize the pixel buffer int count = -(3 * COLS + 3); //Iterator of the output index for greyscale PGM image int outIndex = 0; while (count != totalpixels) { //Fill pixel buffer # pragma unroll for (int i = 3 * COLS + 2; i > 0; --i) { pix_buff = pix_buff; } pix_buff = count >= 0 ? frame_in : 0; //Initialise greyscale variable. char grey_scale_x = 0x00; char grey_scale_y = 0x00; //Compute one convolution each cycle # pragma unroll for (int filterRow = 0; filterRow < 3; ++filterRow) { # pragma unroll for (int filterCol = 0; filterCol < 3; ++filterCol) { //Get the current pixel unsigned int pixel = pix_buff; unsigned char b = pixel & 0xff0000; unsigned char g = pixel & 0x00ff00; unsigned char r = pixel & 0x0000ff; //Approximate version, simply apply a ratio and shift unsigned char luma_apprx = ((2 * r) + b + (3 * g)) >> 3; //Typecast luma as a char and perform convolution grey_scale_x += (char)luma_apprx*h_n_x; grey_scale_y += (char)luma_apprx*h_n_y; } } int sobel_mag = abs(grey_scale_x) + abs(grey_scale_y); if(sobel_mag > 0xff) sobel_mag = set_white; else if(sobel_mag < 0x00) sobel_mag = set_black; if(outIndex != totalpixels) { frame_out = sobel_mag; //Write byte, iterate index } count++; //Iterate overall count } } __kernel void laplacian55(global unsigned int * restrict frame_in, global unsigned char * restrict frame_out, global unsigned char * restrict frame_median, const int totalpixels) { //Define convolution matrix int h_n_x = {{2,2,4,2,2}, {1,1,2,1,1}, {0,0,0,0,0}, {-1,-1,-2,-1,-1}, {-2,-2,-4,-2,-2}}; int h_n_y = {{2,1,0,-1,-2}, {2,1,0,-1,-2}, {4,2,0,-2,-4}, {2,1,0,-1,-2}, {2,1,0,-1,-2}}; unsigned char median_matrix; int height = 5; int width = 5; int value = 0; int median = 0; unsigned char array; char set_white = 0xff; char set_black = 0x00; // Pixel buffer. int pix_buff; int pix_buff_2; // Initial iterations initialize the pixel buffer int count = -(5 * COLS + 3); int sobelCount = -(5 * COLS + 3); //Iterator of the output index for greyscale PGM image int sobel_outIndex = 0; int outIndex = 0; //Initialise greyscale variable. char grey_scale_x = 0x00; char grey_scale_y = 0x00; int filterRow; int filterCol; unsigned char sobelPixel = 0x00; while (count != totalpixels) { //Fill pixel buffer # pragma unroll for (int i = 5 * COLS + 2; i > 0; --i) { pix_buff = pix_buff; } pix_buff = count >= 0 ? frame_in : 0; //Get values and find median each cycle # pragma unroll for (filterRow = 0; filterRow < 5; ++filterRow) { # pragma unroll for (filterCol = 0; filterCol < 5; ++filterCol) { //Get the current pixel unsigned int pixel = pix_buff; //Retrieve the individual bytes of pixel by masking. unsigned char b = pixel & 0xff0000; unsigned char g = pixel & 0x00ff00; unsigned char r = pixel & 0x0000ff; //Convert 3 bytes of RGB --> 1 byte of greyscale data unsigned char luma_apprx = ((2 * r) + b + (3 * g)) >> 3; //Store appropriate grey level in median matrix median_matrix = luma_apprx; } } //Convert 2d array to 1d vector for(int rows = 0; rows < 5; rows++) { for(int cols = 0; cols < 5; cols++) { array = median_matrix; } } //Sort array for(int x = 0; x < 25; x++){ for(int y = 0; y < 24; y++){ if(array>array){ int temp = array; array = array; array = temp; } } } //Get median median = array; //Write the total number of greyscale bytes if(outIndex != totalpixels) { frame_median = median; //Write byte, iterate index } count++; //Iterate overall count } //Do the same, but this time use sobel filter.. while (sobelCount != totalpixels) { //Initialise pixel buffer for sobel # pragma unroll for (int z = 5 * COLS + 2; z > 0; --z) { pix_buff_2 = pix_buff_2; } //Fill buffer with median data pix_buff_2 = sobelCount >= 0 ? frame_median : 0; //Compute one convolution each cycle # pragma unroll for (filterRow = 0; filterRow < 5; ++filterRow) { # pragma unroll for (filterCol = 0; filterCol < 5; ++filterCol) { //Get the current pixel sobelPixel = pix_buff_2; grey_scale_x += (char)sobelPixel*h_n_x; grey_scale_y += (char)sobelPixel*h_n_y; } } //Threshold int magnitude = abs(grey_scale_x) + abs(grey_scale_y); if(magnitude > 255) magnitude = set_white; else if(magnitude < 0) magnitude = set_black; //Write the total number of greyscale bytes if(sobel_outIndex != totalpixels) { frame_out = magnitude; } sobelCount++; //Iterate overall count } }   I think I have pinpointed the problem to be with the line of code:   sobelPixel = pix_buff_2;   Since changing this makes the output look different each time. I think this is because of the "boundary pixels", so the image reduces in size after the median filter, and there are no longer COLS number of columns in the resulting image to be processed by the sobel filter. However, I am struggling to find the correct algorithm to index the "current pixel" to get the correct filtered output. Could anyone please assist with this bug?  Kind regards