SobelFilter, different results from C and OpenCL Implementations

Altera_Forum · ‎07-27-2015

Hello everyone,

I tried to write an OpenCL code to implement the Sobel Filter. But I got strange results. It does detect the edge, but it looks like a distracted version.

http://www.alteraforum.com/forum/attachment.php?attachmentid=10933&stc=1

Then I tried to verfy it in C with the same code (expect for loop, etc.) and got a perfect result.

http://www.alteraforum.com/forum/attachment.php?attachmentid=10934&stc=1

The OpenCL version seems to be distracted in pixels and is slightly larger than the C version. Also I got different results when I used differnet data type (float ot int) in kernel.

There is a detail when see the OpenCL version in detail. There are a lot of square patterns, which maybe a clue to solve the problem.

http://www.alteraforum.com/forum/attachment.php?attachmentid=10935&stc=1

==========================UPDATE============================

The Kernel code: (the 3 in the code is because for each pixel, RGB information is read one by one seperately.)


__kernel void SobelFilter(__global float* src, __global float* dst, int width, int height) {          
   int x = (int)get_global_id(0);     
   int y = (int)get_global_id(1);          
   int W = width;      
   int H = height;
   
   if(x > 2 && x < 3*W - 2 && y > 2 && y < H - 2){           
      int p00 = (int)src;           
      int p10 = (int)src;          
      int p20 = (int)src;        
      int p01 = (int)src;     
      int p21 = (int)src;    
      int p02 = (int)src;    
      int p12 = (int)src;     
      int p22 = (int)src;       
      
      int gx = -p00 + p20 + 2*(p21 - p01) - p02 + p22;    
      int gy = -p00 - p20 + 2*(p12 - p10) + p02 + p22;             
     
      int g = sqrt(gx*gx+ gy*gy);    
       
      dst = (float)g;     
 }          
   else{      
     dst = 0.0f;    
   } 
}

Can any expert help me and tell where may go wrong? Thank you in advance!

Altera_Forum · ‎07-27-2015

What do you mean by "Then I tried to verfy it in C with the same code (expect for loop, etc.) and got a perfect result" It would help to see what changes you made.

Altera_Forum · ‎07-27-2015

--- Quote Start ---

What do you mean by "Then I tried to verfy it in C with the same code (expect for loop, etc.) and got a perfect result" It would help to see what changes you made.

--- Quote End ---

I just want to verify if the code or the algorithm is right. So I write a C code which is almost identical to the kernel, the only differnce is that I added two loops. Then for the result of the C version, it looks perfect. I didn't change the core of the code.

Altera_Forum · ‎07-27-2015

So a number of reasons might be the problem. The first is how you're sending and reading data to and from the kernel. Make sure the data that you read into the kernel is what you expect, the same with when you're reading from the kernel. Other problems maybe how you partition the work items and work groups. The order of operations if some operations depend on previous data can be out of place and thus give you the wrong results.

Altera_Forum · ‎07-27-2015

I don't think it is caused by the read and write of the data. I tried to make the output of the kernel just equal to the input, then the image is perfectly recreated.

I also tried different worksize, the results remain the same. And the input data comes from the source image, and the result id written to the destination. So I also don't think there are dependency.

These are my thoughts. Please correct me, if I am wrong. Thank you.

Altera_Forum · ‎07-27-2015

It's hard to tell what's wrong from your description. One suggestion is to use printf in the kernel to see the execution path. Make sure the order of the operations through the pipeline is correct.