create CPP from ISPC tutorials

tutorials/curve_geometry/curve_geometry_device.ispc

@@ -138,6 +138,7 @@ export void device_init (uniform int8* uniform cfg)
 
   /* commit changes to scene */
   rtcCommitScene (data.g_scene);
+  data.g_traversable = rtcGetSceneTraversable(data.g_scene);
 }
 
 /* task that renders a single screen tile */
@@ -156,7 +157,7 @@ void renderPixelStandard(const uniform TutorialData& data,
   uniform RTCIntersectArguments iargs;
   rtcInitIntersectArguments(&iargs);
   iargs.feature_mask = (uniform RTCFeatureFlags) (FEATURE_MASK);
-  rtcIntersectV(data.g_scene,RTCRayHit_(ray),&iargs);
+  rtcTraversableIntersectV(data.g_traversable,RTCRayHit_(ray),&iargs);
   RayStats_addRay(stats);
 
   /* shade pixels */
@@ -188,7 +189,7 @@ void renderPixelStandard(const uniform TutorialData& data,
     uniform RTCOccludedArguments sargs;
     rtcInitOccludedArguments(&sargs);
     sargs.feature_mask = (uniform RTCFeatureFlags) (FEATURE_MASK);
-    rtcOccludedV(data.g_scene,RTCRay_(shadow),&sargs);
+    rtcTraversableOccludedV(data.g_traversable,RTCRay_(shadow),&sargs);
     RayStats_addShadowRay(stats);
 
     /* add light contribution */
@@ -286,6 +287,5 @@ export void device_render (uniform int* uniform pixels,
 /* called by the C++ code for cleanup */
 export void device_cleanup ()
 {
-  rtcReleaseScene (data.g_scene); data.g_scene = NULL;
   TutorialData_Destructor(&data);
 }

tutorials/curve_geometry/curve_geometry_device.isph

@@ -68,6 +68,7 @@ uniform int8 static_hair_flags_linear[NUM_CURVES] = {
 struct TutorialData
 {
   RTCScene g_scene;
+  RTCTraversable g_traversable;
   Vec4f* hair_vertices;
   Vec3fa* hair_normals;
   Vec3fa* hair_vertex_colors;
@@ -79,6 +80,7 @@ struct TutorialData
 inline void TutorialData_Constructor(uniform TutorialData* uniform This)
 {
   This->g_scene = NULL;
+  This->g_traversable = NULL;
   This->hair_vertices = uniform new uniform Vec4f[NUM_VERTICES];
   This->hair_normals = uniform new uniform Vec3fa[NUM_VERTICES];
   This->hair_vertex_colors = uniform new uniform Vec3fa[NUM_VERTICES];

tutorials/dynamic_scene/dynamic_scene_device.ispc

@@ -163,7 +163,7 @@ void renderPixelStandard(const uniform TutorialData& data,
   rtcInitIntersectArguments(&iargs);
   iargs.feature_mask = (uniform RTCFeatureFlags) (FEATURE_MASK);
   
-  rtcIntersectV(data.g_scene,RTCRayHit_(ray),&iargs);
+  rtcTraversableIntersectV(data.g_traversable,RTCRayHit_(ray),&iargs);
   RayStats_addRay(stats);
 
   /* shade pixels */
@@ -181,7 +181,7 @@ void renderPixelStandard(const uniform TutorialData& data,
     uniform RTCOccludedArguments sargs;
     rtcInitOccludedArguments(&sargs);
     sargs.feature_mask = (uniform RTCFeatureFlags) (FEATURE_MASK);
-    rtcOccludedV(data.g_scene,RTCRay_(shadow),&sargs);
+    rtcTraversableOccludedV(data.g_traversable,RTCRay_(shadow),&sargs);
     RayStats_addShadowRay(stats);
 
     /* add light contribution */
@@ -309,6 +309,7 @@ export void device_render (uniform int* uniform pixels,
 
   /* commit changes to scene */
   rtcCommitScene (data.g_scene);
+  data.g_traversable = rtcGetSceneTraversable(data.g_scene);
 }
 
 /* called by the C++ code for cleanup */

tutorials/dynamic_scene/dynamic_scene_device.isph

@@ -15,6 +15,7 @@ struct TutorialData
 
   /* scene data */
   RTCScene g_scene;
+  RTCTraversable g_traversable;
 
   uniform Vec3fa* uniform position;
   uniform Vec3fa* uniform colors;
@@ -28,6 +29,7 @@ void TutorialData_Constructor(uniform TutorialData* uniform This)
   This->numPhi = NUM_PHI;
   This->numTheta = NUM_THETA;
   This->g_scene  = NULL;
+  This->g_traversable  = NULL;
   This->position = uniform new uniform Vec3fa[NUM_SPHERES];
   This->colors = uniform new uniform Vec3fa[NUM_SPHERES+1];
   This->radius = uniform new uniform float[NUM_SPHERES];

tutorials/grid_geometry/grid_geometry_device.cpp

@@ -470,6 +470,7 @@ void createGridGeometry (GridMesh& gmesh)
   }
 
   /* create normal debug geometry */
+#if 0
   gmesh.geomNormals = rtcNewGeometry (g_device, RTC_GEOMETRY_TYPE_ROUND_BEZIER_CURVE);
   Vec3ff *nvertices = (Vec3ff *) rtcSetNewGeometryBuffer(gmesh.geomNormals,RTC_BUFFER_TYPE_VERTEX,0,RTC_FORMAT_FLOAT4,sizeof(Vec3ff),4*numVertices);
   int*    curves    = (int    *) rtcSetNewGeometryBuffer(gmesh.geomNormals,RTC_BUFFER_TYPE_INDEX,0,RTC_FORMAT_UINT   ,sizeof(int)   ,numVertices);
@@ -497,6 +498,7 @@ void createGridGeometry (GridMesh& gmesh)
     }
     h+=sphere_faces[f];
   }
+#endif
   
   /* we do not need this temporary data anymore */
   rtcReleaseGeometry(geomSubdiv);

tutorials/grid_geometry/grid_geometry_device.h

@@ -27,8 +27,8 @@ struct TutorialData
 
 void TutorialData_Constructor(TutorialData* This)
 {
-  This->g_scene  = nullptr;
+  This->g_scene = nullptr;
-  This->g_traversable  = nullptr;
+  This->g_traversable = nullptr;
   This->gmesh.geom = nullptr;
   This->gmesh.geomNormals = nullptr;
   This->gmesh.egrids = nullptr;

tutorials/grid_geometry/grid_geometry_device.ispc

@@ -545,6 +545,7 @@ export void device_init (uniform int8* uniform cfg)
 
   /* commit changes to scene */
   rtcCommitScene (data.g_scene);
+  data.g_traversable = rtcGetSceneTraversable(data.g_scene);
 }
 
 Vec3f mylerp(float f, const Vec3f& a, const Vec3f& b) { // FIXME: use lerp, need to make ISPC lerp and C++ lerp compatible first
@@ -567,7 +568,7 @@ void renderPixelStandard(const uniform TutorialData& data,
   uniform RTCIntersectArguments iargs;
   rtcInitIntersectArguments(&iargs);
   iargs.feature_mask = (uniform RTCFeatureFlags) (FEATURE_MASK);
-  rtcIntersectV(data.g_scene,RTCRayHit_(ray),&iargs);
+  rtcTraversableIntersectV(data.g_traversable,RTCRayHit_(ray),&iargs);
   RayStats_addRay(stats);
   
   /* shade pixels */
@@ -608,7 +609,7 @@ void renderPixelStandard(const uniform TutorialData& data,
     uniform RTCOccludedArguments sargs;
     rtcInitOccludedArguments(&sargs);
     sargs.feature_mask = (uniform RTCFeatureFlags) (FEATURE_MASK);
-    rtcOccludedV(data.g_scene,RTCRay_(shadow),&sargs);
+    rtcTraversableOccludedV(data.g_traversable,RTCRay_(shadow),&sargs);
     RayStats_addShadowRay(stats);
 
     /* add light contribution */

tutorials/grid_geometry/grid_geometry_device.isph

@@ -18,13 +18,15 @@ struct TutorialData
 {
   /* scene data */
   RTCScene g_scene;
+  RTCTraversable g_traversable;
 
   uniform GridMesh gmesh;
 };
 
 void TutorialData_Constructor(uniform TutorialData* uniform This)
 {
-  This->g_scene  = NULL;
+  This->g_scene = NULL;
+  This->g_traversable = NULL;
   This->gmesh.geom = NULL;
   This->gmesh.geomNormals = NULL;
   This->gmesh.egrids = NULL;

tutorials/hair_geometry/hair_geometry_device.cpp

@@ -250,7 +250,7 @@ Vec3fa occluded(RTCTraversable traversable, RayQueryContext* context, Ray& ray)
   args.filter = occlusionFilter;
   
   rtcTraversableOccluded1(traversable,RTCRay_(ray),&args);
-  context->userRayExt = NULL;
+  context->userRayExt = nullptr;
 
   if (ray.tfar < 0) return Vec3fa(0.0f);
   else              return transparency;

tutorials/hair_geometry/hair_geometry_device.ispc

@@ -232,7 +232,7 @@ RTC_SYCL_INDIRECTLY_CALLABLE unmasked void occlusionFilter(const RTCFilterFuncti
     valid_i[programIndex] = 0;
 }
 
-Vec3f occluded(RTCScene scene, uniform RayQueryContext* uniform context, Ray& ray)
+Vec3f occluded(RTCTraversable traversable, uniform RayQueryContext* uniform context, Ray& ray)
 {
   Vec3f transparency = make_Vec3f(1.0f);
   context->userRayExt = &transparency;
@@ -247,7 +247,7 @@ Vec3f occluded(RTCScene scene, uniform RayQueryContext* uniform context, Ray& ra
   args.feature_mask = (uniform RTCFeatureFlags) (FEATURE_MASK);
   args.filter = occlusionFilter;
   
-  rtcOccludedV(scene,RTCRay_(ray),&args);
+  rtcTraversableOccludedV(traversable,RTCRay_(ray),&args);
   context->userRayExt = NULL;
 
   if (ray.tfar < 0) return make_Vec3f(0.0f);
@@ -286,7 +286,7 @@ Vec3f renderPixel(const uniform TutorialData& data, float x, float y, const unif
       return color;
 
     /* intersect ray with scene and gather all hits */
-    rtcIntersectV(data.scene,RTCRayHit_(ray),&args);
+    rtcTraversableIntersectV(data.traversable,RTCRayHit_(ray),&args);
     RayStats_addRay(stats);
 
     /* exit if we hit environment */
@@ -326,7 +326,7 @@ Vec3f renderPixel(const uniform TutorialData& data, float x, float y, const unif
 
     /* sample directional light */
     Ray shadow = make_Ray(ray.org + ray.tfar*ray.dir, neg(make_Vec3f(data.dirlight_direction)), eps, inf, time);
-    Vec3f T = occluded(data.scene,&context,shadow);
+    Vec3f T = occluded(data.traversable,&context,shadow);
     RayStats_addShadowRay(stats);
     Vec3f c = AnisotropicBlinn__eval(&brdf,neg(ray.dir),neg(make_Vec3f(data.dirlight_direction)));
     color = color + weight*c*T*make_Vec3f(data.dirlight_intensity);
@@ -416,6 +416,7 @@ export void device_init (uniform int8* uniform cfg)
   
   /* create scene */
   g_scene = data.scene = convertScene(data.ispc_scene);
+  data.traversable = rtcGetSceneTraversable(data.scene);
 }
 
 export void renderFrameStandard (uniform int* uniform pixels,

tutorials/hair_geometry/hair_geometry_device.isph

@@ -12,6 +12,7 @@ extern uniform Vec3fa g_ambient_intensity;
 struct TutorialData
 {
   RTCScene scene;
+  RTCTraversable traversable;
   uniform ISPCScene* uniform ispc_scene;
   uniform Vec3f dirlight_direction;
   uniform Vec3f dirlight_intensity;
@@ -47,6 +48,7 @@ namespace embree {
 void TutorialData_Constructor(uniform TutorialData* uniform This)
 {
   This->scene = NULL;
+  This->traversable = NULL;
   This->ispc_scene = g_ispc_scene;
   This->dirlight_direction = make_Vec3f(g_dirlight_direction);
   This->dirlight_intensity = make_Vec3f(g_dirlight_intensity);

tutorials/instanced_geometry/instanced_geometry_device.ispc

@@ -174,7 +174,7 @@ Vec3f renderPixel(const uniform TutorialData& data, float x, float y, const unif
   rtcInitIntersectArguments(&iargs);
   iargs.feature_mask = (uniform RTCFeatureFlags) (FEATURE_MASK);
   
-  rtcIntersectV(data.g_scene,RTCRayHit_(ray),&iargs);
+  rtcTraversableIntersectV(data.g_traversable,RTCRayHit_(ray),&iargs);
   RayStats_addRay(stats);
 
   /* shade pixels */
@@ -186,7 +186,7 @@ Vec3f renderPixel(const uniform TutorialData& data, float x, float y, const unif
     if (ray.instID[0] != RTC_INVALID_GEOMETRY_ID)
     {
       AffineSpace3f xfm;
-      rtcGetGeometryTransformFromScene(data.g_scene,ray.instID[0],0.0f,RTC_FORMAT_FLOAT3X4_COLUMN_MAJOR,&xfm);
+      rtcGetGeometryTransformFromTraversable(data.g_traversable,ray.instID[0],0.0f,RTC_FORMAT_FLOAT3X4_COLUMN_MAJOR,&xfm);
       Ns = xfmNormal(xfm,Ns);
       //Ns = xfmVector(data.normal_xfm[ray.instID[0]],Ns);
     }
@@ -207,7 +207,7 @@ Vec3f renderPixel(const uniform TutorialData& data, float x, float y, const unif
     rtcInitOccludedArguments(&sargs);
     sargs.feature_mask = (uniform RTCFeatureFlags) (FEATURE_MASK);
     
-    rtcOccludedV(data.g_scene,RTCRay_(shadow),&sargs);
+    rtcTraversableOccludedV(data.g_traversable,RTCRay_(shadow),&sargs);
     RayStats_addShadowRay(stats);
 
     /* add light contribution */
@@ -341,6 +341,7 @@ export void device_render (uniform int* uniform pixels,
   rtcCommitGeometry(g_instance2);
   rtcCommitGeometry(g_instance3);
   rtcCommitScene (data.g_scene);
+  data.g_traversable = rtcGetSceneTraversable(data.g_scene);
 }
 
 /* called by the C++ code for cleanup */

tutorials/instanced_geometry/instanced_geometry_device.isph

@@ -11,6 +11,7 @@ struct TutorialData
   /* scene data */
   RTCScene g_scene;
   RTCScene g_scene1;
+  RTCTraversable g_traversable;
   
   uniform AffineSpace3f* uniform instance_xfm;
   uniform LinearSpace3f* uniform normal_xfm;
@@ -31,6 +32,7 @@ void TutorialData_Constructor(uniform TutorialData* uniform This)
 {
   This->g_scene  = NULL;
   This->g_scene1 = NULL;
+  This->g_traversable = NULL;
   This->instance_xfm = uniform new uniform AffineSpace3f[4];
   This->normal_xfm = uniform new uniform LinearSpace3f[4];
   This->colors = uniform new uniform Vec3f[4*4];

tutorials/intersection_filter/intersection_filter_device.cpp

@@ -380,7 +380,6 @@ extern "C" void device_init (char* cfg)
 
   /* commit changes to scene */
   rtcCommitScene (data.g_scene);
-  
   data.g_traversable = rtcGetSceneTraversable(data.g_scene);
 }
 

tutorials/intersection_filter/intersection_filter_device.ispc

@@ -74,7 +74,7 @@ void renderPixelStandard(const uniform TutorialData& data,
 #if USE_ARGUMENT_CALLBACKS
     iargs.filter = intersectionFilter;
 #endif
-    rtcIntersectV(data.g_scene,RTCRayHit_(primary),&iargs);
+    rtcTraversableIntersectV(data.g_traversable,RTCRayHit_(primary),&iargs);
     RayStats_addRay(stats);
 
     /* shade pixels */
@@ -101,7 +101,7 @@ void renderPixelStandard(const uniform TutorialData& data,
 #if USE_ARGUMENT_CALLBACKS
     sargs.filter = occlusionFilter;
 #endif
-    rtcOccludedV(data.g_scene,RTCRay_(shadow),&sargs);
+    rtcTraversableOccludedV(data.g_traversable,RTCRay_(shadow),&sargs);
     RayStats_addShadowRay(stats);
 
     /* add light contribution */
@@ -378,6 +378,7 @@ export void device_init (uniform int8* uniform cfg)
 
   /* commit changes to scene */
   rtcCommitScene (data.g_scene);
+  data.g_traversable = rtcGetSceneTraversable(data.g_scene);
 }
 
 /* task that renders a single screen tile */

tutorials/intersection_filter/intersection_filter_device.isph

@@ -6,6 +6,7 @@
 struct TutorialData
 {
   RTCScene g_scene;
+  RTCTraversable g_traversable;
   uniform Vec3f* uniform colors;
 
   RTCFilterFunctionN intersectionFilter;
@@ -15,6 +16,7 @@ struct TutorialData
 inline void TutorialData_Constructor(uniform TutorialData* uniform This)
 {
   This->g_scene = NULL;
+  This->g_traversable = NULL;
   This->colors = NULL;
   This->intersectionFilter = NULL;
   This->occlusionFilter = NULL;

tutorials/ispc2cpp.sh

@@ -140,9 +140,15 @@ sed -i.backup  's/RTCRay1/RTCRay/g' $2
 sed -i.backup  's/rtcIntersectV/rtcIntersect1/g' $2
 sed -i.backup  's/rtcOccludedV/rtcOccluded1/g' $2
 
+sed -i.backup  's/rtcTraversableIntersectV/rtcTraversableIntersect1/g' $2
+sed -i.backup  's/rtcTraversableOccludedV/rtcTraversableOccluded1/g' $2
+
 sed -i.backup  's/rtcForwardIntersectV/rtcForwardIntersect1/g' $2
 sed -i.backup  's/rtcForwardOccludedV/rtcForwardOccluded1/g' $2
 
+sed -i.backup  's/rtcTraversableForwardIntersectV/rtcTraversableForwardIntersect1/g' $2
+sed -i.backup  's/rtcTraversableForwardOccludedV/rtcTraversableForwardOccluded1/g' $2
+
 sed -i.backup  's/rtcInterpolateV/rtcInterpolate/g' $2
 
 sed -i.backup  's/Texture_FLOAT32/Texture::FLOAT32/g' $2

tutorials/motion_blur_geometry/motion_blur_geometry_device.ispc

@@ -606,6 +606,7 @@ export void device_init (uniform int8* uniform cfg)
 
   /* commit changes to scene */
   rtcCommitScene (g_scene);
+  data.g_traversable = rtcGetSceneTraversable(data.g_scene);
 }
 
 /* task that renders a single screen tile */
@@ -625,7 +626,7 @@ Vec3f renderPixel(const uniform TutorialData& data, float x, float y, const unif
   iargs.intersect = sphereIntersectFuncN;
 #endif
   
-  rtcIntersectV(data.g_scene,RTCRayHit_(ray),&iargs);
+  rtcTraversableIntersectV(data.g_traversable,RTCRayHit_(ray),&iargs);
   RayStats_addRay(stats);
 
   /* shade pixels */
@@ -662,7 +663,7 @@ Vec3f renderPixel(const uniform TutorialData& data, float x, float y, const unif
 #if USE_ARGUMENT_CALLBACKS
     sargs.occluded = sphereOccludedFuncN;
 #endif
-    rtcOccludedV(data.g_scene,RTCRay_(shadow),&sargs);
+    rtcTraversableOccludedV(data.g_traversable,RTCRay_(shadow),&sargs);
     RayStats_addShadowRay(stats);
 
     /* add light contribution */

tutorials/motion_blur_geometry/motion_blur_geometry_device.isph

@@ -21,6 +21,7 @@ struct TutorialData
   uniform int frameID;
   
   RTCScene g_scene;
+  RTCTraversable g_traversable;
   uniform Vec3f* uniform face_colors;
   uniform float g_time;
   
@@ -59,6 +60,7 @@ inline void TutorialData_Constructor(uniform TutorialData* uniform This)
 {
   This->frameID = 50;
   This->g_scene = NULL;
+  This->g_traversable = NULL;
   This->face_colors = uniform new uniform Vec3f[12];
   This->g_time = 0;
   This->g_accu = NULL;

tutorials/multi_instanced_geometry/multi_instanced_geometry_device.cpp

@@ -9,7 +9,7 @@ namespace embree {
 #define FEATURE_MASK \
   RTC_FEATURE_FLAG_TRIANGLE | \
   RTC_FEATURE_FLAG_INSTANCE
-  
+
 RTCScene g_scene = nullptr;
 TutorialData g_data;
 extern "C" bool g_changed;

tutorials/multi_instanced_geometry/multi_instanced_geometry_device.h

@@ -51,8 +51,8 @@ extern "C" void cleanupScene(TutorialData& data);
 
 inline void TutorialData_Constructor(TutorialData* This)
 {
-  This->g_scene  = nullptr;
+  This->g_scene = nullptr;
-  This->g_traversable  = nullptr;
+  This->g_traversable = nullptr;
   This->g_accu = nullptr;
   This->g_accu_width = 0;
   This->g_accu_height = 0;
@@ -71,7 +71,7 @@ inline void TutorialData_Destructor(TutorialData* This)
 {
   rtcReleaseScene (This->g_scene); This->g_scene = nullptr;
   cleanupScene(*This);
-  alignedUSMFree(This->g_accu); This->g_accu = nullptr;
+  delete[] This->g_accu; This->g_accu = nullptr;
   This->g_accu_width = 0;
   This->g_accu_height = 0;
   This->g_accu_count = 0;

tutorials/multi_instanced_geometry/multi_instanced_geometry_device.ispc

@@ -240,7 +240,7 @@ void renderPixelStandard(const uniform TutorialData& data, int x, int y,
     Vec3f emission;
     sampleLightDirection(RandomSampler_get3D(sampler), lightDir, emission);
     Ray shadowRay = makeShadowRay(primaryRay, lightDir, stats);
-    rtcOccludedV(data.g_scene, RTCRay_(shadowRay), &sargs);
+    rtcTraversableOccludedV(data.g_traversable, RTCRay_(shadowRay), &sargs);
     color = shade(data, primaryRay, shadowRay, lightDir, emission);
   }
   

tutorials/multi_instanced_geometry/multi_instanced_geometry_device.isph

@@ -50,6 +50,7 @@ extern "C" void cleanupScene(uniform TutorialData& data);
 inline void TutorialData_Constructor(uniform TutorialData* uniform This)
 {
   This->g_scene = NULL;
+  This->g_traversable = NULL;
   This->g_accu = NULL;
   This->g_accu_width = 0;
   This->g_accu_height = 0;

tutorials/multi_instanced_geometry/scene.cpp

@@ -194,12 +194,12 @@ extern "C" RTCScene initializeScene(TutorialData& data, RTCDevice device)
   cleanupScene(data);
 
   data.g_instanceLevels.numLevels = 2;
-  data.g_instanceLevels.numInstancesOnLevel = (unsigned int*)alignedMalloc(2*sizeof(unsigned int), 16);
+  data.g_instanceLevels.numInstancesOnLevel = (unsigned int*)alignedUSMMalloc(2*sizeof(unsigned int), 16);
   data.g_instanceLevels.numInstancesOnLevel[0] = Trees::instances.numInstances;
   data.g_instanceLevels.numInstancesOnLevel[1] = Twigs01::instances.numInstances;
-  data.g_normalTransforms = (LinearSpace3fa**)alignedMalloc(2*sizeof(LinearSpace3fa*), 16);
+  data.g_normalTransforms = (LinearSpace3fa**)alignedUSMMalloc(2*sizeof(LinearSpace3fa*), 16);
-  data.g_normalTransforms[0] = (LinearSpace3fa*)alignedMalloc(data.g_instanceLevels.numInstancesOnLevel[0]*sizeof(LinearSpace3fa), 16);
+  data.g_normalTransforms[0] = (LinearSpace3fa*)alignedUSMMalloc(data.g_instanceLevels.numInstancesOnLevel[0]*sizeof(LinearSpace3fa), 16);
-  data.g_normalTransforms[1] = (LinearSpace3fa*)alignedMalloc(data.g_instanceLevels.numInstancesOnLevel[1]*sizeof(LinearSpace3fa), 16);
+  data.g_normalTransforms[1] = (LinearSpace3fa*)alignedUSMMalloc(data.g_instanceLevels.numInstancesOnLevel[1]*sizeof(LinearSpace3fa), 16);
   data.g_instanceLevels.normalTransforms = data.g_normalTransforms;
 
   RandomSampler sampler;

tutorials/multi_instanced_geometry/scene.h

@@ -7,14 +7,18 @@
 #include "../common/math/linearspace.h"
 #include "../../include/embree4/rtcore.h"
 
+namespace embree {
+
 extern "C" struct InstanceLevels
 {
   unsigned int numLevels;
   const unsigned int* numInstancesOnLevel;
-  embree::LinearSpace3fa** normalTransforms;
+  LinearSpace3fa* * normalTransforms;
 };
 
 extern "C" RTCScene initializeScene(RTCDevice device,
-                                    struct InstanceLevels* levels);
+                                            struct InstanceLevels* levels);
 
-extern "C" void cleanupScene();
+extern "C" void cleanupScene();
+
+} // namespace embree

tutorials/pathtracer/pathtracer_device.ispc

@@ -1504,7 +1504,7 @@ Vec3f renderPixelFunction(const uniform TutorialData& data, float x, float y, Ra
     args.filter = NULL;
 #endif
   
-    rtcIntersectV(data.scene,RTCRayHit_(ray),&args);
+    rtcTraversableIntersectV(data.traversable,RTCRayHit_(ray),&args);
     RayStats_addRay(stats);
     const Vec3f wo = neg(ray.dir);
 
@@ -1577,7 +1577,7 @@ Vec3f renderPixelFunction(const uniform TutorialData& data, float x, float y, Ra
 #if USE_ARGUMENT_CALLBACKS && ENABLE_FILTER_FUNCTION
       sargs.filter = contextFilterFunction;
 #endif
-      rtcOccludedV(data.scene,RTCRay_(shadow),&sargs);
+      rtcTraversableOccludedV(data.traversable,RTCRay_(shadow),&sargs);
       RayStats_addShadowRay(stats);
 #if !ENABLE_FILTER_FUNCTION
       if (shadow.tfar > 0.0f)
@@ -1792,6 +1792,7 @@ export void device_render (uniform int* uniform pixels,
     data.scene = convertScene(data.ispc_scene);
     if (g_subdiv_mode) updateEdgeLevels(data.ispc_scene,camera.xfm.p);
     rtcCommitScene (data.scene);
+    data.traversable = rtcGetSceneTraversable(data.scene);
   }
 
   /* create accumulator */
@@ -1820,6 +1821,7 @@ export void device_render (uniform int* uniform pixels,
     if (g_subdiv_mode) {
       updateEdgeLevels(data.ispc_scene,camera.xfm.p);
       rtcCommitScene (data.scene);
+      data.traversable = rtcGetSceneTraversable(data.scene);
     }
   }
   else

tutorials/pathtracer/pathtracer_device.isph

@@ -39,6 +39,7 @@ struct TutorialData
   
   /* scene data */
   RTCScene scene;
+  RTCTraversable traversable;
 };
 
 void TutorialData_Constructor(uniform TutorialData* uniform This)
@@ -63,6 +64,7 @@ void TutorialData_Constructor(uniform TutorialData* uniform This)
   This->use_smooth_normals = false;
         
   This->scene = NULL;
+  This->traversable = NULL;
 }
 
 void TutorialData_Destructor(uniform TutorialData* uniform This)

tutorials/point_geometry/point_geometry_device.ispc

@@ -97,6 +97,7 @@ export void device_init (uniform int8* uniform cfg)
 
   /* commit changes to scene */
   rtcCommitScene (g_scene);
+  data.g_traversable = rtcGetSceneTraversable(data.g_scene);
 }
 
 /* task that renders a single screen tile */
@@ -116,7 +117,7 @@ void renderPixelStandard(const uniform TutorialData& data,
   rtcInitIntersectArguments(&iargs);
   iargs.feature_mask = (RTCFeatureFlags)(FEATURE_MASK);
   
-  rtcIntersectV(data.g_scene,RTCRayHit_(ray),&iargs);
+  rtcTraversableIntersectV(data.g_traversable,RTCRayHit_(ray),&iargs);
   RayStats_addRay(stats);
 
   /* shade pixels */
@@ -139,7 +140,7 @@ void renderPixelStandard(const uniform TutorialData& data,
     rtcInitOccludedArguments(&sargs);
     sargs.feature_mask = (RTCFeatureFlags)(FEATURE_MASK);
   
-    rtcOccludedV(data.g_scene,RTCRay_(shadow),&sargs);
+    rtcTraversableOccludedV(data.g_traversable,RTCRay_(shadow),&sargs);
     RayStats_addShadowRay(stats);
 
     /* add light contribution */

tutorials/point_geometry/point_geometry_device.isph

@@ -10,12 +10,14 @@
 struct TutorialData
 {
   RTCScene g_scene;
+  RTCTraversable g_traversable;
   uniform Vec3f* uniform point_colors;
 };
 
 inline void TutorialData_Constructor(uniform TutorialData* uniform This)
 {
   This->g_scene = NULL;
+  This->g_traversable = NULL;
   This->point_colors = uniform new uniform Vec3f[NUM_POINTS];
 }
 

tutorials/quaternion_motion_blur/quaternion_motion_blur_device.h

@@ -49,8 +49,8 @@ inline void TutorialData_Constructor(TutorialData* This)
 {
   This->g_scene = nullptr;
   This->g_scene0 = nullptr;
-  This->g_traversable = nullptr;
   This->g_spheres = nullptr;
+  This->g_traversable = nullptr;
   This->g_spp = 0;
   This->g_motion_blur = true;
   This->g_time = 0.f;

tutorials/quaternion_motion_blur/quaternion_motion_blur_device.ispc

@@ -250,6 +250,7 @@ export void device_init (uniform int8* uniform cfg)
   rtcCommitGeometry(g_instance_quaternion_1);
 
   rtcCommitScene (data.g_scene);
+  data.g_traversable = rtcGetSceneTraversable(data.g_scene);
 }
 
 inline Vec3f face_forward(const Vec3f& dir, const Vec3f& _Ng) {
@@ -279,7 +280,7 @@ Vec3f renderPixelFunction(const uniform TutorialData& data,
                      RTC_INVALID_GEOMETRY_ID, RTC_INVALID_GEOMETRY_ID);
 
   /* intersect ray with scene */
-  rtcIntersectV(data.g_scene,RTCRayHit_(ray),&args);
+  rtcTraversableIntersectV(data.g_traversable,RTCRayHit_(ray),&args);
   RayStats_addRay(stats);
 
   /* shade pixels */
@@ -442,6 +443,7 @@ export void device_render (uniform int* uniform pixels,
   {
     updateTransformation();
     rtcCommitScene(data.g_scene);
+    data.g_traversable = rtcGetSceneTraversable(data.g_scene);
   }
 
 

tutorials/quaternion_motion_blur/quaternion_motion_blur_device.isph

@@ -16,6 +16,7 @@ struct TutorialData
 {  
   RTCScene g_scene;
   RTCScene g_scene0;
+  RTCTraversable g_traversable;
   uniform Sphere* uniform g_spheres;
   int g_spp;
   bool g_motion_blur;
@@ -47,6 +48,7 @@ inline void TutorialData_Constructor(uniform TutorialData* uniform This)
   This->g_scene = NULL;
   This->g_scene0 = NULL;
   This->g_spheres = NULL;
+  This->g_traversable = NULL;
   This->g_spp = 0;
   This->g_motion_blur = true;
   This->g_time = 0.f;

tutorials/ray_mask/ray_mask_device.ispc

@@ -117,6 +117,7 @@ export void device_init (uniform int8* uniform cfg)
 
   /* commit changes to scene */
   rtcCommitScene (data.g_scene);
+  data.g_traversable = rtcGetSceneTraversable(data.g_scene);
 }
 
 /* task that renders a single screen tile */
@@ -137,7 +138,7 @@ void renderPixelStandard(const uniform TutorialData& data,
   uniform RTCIntersectArguments iargs;
   rtcInitIntersectArguments(&iargs);
   iargs.feature_mask = (uniform RTCFeatureFlags) (FEATURE_MASK);
-  rtcIntersectV(data.g_scene,RTCRayHit_(ray),&iargs);
+  rtcTraversableIntersectV(data.g_traversable,RTCRayHit_(ray),&iargs);
   RayStats_addRay(stats);
 
   /* shade pixels */
@@ -157,7 +158,7 @@ void renderPixelStandard(const uniform TutorialData& data,
     uniform RTCOccludedArguments sargs;
     rtcInitOccludedArguments(&sargs);
     sargs.feature_mask = (uniform RTCFeatureFlags) (FEATURE_MASK);
-    rtcOccludedV(data.g_scene,RTCRay_(shadow),&sargs);
+    rtcTraversableOccludedV(data.g_traversable,RTCRay_(shadow),&sargs);
     RayStats_addShadowRay(stats);
 
     /* add light contribution */

tutorials/ray_mask/ray_mask_device.isph

@@ -6,6 +6,7 @@
 struct TutorialData
 {
   RTCScene g_scene;
+  RTCTraversable g_traversable;
   uniform Vec3f* uniform face_colors;
   uniform Vec3fa* uniform vertex_colors;
   uniform bool enable_ray_mask;
@@ -14,6 +15,7 @@ struct TutorialData
 inline void TutorialData_Constructor(uniform TutorialData* uniform This)
 {
   This->g_scene = NULL;
+  This->g_traversable = NULL;
   This->face_colors = NULL;
   This->vertex_colors = NULL;
 }

tutorials/triangle_geometry/triangle_geometry_device.cpp

@@ -100,7 +100,6 @@ extern "C" void device_init (char* cfg)
   /* create scene */
   TutorialData_Constructor(&data);
   g_scene = data.g_scene = rtcNewScene(g_device);
-  rtcSetSceneFlags(data.g_scene, RTC_SCENE_FLAG_PREFETCH_USM_SHARED_ON_GPU);
 
   /* add cube */
   addCube(data.g_scene);
@@ -193,7 +192,7 @@ extern "C" void renderFrameStandard (int* pixels,
                           const float time,
                           const ISPCCamera& camera)
 {
-#if defined(EMBREE_SYCL_TUTORIAL) && !defined(EMBREE_SYCL_RT_SIMULATION) && !defined(EMBREE_SYCL_RT_SIMULATION)
+#if defined(EMBREE_SYCL_TUTORIAL) && !defined(EMBREE_SYCL_RT_SIMULATION)
   TutorialData ldata = data;
   sycl::event event = global_gpu_queue->submit([=](sycl::handler& cgh){
     const sycl::nd_range<2> nd_range = make_nd_range(height,width);

tutorials/triangle_geometry/triangle_geometry_device.ispc

@@ -107,6 +107,9 @@ export void device_init (uniform int8* uniform cfg)
 
   /* commit changes to scene */
   rtcCommitScene (data.g_scene);
+
+  /* get traversable handle for the scene */
+  data.g_traversable = rtcGetSceneTraversable(data.g_scene);
 }
 
 /* task that renders a single screen tile */
@@ -125,7 +128,7 @@ void renderPixelStandard(const uniform TutorialData& data,
   uniform RTCIntersectArguments iargs;
   rtcInitIntersectArguments(&iargs);
   iargs.feature_mask = (uniform RTCFeatureFlags) (FEATURE_MASK);
-  rtcIntersectV(data.g_scene,RTCRayHit_(ray),&iargs);
+  rtcTraversableIntersectV(data.g_traversable,RTCRayHit_(ray),&iargs);
   RayStats_addRay(stats);
 
   /* shade pixels */
@@ -143,7 +146,7 @@ void renderPixelStandard(const uniform TutorialData& data,
     uniform RTCOccludedArguments sargs;
     rtcInitOccludedArguments(&sargs);
     sargs.feature_mask = (uniform RTCFeatureFlags) (FEATURE_MASK);
-    rtcOccludedV(data.g_scene,RTCRay_(shadow),&sargs);
+    rtcTraversableOccludedV(data.g_traversable,RTCRay_(shadow),&sargs);
     RayStats_addShadowRay(stats);
 
     /* add light contribution */

tutorials/triangle_geometry/triangle_geometry_device.isph

@@ -6,6 +6,7 @@
 struct TutorialData
 {
   RTCScene g_scene;
+  RTCTraversable g_traversable;
   uniform Vec3f* uniform face_colors;
   uniform Vec3fa* uniform vertex_colors;
 };
@@ -13,6 +14,7 @@ struct TutorialData
 inline void TutorialData_Constructor(uniform TutorialData* uniform This)
 {
   This->g_scene = NULL;
+  This->g_traversable = NULL;
   This->face_colors = NULL;
   This->vertex_colors = NULL;
 }

tutorials/user_geometry/user_geometry_device.cpp

@@ -502,9 +502,7 @@ RTC_SYCL_INDIRECTLY_CALLABLE void sphereFilterFunction(const RTCFilterFunctionNA
   const RayQueryContext* context = (const RayQueryContext*) args->context;
   struct Ray* ray    = (struct Ray*)args->ray;
   //struct RTCHit* hit = (struct RTCHit*)args->hit;
-  const unsigned int N = args->N;
+  assert(args->N == 1);
-  assert(N == 1);
-  _unused(N);
 
   /* avoid crashing when debug visualizations are used */
   if (context == nullptr)
@@ -778,7 +776,7 @@ Vec3fa renderPixelStandard(const TutorialData& data,
       Ns = xfmVector(data.g_instance[ray.instID[0]]->normal2world,Vec3fa(Ns));
     }
     Ns = face_forward(ray.dir,normalize(Ns));
-    
+
     /* calculate diffuse color of geometries */
     Vec3fa diffuse = Vec3fa(0.0f);
     if      (ray.instID[0] ==  0) diffuse = data.colors[4*ray.instID[0]+ray.primID];

tutorials/user_geometry/user_geometry_device.h

@@ -41,13 +41,14 @@ struct TutorialData
   RTCScene g_scene0;
   RTCScene g_scene1;
   RTCScene g_scene2;
+
+  RTCTraversable g_traversable;
+
   Sphere* g_spheres;
   Sphere* g_sphere0;
   Sphere* g_sphere1;
-  
+
   Instance* g_instance[4];
-  
-  RTCTraversable g_traversable;
 
   Vec3fa* colors;
 };
@@ -67,6 +68,7 @@ inline void TutorialData_Constructor(TutorialData* This)
   This->g_scene0 = nullptr;
   This->g_scene1 = nullptr;
   This->g_scene2 = nullptr;
+  This->g_traversable = nullptr;
   This->g_spheres = nullptr;
   This->g_sphere0 = nullptr;
   This->g_sphere1 = nullptr;
@@ -74,7 +76,6 @@ inline void TutorialData_Constructor(TutorialData* This)
   This->g_instance[1] = nullptr;
   This->g_instance[2] = nullptr;
   This->g_instance[3] = nullptr;
-  This->g_traversable  = nullptr;
   This->colors = (Vec3fa*) alignedUSMMalloc((5*4)*sizeof(Vec3fa),16);
 }
 

tutorials/user_geometry/user_geometry_device.ispc

@@ -121,7 +121,7 @@ RTC_SYCL_INDIRECTLY_CALLABLE unmasked void instanceIntersectFunc(const RTCInters
   uniform RTCIntersectArguments args;
   rtcInitIntersectArguments(&args);
   args.context = context;
-  rtcIntersectV(instance->object,RTCRayHit_(*ray),&args);
+  rtcTraversableIntersectV(instance->object,RTCRayHit_(*ray),&args);
   popInstanceId(context);
   const float updated_tfar = ray->tfar;
   ray->org = ray_org;
@@ -142,7 +142,7 @@ RTC_SYCL_INDIRECTLY_CALLABLE unmasked void instanceIntersectFunc(const RTCInters
   xray.id = 0;
   xray.flags = 0;
   
-  rtcForwardIntersectV(args,instance->object,&xray,args->geomID);
+  rtcTraversableForwardIntersectV(args,instance->object,&xray,args->geomID);
   
 #endif
 }
@@ -174,7 +174,7 @@ RTC_SYCL_INDIRECTLY_CALLABLE unmasked void instanceOccludedFunc(const RTCOcclude
   uniform RTCOccludedArguments args;
   rtcInitOccludedArguments(&args);
   args.context = context;
-  rtcOccludedV(instance->object,RTCRay_(*ray),&args);
+  rtcTraversableOccludedV(instance->object,RTCRay_(*ray),&args);
   popInstanceId(context);
   const float updated_tfar = ray->tfar;
   ray->org    = ray_org;
@@ -196,7 +196,7 @@ RTC_SYCL_INDIRECTLY_CALLABLE unmasked void instanceOccludedFunc(const RTCOcclude
   xray.id = 0;
   xray.flags = 0;
   
-  rtcForwardOccludedV(args,instance->object,&xray,args->geomID);
+  rtcTraversableForwardOccludedV(args,instance->object,&xray,args->geomID);
   
 #endif
 }
@@ -206,7 +206,7 @@ uniform Instance* uniform createInstance (RTCScene scene, RTCScene object, unifo
 #if !ENABLE_NATIVE_INSTANCING
   uniform Instance* uniform instance = uniform new uniform Instance;
   instance->type = USER_GEOMETRY_INSTANCE;
-  instance->object = object;
+  instance->object = rtcGetSceneTraversable(object);
   instance->lower = lower;
   instance->upper = upper;
   instance->local2world.l.vx = make_Vec3f(1,0,0);
@@ -500,8 +500,7 @@ RTC_SYCL_INDIRECTLY_CALLABLE unmasked void sphereFilterFunction(const RTCFilterF
   const uniform RayQueryContext* uniform context = (const uniform RayQueryContext* uniform) args->context;
   varying struct Ray* uniform ray    = (varying struct Ray* uniform)args->ray;
   //varying struct RTCHit* uniform hit = (varying struct RTCHit* uniform)args->hit;
-  const uniform unsigned int N = args->N;
+  assert(args->N == programCount);
-  assert(N == programCount);
 
   /* avoid crashing when debug visualizations are used */
   if (context == NULL)
@@ -761,7 +760,7 @@ Vec3f renderPixelStandard(const uniform TutorialData& data,
 #endif
   iargs.feature_mask = (uniform RTCFeatureFlags) (FEATURE_MASK);
   
-  rtcIntersectV(data.g_scene,RTCRayHit_(ray),&iargs);
+  rtcTraversableIntersectV(data.g_traversable,RTCRayHit_(ray),&iargs);
   RayStats_addRay(stats);
 
   /* shade pixels */
@@ -798,7 +797,7 @@ Vec3f renderPixelStandard(const uniform TutorialData& data,
 #endif
     sargs.feature_mask = (uniform RTCFeatureFlags) (FEATURE_MASK);
     
-    rtcOccludedV(data.g_scene,RTCRay_(shadow),&sargs);
+    rtcTraversableOccludedV(data.g_traversable,RTCRay_(shadow),&sargs);
     RayStats_addShadowRay(stats);
 
     /* add light contribution */
@@ -922,6 +921,7 @@ export void device_render (uniform int* uniform pixels,
   updateInstance(data.g_scene,data.g_instance[2]);
   updateInstance(data.g_scene,data.g_instance[3]);
   rtcCommitScene (data.g_scene);
+  data.g_traversable = rtcGetSceneTraversable(data.g_scene);
 }
 
 /* called by the C++ code for cleanup */

tutorials/user_geometry/user_geometry_device.isph

@@ -24,7 +24,7 @@ struct Instance
   ALIGNED_STRUCT_(16)
   UserGeometryType type;
   RTCGeometry geometry;
-  RTCScene object;
+  RTCTraversable object;
   AffineSpace3f local2world;
   AffineSpace3f world2local;
   LinearSpace3f normal2world;
@@ -39,12 +39,15 @@ struct TutorialData
   RTCScene g_scene0;
   RTCScene g_scene1;
   RTCScene g_scene2;
+
+  RTCTraversable g_traversable;
+
   uniform Sphere* uniform g_spheres;
   uniform Sphere* uniform g_sphere0;
   uniform Sphere* uniform g_sphere1;
-  
+
   uniform Instance* uniform g_instance[4];
-  
+
   uniform Vec3f* uniform colors;
 };
 
@@ -63,6 +66,7 @@ inline void TutorialData_Constructor(uniform TutorialData* uniform This)
   This->g_scene0 = NULL;
   This->g_scene1 = NULL;
   This->g_scene2 = NULL;
+  This->g_traversable = NULL;
   This->g_spheres = NULL;
   This->g_sphere0 = NULL;
   This->g_sphere1 = NULL;

tutorials/viewer/viewer_device.cpp

@@ -142,6 +142,8 @@ AffineSpace3fa calculate_interpolated_space (ISPCInstance* instance, float gtime
   return (1.0f-ftime)*AffineSpace3fa(instance->spaces[itime+0]) + ftime*AffineSpace3fa(instance->spaces[itime+1]);
 }
 
+typedef ISPCInstance* ISPCInstancePtr;
+
 AffineSpace3fa calculate_interpolated_space (ISPCInstanceArray* instanceArray, unsigned int primID, float gtime)
 {
   if (instanceArray->numTimeSteps == 1)
@@ -155,7 +157,7 @@ AffineSpace3fa calculate_interpolated_space (ISPCInstanceArray* instanceArray, u
   return (1.0f-ftime)*AffineSpace3fa(instanceArray->spaces_array[itime+0][primID]) + ftime*AffineSpace3fa(instanceArray->spaces_array[itime+1][primID]);
 }
 
-typedef ISPCInstance* ISPCInstancePtr;
+typedef ISPCInstanceArray* ISPCInstanceArrayPtr;
 
 unsigned int postIntersect(const TutorialData& data, const Ray& ray, DifferentialGeometry& dg)
 {
@@ -170,18 +172,18 @@ unsigned int postIntersect(const TutorialData& data, const Ray& ray, Differentia
     if (geometries[instID]->type == INSTANCE) {
       ISPCInstance* instance = (ISPCInstancePtr) geometries[instID];
       local2world = local2world * calculate_interpolated_space(instance,ray.time());
+
       assert(instance->child->type == GROUP);
       geometries = ((ISPCGroup*)instance->child)->geometries;
     }
 #if defined(RTC_GEOMETRY_INSTANCE_ARRAY)
     else if (geometries[instID]->type == INSTANCE_ARRAY) {
-      ISPCInstanceArray* instanceArray = (ISPCInstanceArray*) geometries[instID];
+      ISPCInstanceArray* instanceArray = (ISPCInstanceArrayPtr) geometries[instID];
-      local2world = local2world * calculate_interpolated_space(instanceArray, ray.instPrimID[i],ray.time());
+      local2world = local2world * calculate_interpolated_space(instanceArray,ray.instPrimID[i],ray.time());
       assert(instanceArray->child->type == GROUP);
       geometries = ((ISPCGroup*)instanceArray->child)->geometries;
     }
 #endif
-
   }
 
   ISPCGeometry* mesh = geometries[ray.geomID];
@@ -229,7 +231,6 @@ void renderPixelStandard(const TutorialData& data,
   rtcTraversableIntersect1(data.traversable,RTCRayHit_(ray),&args);
   RayStats_addRay(stats);
 
-
   /* shade background black */
   if (ray.geomID == RTC_INVALID_GEOMETRY_ID) {
     pixels[y*width+x] = 0;

tutorials/viewer/viewer_device.h

@@ -43,8 +43,7 @@ void TutorialData_Constructor(TutorialData* This)
 
 void TutorialData_Destructor(TutorialData* This)
 {
-  // rtcReleaseScene is calld in ISPCScene destructor
+  rtcReleaseScene (This->scene); This->scene = nullptr;
-  // rtcReleaseScene (This->scene); This->scene = nullptr;
 }
 
 } // namespace embree

tutorials/viewer/viewer_device.ispc

@@ -134,6 +134,21 @@ AffineSpace3f calculate_interpolated_space (uniform ISPCInstance* instance, floa
 
 typedef ISPCInstance* ISPCInstancePtr;
 
+AffineSpace3f calculate_interpolated_space (uniform ISPCInstanceArray* instanceArray, unsigned int primID, float gtime)
+{
+  if (instanceArray->numTimeSteps == 1)
+    return make_AffineSpace3f(instanceArray->spaces_array[0][primID]);
+
+   /* calculate time segment itime and fractional time ftime */
+  const int time_segments = instanceArray->numTimeSteps-1;
+  const float time = gtime*(float)(time_segments);
+  const int itime = clamp((int)(floor(time)),(varying int)0,time_segments-1);
+  const float ftime = time - (float)(itime);
+  return (1.0f-ftime)*make_AffineSpace3f(instanceArray->spaces_array[itime+0][primID]) + ftime*make_AffineSpace3f(instanceArray->spaces_array[itime+1][primID]);
+}
+
+typedef ISPCInstanceArray* ISPCInstanceArrayPtr;
+
 unsigned int postIntersect(const uniform TutorialData& data, const Ray& ray, DifferentialGeometry& dg)
 {
   AffineSpace3f local2world = make_AffineSpace3f_scale(make_Vec3f(1));
@@ -144,11 +159,21 @@ unsigned int postIntersect(const uniform TutorialData& data, const Ray& ray, Dif
     const unsigned int instID = ray.instID[i];
     if (instID == -1) break;
 
-    ISPCInstance* instance = (ISPCInstancePtr) geometries[instID];
+    if (geometries[instID]->type == INSTANCE) {
-    local2world = local2world * calculate_interpolated_space(instance,ray.time);
+      ISPCInstance* instance = (ISPCInstancePtr) geometries[instID];
+      local2world = local2world * calculate_interpolated_space(instance,ray.time);
 
-    assert(instance->child->type == GROUP);
+      assert(instance->child->type == GROUP);
-    geometries = ((ISPCGroup*)instance->child)->geometries;
+      geometries = ((ISPCGroup*)instance->child)->geometries;
+    }
+#if defined(RTC_GEOMETRY_INSTANCE_ARRAY)
+    else if (geometries[instID]->type == INSTANCE_ARRAY) {
+      ISPCInstanceArray* instanceArray = (ISPCInstanceArrayPtr) geometries[instID];
+      local2world = local2world * calculate_interpolated_space(instanceArray,ray.instPrimID[i],ray.time);
+      assert(instanceArray->child->type == GROUP);
+      geometries = ((ISPCGroup*)instanceArray->child)->geometries;
+    }
+#endif
   }
 
   ISPCGeometry* mesh = geometries[ray.geomID];
@@ -193,7 +218,7 @@ void renderPixelStandard(const uniform TutorialData& data,
 #endif
   args.feature_mask = feature_mask;
   
-  rtcIntersectV(data.scene,RTCRayHit_(ray),&args);
+  rtcTraversableIntersectV(data.traversable,RTCRayHit_(ray),&args);
   RayStats_addRay(stats);
 
   /* shade background black */
@@ -341,6 +366,7 @@ export void device_render (uniform int* uniform pixels,
     g_scene = data.scene = convertScene(g_ispc_scene);
     if (data.subdiv_mode) updateEdgeLevels(g_ispc_scene, camera.xfm.p);
     rtcCommitScene (data.scene);
+    data.traversable = rtcGetSceneTraversable(data.scene);
     old_p = camera.xfm.p;
   }
 
@@ -356,6 +382,7 @@ export void device_render (uniform int* uniform pixels,
     if (camera_changed && data.subdiv_mode) {
       updateEdgeLevels(g_ispc_scene,camera.xfm.p);
       rtcCommitScene (data.scene);
+      data.traversable = rtcGetSceneTraversable(data.scene);
     }
 
     if (g_animation_mode)

tutorials/viewer/viewer_device.isph

@@ -20,6 +20,7 @@ struct TutorialData
   
   /* scene data */
   RTCScene scene;
+  RTCTraversable traversable;
   uniform bool subdiv_mode;
   uniform bool motion_blur;
 
@@ -32,6 +33,7 @@ void TutorialData_Constructor(uniform TutorialData* uniform This)
   This->instancing_mode = g_instancing_mode;
   This->iflags_coherent = g_iflags_coherent;
   This->scene = NULL;
+  This->traversable = NULL;
   This->subdiv_mode = false;
   This->motion_blur = g_motion_blur;
   This->min_width = g_min_width;

tutorials/viewer/viewer_device_debug.cpp

@@ -168,12 +168,12 @@ void DebugShaderData_Constructor(DebugShaderData* This)
   }
 
 #define RENDER_FRAME_FUNCTION_SYCL(Name)                                \
-extern "C" void renderFrame##Name (int* pixels,       \
+  extern "C" void renderFrame##Name (int* pixels,       \
-                                   const unsigned int width,    \
+                          const unsigned int width,             \
-                                   const unsigned int height,   \
+                          const unsigned int height,            \
-                                   const float time,            \
+                          const float time,                     \
-                                   const ISPCCamera& camera)          \
+                          const ISPCCamera& camera)             \
-{                                                                       \
+  {                                                                     \
     DebugShaderData data;                                               \
     DebugShaderData_Constructor(&data);                                 \
     sycl::event event;                                                  \
@@ -191,8 +191,8 @@ extern "C" void renderFrame##Name (int* pixels,       \
       unsigned int g = (unsigned int) (255.0f * clamp(color.y,0.0f,1.0f)); \
       unsigned int b = (unsigned int) (255.0f * clamp(color.z,0.0f,1.0f)); \
       pixels[y*width+x] = (b << 16) + (g << 8) + r;                   \
-   });                                     \
+    });                                     \
- });                                \
+  });                                \
   global_gpu_queue->wait_and_throw();\
   \
   const auto t0 = event.template get_profiling_info<sycl::info::event_profiling::command_start>();\
@@ -240,7 +240,7 @@ Vec3fa renderPixelDebugShader(const DebugShaderData& data, float x, float y, con
     args.feature_mask = feature_mask;
     rtcTraversableIntersect1(data.traversable,RTCRayHit_(ray),&args);
   }
-  
+
   int64_t c1 = get_tsc();
   RayStats_addRay(stats);
 
@@ -267,7 +267,7 @@ Vec3fa renderPixelDebugShader(const DebugShaderData& data, float x, float y, con
 
   case SHADER_TEXCOORDS:
   case SHADER_TEXCOORDS_GRID:
-    
+
 #if !defined(__SYCL_DEVICE_ONLY__)
     if (ray.geomID == RTC_INVALID_GEOMETRY_ID)
       return Vec3fa(0.0f,0.0f,1.0f);

tutorials/viewer/viewer_device_debug.ispc

@@ -25,6 +25,7 @@ extern uniform RTCFeatureFlags g_feature_mask;
 struct DebugShaderData
 {
   RTCScene scene;
+  RTCTraversable traversable;
   ISPCScene* uniform ispc_scene;
 
   /* intensity scaling for traversal cost visualization */
@@ -38,6 +39,7 @@ struct DebugShaderData
 void DebugShaderData_Constructor(uniform DebugShaderData* uniform This)
 {
   This->scene = g_scene;
+  This->traversable = rtcGetSceneTraversable(g_scene);
   This->ispc_scene = g_ispc_scene;
   This->scale = scale;
   This->debug = g_debug;
@@ -227,14 +229,14 @@ Vec3f renderPixelDebugShader(const uniform DebugShaderData& data, float x, float
     uniform RTCOccludedArguments args;
     rtcInitOccludedArguments(&args);
     args.feature_mask = feature_mask;
-    rtcOccludedV(data.scene,RTCRay_(ray),&args);
+    rtcTraversableOccludedV(data.traversable,RTCRay_(ray),&args);
   }
   else
   {
     uniform RTCIntersectArguments args;
     rtcInitIntersectArguments(&args);
     args.feature_mask = feature_mask;
-    rtcIntersectV(data.scene,RTCRayHit_(ray),&args);
+    rtcTraversableIntersectV(data.traversable,RTCRayHit_(ray),&args);
   }
 
   uniform int64 c1 = get_tsc();
@@ -303,6 +305,9 @@ Vec3f renderPixelDebugShader(const uniform DebugShaderData& data, float x, float
   case SHADER_CYCLES:
     return make_Vec3f((uniform float)(c1-c0)*data.scale,0.0f,0.0f);
     
+  case SHADER_AO:
+    return make_Vec3f(0,0,0);
+
   case SHADER_DEFAULT:
     return make_Vec3f(0,0,0);
   }
@@ -328,7 +333,7 @@ Vec3f renderPixelAOShader(const uniform DebugShaderData& data, float x, float y,
   uniform RTCIntersectArguments args;
   rtcInitIntersectArguments(&args);
   args.feature_mask = feature_mask;
-  rtcIntersectV(data.scene,RTCRayHit_(ray),&args);
+  rtcTraversableIntersectV(data.traversable,RTCRayHit_(ray),&args);
   RayStats_addRay(stats);
 
   /* shade pixel */
@@ -366,7 +371,7 @@ Vec3f renderPixelAOShader(const uniform DebugShaderData& data, float x, float y,
     uniform RTCOccludedArguments args;
     rtcInitOccludedArguments(&args);
     args.feature_mask = feature_mask;
-    rtcOccludedV(data.scene,RTCRay_(shadow),&args);
+    rtcTraversableOccludedV(data.traversable,RTCRay_(shadow),&args);
     RayStats_addShadowRay(stats);
 
     /* add light contribution */
@@ -379,5 +384,10 @@ Vec3f renderPixelAOShader(const uniform DebugShaderData& data, float x, float y,
   return col * intensity;
 }
 
+#if defined(EMBREE_SYCL_TUTORIAL) && !defined(EMBREE_SYCL_RT_SIMULATION)
+RENDER_FRAME_FUNCTION_SYCL(DebugShader)
+RENDER_FRAME_FUNCTION_SYCL(AOShader)
+#else
 RENDER_FRAME_FUNCTION_ISPC(DebugShader)
 RENDER_FRAME_FUNCTION_ISPC(AOShader)
+#endif