- Mark as New
- Bookmark
- Subscribe
- Mute
- Subscribe to RSS Feed
- Permalink
- Report Inappropriate Content
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
Link Copied
- Mark as New
- Bookmark
- Subscribe
- Mute
- Subscribe to RSS Feed
- Permalink
- Report Inappropriate Content
![](/skins/images/872293744008A34B36F8ABF94A46CC66/responsive_peak/images/icon_anonymous_message.png)
- Subscribe to RSS Feed
- Mark Topic as New
- Mark Topic as Read
- Float this Topic for Current User
- Bookmark
- Subscribe
- Printer Friendly Page