[WIP] Add ISPC API using RTCTraversable

common/cmake/dpcpp.cmake

@@ -52,7 +52,7 @@ IF (EMBREE_SYCL_SUPPORT)
     SET(CMAKE_CXX_FLAGS_SYCL "${CMAKE_CXX_FLAGS_SYCL} /debug:none")    # FIXME: debug information generation takes forever in SYCL
     SET(CMAKE_CXX_FLAGS_SYCL "${CMAKE_CXX_FLAGS_SYCL} /DNDEBUG")    # FIXME: debug information generation takes forever in SYCL
   ELSE()
-    SET(CMAKE_CXX_FLAGS_SYCL "${CMAKE_CXX_FLAGS_SYCL} -g0")            # FIXME: debug information generation takes forever in SYCL
+    #SET(CMAKE_CXX_FLAGS_SYCL "${CMAKE_CXX_FLAGS_SYCL} -g0")            # FIXME: debug information generation takes forever in SYCL
     SET(CMAKE_CXX_FLAGS_SYCL "${CMAKE_CXX_FLAGS_SYCL} -UDEBUG -DNDEBUG") # FIXME: assertion still not working in SYCL
   ENDIF()
 

include/embree4/rtcore_buffer.h

@@ -38,14 +38,16 @@ typedef struct RTCBufferTy* RTCBuffer;
 /* Creates a new buffer. */
 RTC_API RTCBuffer rtcNewBuffer(RTCDevice device, size_t byteSize);
 
+/* Creates a new buffer using explicit host device memory. */
 RTC_API RTCBuffer rtcNewBufferHostDevice(RTCDevice device, size_t byteSize);
 
 /* Creates a new shared buffer. */
 RTC_API RTCBuffer rtcNewSharedBuffer(RTCDevice device, void* ptr, size_t byteSize);
 
-/* Creates a new shared buffer. */
+/* Creates a new shared buffer using explicit host device memory. */
 RTC_API RTCBuffer rtcNewSharedBufferHostDevice(RTCDevice device, void* ptr, size_t byteSize);
 
+/* Synchronize host and device memory by copying data from host to device. */
 RTC_API void rtcCommitBuffer(RTCBuffer buffer);
 
 #if defined(EMBREE_SYCL_SUPPORT) && defined(SYCL_LANGUAGE_VERSION)

include/embree4/rtcore_buffer.isph

@@ -37,12 +37,25 @@ typedef uniform struct RTCBufferTy* uniform RTCBuffer;
 /* Creates a new buffer. */
 RTC_API RTCBuffer rtcNewBuffer(RTCDevice device, uniform uintptr_t byteSize);
 
+/* Creates a new buffer using explicit host device memory. */
+RTC_API RTCBuffer rtcNewBufferHostDevice(RTCDevice device, uniform size_t byteSize);
+
 /* Creates a new shared buffer. */
 RTC_API RTCBuffer rtcNewSharedBuffer(RTCDevice device, void* uniform ptr, uniform uintptr_t byteSize);
 
+/* Creates a new shared buffer using explicit host device memory. */
+RTC_API RTCBuffer rtcNewSharedBufferHostDevice(RTCDevice device, void* uniform ptr, uniform uintptr_t byteSize);
+
+RTC_API void rtcCommitBuffer(RTCBuffer buffer);
+
 /* Returns a pointer to the buffer data. */
 RTC_API void* uniform rtcGetBufferData(RTCBuffer buffer);
 
+/* Returns a pointer to the buffer data on the device. Returns the same pointer as
+  rtcGetBufferData if the device is no SYCL device or if Embree is executed on a
+  system with unified memory (e.g., iGPUs). */
+RTC_API void* uniform rtcGetBufferDataDevice(RTCBuffer buffer);
+
 /* Retains the buffer (increments the reference count). */
 RTC_API void rtcRetainBuffer(RTCBuffer buffer);
 

include/embree4/rtcore_device.h

@@ -32,10 +32,6 @@ RTC_API int rtcSYCLDeviceSelector(const sycl::device sycl_device);
 /* Set the SYCL device to be used to allocate data */
 RTC_API void rtcSetDeviceSYCLDevice(RTCDevice device, const sycl::device sycl_device);
 
-/* rtcCommitGeometryWithQueue is asynchronous, user has to call queue.wait()
-   for synchronization. rtcCommitGemometry is blocking. */
-RTC_API void rtcCommitGeometryWithQueue(RTCScene scene, sycl::queue queue);
-
 /* rtcCommitSceneWithQueue is asynchronous, user has to call queue.wait()
    for synchronization. rtcCommitScene is blocking. */
 RTC_API void rtcCommitSceneWithQueue(RTCScene scene, sycl::queue queue);

