diff --git a/.gitignore b/.gitignore
index 98b65133..a94075b1 100644
--- a/.gitignore
+++ b/.gitignore
@@ -1,3 +1,4 @@
+build
 *.vcxproj
 App/App.vcxproj.user
 *.filters
diff --git a/CLW/CLWImage2D.cpp b/CLW/CLWImage2D.cpp
index 8b823788..026f225b 100644
--- a/CLW/CLWImage2D.cpp
+++ b/CLW/CLWImage2D.cpp
@@ -53,12 +53,12 @@ CLWImage2D CLWImage2D::Create(cl_context context, cl_image_format const* imgForm
 }
 
 
-CLWImage2D CLWImage2D::CreateFromGLTexture(cl_context context, cl_GLint texture)
+CLWImage2D CLWImage2D::CreateFromGLTexture(cl_context context, cl_GLint texture, cl_mem_flags flags)
 {
     cl_int status = CL_SUCCESS;
 
     // TODO: handle that gracefully: GL_TEXTURE_2D
-    cl_mem deviceImg = clCreateFromGLTexture(context, CL_MEM_WRITE_ONLY, 0x0DE1, 0, texture, &status);
+    cl_mem deviceImg = clCreateFromGLTexture(context, flags, 0x0DE1, 0, texture, &status);
     
     ThrowIf(status != CL_SUCCESS, status, "clCreateFromGLTexture failed");
     
@@ -69,6 +69,20 @@ CLWImage2D CLWImage2D::CreateFromGLTexture(cl_context context, cl_GLint texture)
     return image;
 }
 
