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

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