include/embree4/rtcore_scene.h

@@ -6,7 +6,7 @@
 #include "rtcore_device.h"
 
 RTC_NAMESPACE_BEGIN
-  
+
 /* Opaque traversable type */
 typedef struct RTCTraversableTy* RTCTraversable;
 

include/embree4/rtcore_scene.isph

@@ -6,6 +6,9 @@
 
 #include "rtcore_device.isph"
 
+/* Opaque traversable type */
+typedef uniform struct RTCTraversableTy* uniform RTCTraversable;
+
 /* Forward declarations for ray structures */
 struct RTCRayHit;
 
@@ -82,6 +85,8 @@ RTC_API void rtcRetainScene(RTCScene scene);
 /* Releases the scene (decrements the reference count). */
 RTC_API void rtcReleaseScene(RTCScene scene);
 
+/* Returns the traversable object of the scene which can be passed to ray queries. */
+RTC_API RTCTraversable rtcGetSceneTraversable(RTCScene scene);
 
 /* Attaches the geometry to a scene. */
 RTC_API uniform unsigned int rtcAttachGeometry(RTCScene scene, RTCGeometry geometry);
@@ -147,6 +152,18 @@ inline void rtcGetGeometryTransformFromScene(RTCScene scene, varying unsigned in
   }
 }
 
+/* Gets the user-defined data pointer of the geometry. This function is not thread safe and should get used during rendering. */
+RTC_API void* rtcGetGeometryUserDataFromTraversable(RTCTraversable traversable, uniform unsigned int geomID);
+
+/* Returns the interpolated transformation of an instance for the specified time. */
+RTC_API void rtcGetGeometryTransformFromTraversable(RTCTraversable traversable, uniform unsigned int geomID, uniform float time, uniform RTCFormat format, void* uniform xfm);
+
+/* Returns the interpolated transformation of an instance for the specified time. Varying version. */
+inline void rtcGetGeometryTransformFromTraversable(RTCTraversable traversable, varying unsigned int geomID, varying float time, uniform RTCFormat format, void* uniform xfm)
+{
+  rtcGetGeometryTransformFromScene((RTCScene)traversable, geomID, time, format, xfm);
+}
+
 /* Commits the scene. */
 RTC_API void rtcCommitScene(RTCScene scene);
 
@@ -374,6 +391,133 @@ RTC_FORCEINLINE void rtcForwardOccludedVEx(const uniform RTCOccludedFunctionNArg
 }
 
 