+CLWImage2D CLWImage2D::CreateFromGLRenderbuffer(cl_context context, cl_GLint buffer, cl_mem_flags flags) {
+    cl_int status = CL_SUCCESS;
+
+    cl_mem deviceImg = clCreateFromGLRenderbuffer(context, flags, buffer, &status);
+
+    ThrowIf(status != CL_SUCCESS, status, "clCreateFromGLRenderbuffer failed");
+
+    CLWImage2D image(deviceImg);
+
+    clReleaseMemObject(deviceImg);
+
+    return image;
+}
+
 CLWImage2D::CLWImage2D(cl_mem image)
 : ReferenceCounter<cl_mem, clRetainMemObject, clReleaseMemObject>(image)
 {
diff --git a/CLW/CLWImage2D.h b/CLW/CLWImage2D.h
index f06e42be..761afec7 100644
--- a/CLW/CLWImage2D.h
+++ b/CLW/CLWImage2D.h
@@ -45,7 +45,8 @@ class CLWImage2D : public ReferenceCounter<cl_mem, clRetainMemObject, clReleaseM
 {
 public:
     static CLWImage2D Create(cl_context context, cl_image_format const* imgFormat, size_t width, size_t height, size_t rowPitch);
-    static CLWImage2D CreateFromGLTexture(cl_context context, cl_GLint texture);
+    static CLWImage2D CreateFromGLTexture(cl_context context, cl_GLint texture, cl_mem_flags flags = CL_MEM_WRITE_ONLY);
+    static CLWImage2D CreateFromGLRenderbuffer(cl_context context, cl_GLint buffer, cl_mem_flags flags = CL_MEM_WRITE_ONLY);
 
     CLWImage2D(){}
     virtual ~CLWImage2D();
diff --git a/RadeonRays/src/kernels/CL/build_hlbvh.cl b/RadeonRays/src/kernels/CL/build_hlbvh.cl
index 6ffd720f..53809d6c 100644
--- a/RadeonRays/src/kernels/CL/build_hlbvh.cl
+++ b/RadeonRays/src/kernels/CL/build_hlbvh.cl
@@ -38,7 +38,7 @@ THE SOFTWARE.
 /*************************************************************************
 INCLUDES
 **************************************************************************/
-#include <../RadeonRays/src/kernels/CL/common.cl>
+#include <common.cl>
 
 /*************************************************************************
 DEFINES
diff --git a/RadeonRays/src/kernels/CL/common.cl b/RadeonRays/src/kernels/CL/common.cl
index 1692feeb..40313b14 100644
--- a/RadeonRays/src/kernels/CL/common.cl
+++ b/RadeonRays/src/kernels/CL/common.cl
@@ -68,6 +68,11 @@ typedef struct
     float4 uvwt;
 } Intersection;
 
+typedef struct
+{
+    int shape_id;
+    int prim_id;
+} Occlusion;
 
 /*************************************************************************
 HELPER FUNCTIONS
diff --git a/RadeonRays/src/kernels/CL/intersect_bvh2_bittrail.cl b/RadeonRays/src/kernels/CL/intersect_bvh2_bittrail.cl
index 8a517ef5..3fa0a189 100644
--- a/RadeonRays/src/kernels/CL/intersect_bvh2_bittrail.cl
+++ b/RadeonRays/src/kernels/CL/intersect_bvh2_bittrail.cl
@@ -89,7 +89,7 @@ THE SOFTWARE.
 /*************************************************************************
 INCLUDES
 **************************************************************************/
-#include <../RadeonRays/src/kernels/CL/common.cl>
+#include <common.cl>
 
 
 /*************************************************************************
@@ -147,7 +147,7 @@ occluded_main(
     // Displacement table size
     int const displacement_table_size,
     // Hit results: 1 for hit and -1 for miss
-    GLOBAL int* hits
+    GLOBAL Occlusion* hits
     )
 {
     int global_id = get_global_id(0);
@@ -193,7 +193,8 @@ occluded_main(
                     // If hit store the result and bail out
                     if (f < t_max)
                     {
-                        hits[global_id] = HIT_MARKER;
+                        hits[global_id].shape_id = face.shape_id;
+                        hits[global_id].prim_id = face.prim_id;
                         return;
                     }
                 }
@@ -267,7 +268,8 @@ occluded_main(
             }
 
             // Finished traversal, but no intersection found
-            hits[global_id] = MISS_MARKER;
+            hits[global_id].shape_id = MISS_MARKER;
+            hits[global_id].prim_id = MISS_MARKER;
         }
     }
 }
diff --git a/RadeonRays/src/kernels/CL/intersect_bvh2_lds.cl b/RadeonRays/src/kernels/CL/intersect_bvh2_lds.cl
index 581c7625..73388bfb 100644
--- a/RadeonRays/src/kernels/CL/intersect_bvh2_lds.cl
+++ b/RadeonRays/src/kernels/CL/intersect_bvh2_lds.cl
@@ -23,7 +23,7 @@ THE SOFTWARE.
 /*************************************************************************
 INCLUDES
 **************************************************************************/
-#include <../RadeonRays/src/kernels/CL/common.cl>
+#include <common.cl>
 
 /*************************************************************************
 TYPE DEFINITIONS
@@ -229,7 +229,7 @@ KERNEL void occluded_main(
     // Stack memory
     GLOBAL uint *stack,
     // Hit results: 1 for hit and -1 for miss
-    GLOBAL int *hits)
+    GLOBAL Occlusion *hits)
 {
     __local uint lds_stack[GROUP_SIZE * LDS_STACK_SIZE];
 
@@ -322,7 +322,8 @@ KERNEL void occluded_main(
 
                     if (t < closest_t)
                     {
-                        hits[index] = HIT_MARKER;
+                        hits[index].shape_id = GetMeshId(node);
+                        hits[index].prim_id = GetPrimId(node);
                         return;
                     }
                 }
@@ -343,7 +344,8 @@ KERNEL void occluded_main(
             }
 
             // Finished traversal, but no intersection found
-            hits[index] = MISS_MARKER;
+            hits[index].shape_id = MISS_MARKER;
+            hits[index].prim_id = MISS_MARKER;
         }
     }
 }
diff --git a/RadeonRays/src/kernels/CL/intersect_bvh2_lds_fp16.cl b/RadeonRays/src/kernels/CL/intersect_bvh2_lds_fp16.cl
index d57812f9..0a24de50 100644
--- a/RadeonRays/src/kernels/CL/intersect_bvh2_lds_fp16.cl
+++ b/RadeonRays/src/kernels/CL/intersect_bvh2_lds_fp16.cl
@@ -25,7 +25,7 @@ THE SOFTWARE.
 /*************************************************************************
 INCLUDES
 **************************************************************************/
-#include <../RadeonRays/src/kernels/CL/common.cl>
+#include <common.cl>
 
 /*************************************************************************
 TYPE DEFINITIONS
@@ -317,7 +317,7 @@ KERNEL void occluded_main(
     // Stack memory
     GLOBAL uint *stack,
     // Hit results: 1 for hit and -1 for miss
-    GLOBAL int *hits)
+    GLOBAL Occlusion *hits)
 {
     __local uint lds_stack[GROUP_SIZE * LDS_STACK_SIZE];
 
@@ -436,7 +436,8 @@ KERNEL void occluded_main(
 
                     if (t < closest_t)
                     {
-                        hits[index] = HIT_MARKER;
+                        hits[index].prim_id = node.aabb23_min_or_v2_and_addr2_or_prim_id.w;
+                        hits[index].shape_id = node.aabb01_max_or_v1_and_addr1_or_mesh_id.w;
                         return;
                     }
                 }
@@ -457,7 +458,8 @@ KERNEL void occluded_main(
             }
 
             // Finished traversal, but no intersection found
-            hits[index] = MISS_MARKER;
+            hits[index].shape_id = MISS_MARKER;
+            hits[index].prim_id = MISS_MARKER;
         }
     }
 }
diff --git a/RadeonRays/src/kernels/CL/intersect_bvh2_short_stack.cl b/RadeonRays/src/kernels/CL/intersect_bvh2_short_stack.cl
index 6a7fac72..4871b1de 100644
--- a/RadeonRays/src/kernels/CL/intersect_bvh2_short_stack.cl
+++ b/RadeonRays/src/kernels/CL/intersect_bvh2_short_stack.cl
@@ -82,7 +82,7 @@ THE SOFTWARE.
 /*************************************************************************
 INCLUDES
 **************************************************************************/
-#include <../RadeonRays/src/kernels/CL/common.cl>
+#include <common.cl>
 
 
 /*************************************************************************
@@ -139,7 +139,7 @@ occluded_main(
     // Stack memory
     GLOBAL int* stack,
     // Hit results: 1 for hit and -1 for miss
-    GLOBAL int* hits
+    GLOBAL Occlusion* hits
     )
 {
     // Allocate stack in LDS
@@ -197,7 +197,8 @@ occluded_main(
                     // If hit update closest hit distance and index
                     if (f < t_max)
                     {
-                        hits[global_id] = HIT_MARKER;
+                        hits[global_id].shape_id = node.shape_id;
+                        hits[global_id].prim_id = node.prim_id;
                         return;
                     }
                 }
@@ -275,7 +276,8 @@ occluded_main(
             }
 
             // Finished traversal, but no intersection found
-            hits[global_id] = MISS_MARKER;
+            hits[global_id].shape_id = MISS_MARKER;
+            hits[global_id].prim_id = MISS_MARKER;
         }
     }
 }
diff --git a/RadeonRays/src/kernels/CL/intersect_bvh2_skiplinks.cl b/RadeonRays/src/kernels/CL/intersect_bvh2_skiplinks.cl
index 469b4f76..4a7c1099 100644
--- a/RadeonRays/src/kernels/CL/intersect_bvh2_skiplinks.cl
+++ b/RadeonRays/src/kernels/CL/intersect_bvh2_skiplinks.cl
@@ -61,7 +61,8 @@ THE SOFTWARE.
 /*************************************************************************
  INCLUDES
  **************************************************************************/
-#include <../RadeonRays/src/kernels/CL/common.cl>
+#include <common.cl>
+
 /*************************************************************************
 EXTENSIONS
 **************************************************************************/
@@ -211,7 +212,7 @@ void occluded_main(
     // Number of rays
     GLOBAL int const* restrict num_rays,
     // Hit data
-    GLOBAL int* hits
+    GLOBAL Occlusion* hits
 )
 {
     int global_id = get_global_id(0);
@@ -256,7 +257,8 @@ void occluded_main(
                         // If hit store the result and bail out
                         if (f < t_max)
                         {
-                            hits[global_id] = HIT_MARKER;
+                            hits[global_id].shape_id = face.shape_id;
+                            hits[global_id].prim_id = face.prim_id;
                             return;
                         }
                     }
@@ -273,7 +275,8 @@ void occluded_main(
             }
 
             // Finished traversal, but no intersection found
-            hits[global_id] = MISS_MARKER;
+            hits[global_id].shape_id = MISS_MARKER;
+            hits[global_id].prim_id = MISS_MARKER;
         }
     }
 }
