opencl compiler crash and ignored returns causing kernel hangs on HD4000

[Michael_D_5]

Hi,

I'm pretty new to the forum, so forgive me if I post something that has been seen before.

I have encounted and found work arounds for two OpenCL bugs when working with HD400. We are working with an OpenCL ray tracer that uses an infinite loop for BVH traversal. When the intersection with a triangle is found the loop is exited using continue for traversal and return for intersection.

1. Compiler crashes when using a continue statement. I have attached a kernel  code which contains a continue statement that causes the ioc32.exe command line compiler to crash if the continue is in the code. I have worked around it using booleans. It's at line 378.

2. The second problem is that the infinite loop used for recursion will never terminate because either if conditions are ignored or return statements are ignored, I don't know which. To diagnose this problem I have switched between running on the CPU and the GPU and compare printf traces and I can see cases where the GPU loops continue on after they should have taken a return statement. I noticed that sometimes the code started to work with the printfs. So I added the following define :

#define CODE if (c) hack = get_global_id(0);  // at line 454

and placed CODE in different parts of the loop and through a process of elimination found two locations where this piece of code makes the returns work, line 587 and 599.

The attached code shows both of these additions. Strangely the CODE addtion wasn't needed for the primary ray loop, just the shadow ray loop, which is very similar. 

I am using the 2013 OpenCL beta on an Corei5. The ray tracer is now working and if you are interested you can read about the performance tuning results here :

http://fileadmin.cs.lth.se/graphics/research/papers/2012/atrt/

Mike

P.S. Why can't you upload files with the .cl extension? That's seems pretty broken for an OpenCL forum. I renamed the .cl file to .c

[Raghupathi_M_Intel]

Thanks for submitting the test case. I have reproduced this and submitted a bug. Will keep you posted if I have any update. Thanks, Raghu