+/* perform a closest point query of the scene. */
+RTC_API bool rtcTraversablePointQuery(RTCTraversable traversable, uniform RTCPointQuery* uniform query, uniform RTCPointQueryContext* uniform context, RTCPointQueryFunction queryFunc, void* uniform userPtr);
+
+/* Perform a closest point query with a packet of 4 points with the scene. */
+RTC_API bool rtcTraversablePointQuery4(const int* uniform valid, RTCTraversable traversable, void* uniform query, uniform RTCPointQueryContext* uniform context, RTCPointQueryFunction queryFunc, void * varying * uniform userPtr);
+
+/* Perform a closest point query with a packet of 4 points with the scene. */
+RTC_API bool rtcTraversablePointQuery8(const int* uniform valid, RTCTraversable traversable, void* uniform query, uniform RTCPointQueryContext* uniform context, RTCPointQueryFunction queryFunc, void * varying * uniform userPtr);
+
+/* Perform a closest point query with a packet of 4 points with the scene. */
+RTC_API bool rtcTraversablePointQuery16(const int* uniform valid, RTCTraversable traversable, void* uniform query, uniform RTCPointQueryContext* uniform context, RTCPointQueryFunction queryFunc, void * varying * uniform userPtr);
+
+/* Intersects a varying ray with the scene. */
+RTC_FORCEINLINE bool rtcTraversablePointQueryV(RTCTraversable traversable, varying RTCPointQuery* uniform query, uniform RTCPointQueryContext* uniform context, RTCPointQueryFunction queryFunc, void * varying * uniform userPtr)
+{
+  return rtcPointQueryV((RTCScene)traversable, query, context, queryFunc, userPtr);
+}
+
+/* Intersects a single ray with the scene. */
+RTC_API void rtcTraversableIntersect1(RTCTraversable traversable, uniform RTCRayHit* uniform rayhit, uniform RTCIntersectArguments* uniform args = NULL);
+
+/* Intersects a packet of 4 rays with the scene. */
+RTC_API void rtcTraversableIntersect4(const int* uniform valid, RTCTraversable traversable, void* uniform rayhit, uniform RTCIntersectArguments* uniform args = NULL);
+
+/* Intersects a packet of 8 rays with the scene. */
+RTC_API void rtcTraversableIntersect8(const int* uniform valid, RTCTraversable traversable, void* uniform rayhit, uniform RTCIntersectArguments* uniform args = NULL);
+
+/* Intersects a packet of 16 rays with the scene. */
+RTC_API void rtcTraversableIntersect16(const int* uniform valid, RTCTraversable traversable, void* uniform rayhit, uniform RTCIntersectArguments* uniform args = NULL);
+
+/* Intersects a varying ray with the scene. */
+RTC_FORCEINLINE void rtcTraversableIntersectV(RTCTraversable traversable, varying RTCRayHit* uniform rayhit, uniform RTCIntersectArguments* uniform args = NULL)
+{
+  rtcIntersectV((RTCScene)traversable, rayhit, args);
+}
+
+/* Forwards ray inside user geometry callback. */
+RTC_API void rtcTraversableForwardIntersect1(const uniform RTCIntersectFunctionNArguments* uniform args, RTCTraversable traversable, uniform RTCRay* uniform ray, uniform unsigned int instID);
+
+/* Forwards ray inside user geometry callback. Extended to handle instance arrays using instPrimID parameter. */
+RTC_API void rtcTraversableForwardIntersect1Ex(const uniform RTCIntersectFunctionNArguments* uniform args, RTCTraversable traversable, uniform RTCRay* uniform ray, uniform unsigned int instID, uniform unsigned int instPrimID);
+
+/* Forwards ray packet of size 4 inside user geometry callback. */
+RTC_API void rtcTraversableForwardIntersect4(const uniform int* uniform valid, const uniform RTCIntersectFunctionNArguments* uniform args, RTCTraversable traversable, void* uniform ray4, uniform unsigned int instID);
+
+/* Forwards ray packet of size 4 inside user geometry callback. Extended to handle instance arrays using instPrimID parameter. */
+RTC_API void rtcTraversableForwardIntersect4Ex(const uniform int* uniform valid, const uniform RTCIntersectFunctionNArguments* uniform args, RTCTraversable traversable, void* uniform ray4, uniform unsigned int instID, uniform unsigned int instPrimID);
+
+/* Forwards ray packet of size 8 inside user geometry callback. */
+RTC_API void rtcTraversableForwardIntersect8(const uniform int* uniform valid, const uniform RTCIntersectFunctionNArguments* uniform args, RTCTraversable traversable, void* uniform ray8, uniform unsigned int instID);
+
+/* Forwards ray packet of size 8 inside user geometry callback. Extended to handle instance arrays using instPrimID parameter. */
+RTC_API void rtcTraversableForwardIntersect8Ex(const uniform int* uniform valid, const uniform RTCIntersectFunctionNArguments* uniform args, RTCTraversable traversable, void* uniform ray8, uniform unsigned int instID, uniform unsigned int instPrimID);
+
+/* Forwards ray packet of size 16 inside user geometry callback. */
+RTC_API void rtcTraversableForwardIntersect16(const uniform int* uniform valid, const uniform RTCIntersectFunctionNArguments* uniform args, RTCTraversable traversable, void* uniform ray16, uniform unsigned int instID);
+
+/* Forwards ray packet of size 16 inside user geometry callback. Extended to handle instance arrays using instPrimID parameter. */
+RTC_API void rtcTraversableForwardIntersect16Ex(const uniform int* uniform valid, const uniform RTCIntersectFunctionNArguments* uniform args, RTCTraversable traversable, void* uniform ray16, uniform unsigned int instID, uniform unsigned int instPrimID);
+
+/* Forwards ray intersection query inside user geometry callback. */
+RTC_FORCEINLINE void rtcTraversableForwardIntersectV(const uniform RTCIntersectFunctionNArguments* uniform args, RTCTraversable traversable, varying RTCRay* uniform iray, uniform unsigned int instID)
+{
+  rtcForwardIntersectV(args, (RTCScene)traversable, iray, instID);
+}
+
+/* Forwards ray intersection query inside user geometry callback. Extended to handle instance arrays using instPrimID parameter. */
+RTC_FORCEINLINE void rtcTraversableForwardIntersectVEx(const uniform RTCIntersectFunctionNArguments* uniform args, RTCTraversable traversable, varying RTCRay* uniform iray, uniform unsigned int instID, uniform unsigned int instPrimID)
+{
+  rtcForwardIntersectVEx(args, (RTCScene)traversable, iray, instID, instPrimID);
+}
+
+/* Tests a single ray for occlusion with the scene. */
+RTC_API void rtcTraversableOccluded1(RTCTraversable traversable, uniform RTCRay* uniform ray, uniform RTCOccludedArguments* uniform args = NULL);
+
+/* Tests a packet of 4 rays for occlusion occluded with the scene. */
+RTC_API void rtcTraversableOccluded4(const uniform int* uniform valid, RTCTraversable traversable, void* uniform ray, uniform RTCOccludedArguments* uniform args = NULL);
+
+/* Tests a packet of 8 rays for occlusion occluded with the scene. */
+RTC_API void rtcTraversableOccluded8(const uniform int* uniform valid, RTCTraversable traversable, void* uniform ray, uniform RTCOccludedArguments* uniform args = NULL);
+
+/* Tests a packet of 16 rays for occlusion occluded with the scene. */
+RTC_API void rtcTraversableOccluded16(const uniform int* uniform valid, RTCTraversable traversable, void* uniform ray, uniform RTCOccludedArguments* uniform args = NULL);
+
+/* Tests a varying ray for occlusion with the scene. */
+RTC_FORCEINLINE void rtcTraversableOccludedV(RTCTraversable traversable, varying RTCRay* uniform ray, uniform RTCOccludedArguments* uniform args = NULL)
+{
+  rtcOccludedV((RTCScene)traversable, ray, args);
+}
+
+
+/* Forwards single occlusion ray inside user geometry callback. */
+RTC_API void rtcTraversableForwardOccluded1(const uniform RTCOccludedFunctionNArguments* uniform args, RTCTraversable traversable, uniform RTCRay* uniform ray, uniform unsigned int instID);
+
+/* Forwards single occlusion ray inside user geometry callback.  Extended to handle instance arrays using instPrimID parameter.*/
+RTC_API void rtcTraversableForwardOccluded1Ex(const uniform RTCOccludedFunctionNArguments* uniform args, RTCTraversable traversable, uniform RTCRay* uniform ray, uniform unsigned int instID, uniform unsigned int instPrimID);
+
+/* Forwards occlusion ray packet of size 4 inside user geometry callback. */
+RTC_API void rtcTraversableForwardOccluded4(const uniform int* uniform valid, const uniform RTCOccludedFunctionNArguments* uniform args, RTCTraversable traversable, void* uniform ray4, uniform unsigned int instID);
+
+/* Forwards occlusion ray packet of size 4 inside user geometry callback. Extended to handle instance arrays using instPrimID parameter. */
+RTC_API void rtcTraversableForwardOccluded4Ex(const uniform int* uniform valid, const uniform RTCOccludedFunctionNArguments* uniform args, RTCTraversable traversable, void* uniform ray4, uniform unsigned int instID, uniform unsigned int instPrimID);
+
+/* Forwards occlusion ray packet of size 8 inside user geometry callback. */
+RTC_API void rtcTraversableForwardOccluded8(const uniform int* uniform valid, const uniform RTCOccludedFunctionNArguments* uniform args, RTCTraversable traversable, void* uniform ray8, uniform unsigned int instID);
+
+/* Forwards occlusion ray packet of size 8 inside user geometry callback. Extended to handle instance arrays using instPrimID parameter. */
+RTC_API void rtcTraversableForwardOccluded8Ex(const uniform int* uniform valid, const uniform RTCOccludedFunctionNArguments* uniform args, RTCTraversable traversable, void* uniform ray8, uniform unsigned int instID, uniform unsigned int instPrimID);
+
+/* Forwards occlusion ray packet of size 16 inside user geometry callback. */
+RTC_API void rtcTraversableForwardOccluded16(const uniform int* uniform valid, const uniform RTCOccludedFunctionNArguments* uniform args, RTCTraversable traversable, void* uniform ray16, uniform unsigned int instID);
+
+/* Forwards occlusion ray packet of size 16 inside user geometry callback. Extended to handle instance arrays using instPrimID parameter. */
+RTC_API void rtcTraversableForwardOccluded16Ex(const uniform int* uniform valid, const uniform RTCOccludedFunctionNArguments* uniform args, RTCTraversable traversable, void* uniform ray16, uniform unsigned int instID, uniform unsigned int instPrimID);
+
+/* Forwards ray occlusion query inside user geometry callback. */
+RTC_FORCEINLINE void rtcTraversableForwardOccludedV(const uniform RTCOccludedFunctionNArguments* uniform args, RTCTraversable traversable, varying RTCRay* uniform iray, uniform unsigned int instID)
+{
+  rtcForwardOccludedV(args, (RTCScene)traversable, iray, instID);
+}
+
+/* Forwards ray occlusion query inside user geometry callback. Extended to handle instance arrays using instPrimID parameter. */
+RTC_FORCEINLINE void rtcTraversableForwardOccludedVEx(const uniform RTCOccludedFunctionNArguments* uniform args, RTCTraversable traversable, varying RTCRay* uniform iray, uniform unsigned int instID, uniform unsigned int instPrimID)
+{
+  rtcForwardOccludedVEx(args, (RTCScene)traversable, iray, instID, instPrimID);
+}
+
 /*! collision callback */
 struct RTCCollision { unsigned int geomID0; unsigned int primID0; unsigned int geomID1; unsigned int primID1; };
 typedef unmasked void (* uniform RTCCollideFunc) (void* uniform userPtr, uniform RTCCollision* uniform collisions, uniform unsigned int num_collisions);