\ No newline at end of file
diff --git a/RadeonRays/src/kernels/CL/intersect_bvh2level_skiplinks.cl b/RadeonRays/src/kernels/CL/intersect_bvh2level_skiplinks.cl
index 7ac62776..4a4e73c6 100644
--- a/RadeonRays/src/kernels/CL/intersect_bvh2level_skiplinks.cl
+++ b/RadeonRays/src/kernels/CL/intersect_bvh2level_skiplinks.cl
@@ -50,7 +50,7 @@ THE SOFTWARE.
 /*************************************************************************
 INCLUDES
 **************************************************************************/
-#include <../RadeonRays/src/kernels/CL/common.cl>
+#include <common.cl>
 
 /*************************************************************************
 EXTENSIONS
@@ -320,7 +320,7 @@ KERNEL void occluded_main(
     // Number of rays in ray buffer
     GLOBAL int const* restrict num_rays,
     // Hits 
-    GLOBAL int* hits
+    GLOBAL Occlusion* hits
 )
 {
     int global_id = get_global_id(0);
@@ -374,7 +374,8 @@ KERNEL void occluded_main(
                             // If hit update closest hit distance and index
                             if (f < t_max)
                             {
-                                hits[global_id] = HIT_MARKER;
+                                hits[global_id].shape_id = face.shape_id;
+                                hits[global_id].prim_id = face.prim_id;
                                 return;
                             }
 
@@ -443,7 +444,8 @@ KERNEL void occluded_main(
                 }
             }
 
-            hits[global_id] = MISS_MARKER;
+            hits[global_id].shape_id = MISS_MARKER;
+            hits[global_id].prim_id = MISS_MARKER;
         }
     }
 }
\ No newline at end of file
diff --git a/RadeonRays/src/kernels/CL/intersect_hlbvh_stack.cl b/RadeonRays/src/kernels/CL/intersect_hlbvh_stack.cl
index c0cd1f5b..9d8fd0c1 100644
--- a/RadeonRays/src/kernels/CL/intersect_hlbvh_stack.cl
+++ b/RadeonRays/src/kernels/CL/intersect_hlbvh_stack.cl
@@ -39,7 +39,7 @@ THE SOFTWARE.
  /*************************************************************************
   INCLUDES
   **************************************************************************/