kernels/sycl/rthwif_embree_builder.cpp

@@ -673,7 +673,7 @@ namespace embree
       throw_RTCError(RTC_ERROR_INVALID_ARGUMENT, "AccelBuffer constructor called with non-GPU device");
     }
 
-    unifiedMemory = gpu_device->has_unified_memory();
+    unifiedMemory = true; //gpu_device->has_unified_memory();
 
     if (unifiedMemory)
     {

tutorials/multi_instanced_geometry/multi_instanced_geometry_device.ispc

@@ -231,7 +231,7 @@ void renderPixelStandard(const uniform TutorialData& data, int x, int y,
 
   RandomSampler sampler;
   Ray primaryRay = samplePrimaryRay(data, x, 0, y, 0, camera, sampler, stats);
-  rtcIntersectV(data.g_scene, RTCRayHit_(primaryRay),&iargs);
+  rtcTraversableIntersectV(data.g_traversable, RTCRayHit_(primaryRay), &iargs);
 
   Vec3f color = make_Vec3f(0.f);
   if (primaryRay.geomID != RTC_INVALID_GEOMETRY_ID)
@@ -303,6 +303,7 @@ export void device_init(uniform int8* uniform cfg)
 {
   TutorialData_Constructor(&g_data);
   g_scene = g_data.g_scene = initializeScene(g_data, g_device);
+  g_data.g_traversable = rtcGetSceneTraversable(g_scene);
 }
 
 

tutorials/multi_instanced_geometry/multi_instanced_geometry_device.isph

@@ -18,6 +18,7 @@ struct InstanceLevels
 struct TutorialData
 {
   RTCScene g_scene;
+  RTCTraversable g_traversable;
   uniform InstanceLevels g_instanceLevels;
 
   /* accumulation buffer */
@@ -48,7 +49,7 @@ extern "C" void cleanupScene(uniform TutorialData& data);
 
 inline void TutorialData_Constructor(uniform TutorialData* uniform This)
 {
-  This->g_scene  = NULL;
+  This->g_scene = 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*)alignedUSMMalloc(2*sizeof(unsigned int), 16);
+  data.g_instanceLevels.numInstancesOnLevel = (unsigned int*)alignedMalloc(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**)alignedUSMMalloc(2*sizeof(LinearSpace3fa*), 16);
-  data.g_normalTransforms[0] = (LinearSpace3fa*)alignedUSMMalloc(data.g_instanceLevels.numInstancesOnLevel[0]*sizeof(LinearSpace3fa), 16);
-  data.g_normalTransforms[1] = (LinearSpace3fa*)alignedUSMMalloc(data.g_instanceLevels.numInstancesOnLevel[1]*sizeof(LinearSpace3fa), 16);
+  data.g_normalTransforms = (LinearSpace3fa**)alignedMalloc(2*sizeof(LinearSpace3fa*), 16);
+  data.g_normalTransforms[0] = (LinearSpace3fa*)alignedMalloc(data.g_instanceLevels.numInstancesOnLevel[0]*sizeof(LinearSpace3fa), 16);
+  data.g_normalTransforms[1] = (LinearSpace3fa*)alignedMalloc(data.g_instanceLevels.numInstancesOnLevel[1]*sizeof(LinearSpace3fa), 16);
   data.g_instanceLevels.normalTransforms = data.g_normalTransforms;
 
   RandomSampler sampler;