-#include <../RadeonRays/src/kernels/CL/common.cl>
+#include <common.cl>
 
  /*************************************************************************
    EXTENSIONS
@@ -94,7 +94,7 @@ occluded_main(
     // Stack memory
     GLOBAL int* stack,
     // Hit results: 1 for hit and -1 for miss
-    GLOBAL int* hits
+    GLOBAL Occlusion* hits
     )
 {
     int global_id = get_global_id(0);
@@ -151,7 +151,8 @@ occluded_main(
                     // If hit update closest hit distance and index
                     if (f < t_max)
                     {
-                        hits[global_id] = HIT_MARKER;
+                        hits[global_id].shape_id = face.shape_id;
+                        hits[global_id].prim_id = face.prim_id;
                         return;
                     }
                 }
@@ -229,7 +230,8 @@ occluded_main(
             }
 
             // Finished traversal, but no intersection found
-            hits[global_id] = MISS_MARKER;
+            hits[global_id].shape_id = MISS_MARKER;
+            hits[global_id].prim_id = MISS_MARKER;
         }
     }
 }
diff --git a/Tools/scripts/stringify.py b/Tools/scripts/stringify.py
index cee35893..cb9abc5a 100644
--- a/Tools/scripts/stringify.py
+++ b/Tools/scripts/stringify.py
@@ -8,9 +8,12 @@ def printfile(filename, dir):
     fh = open(dir + "/" + filename)
     for line in fh.readlines():
         a = line.strip('\r\n')
-        inl = re.search("#include\s*<.*/(.+)>", a)
+        inl = re.search("#include\s*<(.+)>", a)
         if inl:
-            printfile( inl.group(1), dir)
+            a = inl.group(1)
+            fn = a.split('/')[-1]
+            a = a.replace( fn, "" )
+            printfile( fn, dir + "/" + a)
         else:
             print( '"' + a.replace("\\","\\\\").replace("\"", "\\\"") + ' \\n"\\' )