[WIP] scene is now in device memory, too. USM shared memory is now only

used in tutorials.
RenderKit · Nov 6, 2024 · b8e4966 · b8e4966
1 parent 8eb2895
commit b8e4966
Show file tree

Hide file tree

Showing 28 changed files with 66 additions and 11 deletions.
diff --git a/common/sys/alloc.h b/common/sys/alloc.h
@@ -38,11 +38,6 @@ namespace embree
     ALIGNED_STRUCT_(align)                                             \
  private:
 
-#define ALIGNED_CLASS_USM_(align)                                          \
- public:                                                               \
-    ALIGNED_STRUCT_USM_(align)                                             \
- private:
-
   enum EmbreeUSMMode {
     EMBREE_USM_SHARED = 0,
     EMBREE_USM_SHARED_DEVICE_READ_WRITE = 0,

diff --git a/include/embree4/rtcore_scene.h b/include/embree4/rtcore_scene.h
@@ -92,6 +92,8 @@ RTC_API void rtcRetainScene(RTCScene scene);
 /* Releases the scene (decrements the reference count). */
 RTC_API void rtcReleaseScene(RTCScene scene);
 
+/* Returns the representation of the scene which is accessible on the device */
+RTC_API RTCScene rtcGetSceneDevicePointer(RTCScene scene);
 
 /* Attaches the geometry to a scene. */
 RTC_API unsigned int rtcAttachGeometry(RTCScene scene, RTCGeometry geometry);

diff --git a/kernels/common/rtcore.cpp b/kernels/common/rtcore.cpp
@@ -266,6 +266,19 @@ RTC_NAMESPACE_BEGIN;
     return (RTCDevice)nullptr;
   }
 
+  RTC_API RTCScene rtcGetSceneDevicePointer(RTCScene hscene)
+  {
+    Scene* scene = (Scene*) hscene;
+    RTC_CATCH_BEGIN;
+    RTC_TRACE(rtcGetSceneDevice);
+    RTC_VERIFY_HANDLE(hscene);
+    if (scene->getDevicePointer() == nullptr)
+      throw_RTCError(RTC_ERROR_INVALID_OPERATION,"scene device pointer is NULL");
+    return (RTCScene)scene->getDevicePointer();
+    RTC_CATCH_END2(scene);
+    return (RTCScene)nullptr;
+  }
+
   RTC_API void rtcSetSceneProgressMonitorFunction(RTCScene hscene, RTCProgressMonitorFunction progress, void* ptr) 
   {
     Scene* scene = (Scene*) hscene;

diff --git a/kernels/common/scene.cpp b/kernels/common/scene.cpp
@@ -42,6 +42,7 @@ namespace embree
 
   Scene::Scene (Device* device)
     : device(device),
+      scene_device(nullptr),
       flags_modified(true), enabled_geometry_types(0),
       scene_flags(RTC_SCENE_FLAG_NONE),
       quality_flags(RTC_BUILD_QUALITY_MEDIUM),
@@ -79,6 +80,9 @@ namespace embree
     if (geometries_device) {
       device->free(geometries_device);
     }
+    if (scene_device) {
+      device->free(scene_device);
+    }
 #endif
 
     device->refDec();

diff --git a/kernels/common/scene.h b/kernels/common/scene.h
@@ -37,7 +37,7 @@ namespace embree
   /*! Base class all scenes are derived from */
   class Scene : public AccelN
   {
-    ALIGNED_CLASS_USM_(std::alignment_of<Scene>::value);
+    ALIGNED_CLASS_(std::alignment_of<Scene>::value);
 
   public:
     template<typename Ty, bool mblur = false>
@@ -317,11 +317,21 @@ namespace embree
 
 #if defined(EMBREE_SYCL_SUPPORT)
     void syncWithDevice(sycl::queue* queue);
+
+  public:
+    Scene* getDevicePointer() {
+      DeviceGPU* gpu_device = dynamic_cast<DeviceGPU*>(device);
+      if(!gpu_device) {
+        return this;
+      }
+      return scene_device;
+    }
 #endif
 
 
   public:
     Device* device;
+    Scene* scene_device;
 
   public:
     IDPool<unsigned,0xFFFFFFFE> id_pool;

diff --git a/kernels/common/scene_instance.cpp b/kernels/common/scene_instance.cpp
@@ -172,6 +172,7 @@ namespace embree
 
     // override local2world value with device ptr in geometries_data_host
     Instance* instance = (Instance*)(data_host + offsetInstance);
+    instance->object = ((Scene*)(instance->object))->getDevicePointer();
     instance->local2world = (AffineSpace3ff*)(data_device + offsetInstance + sizeof(Instance));
   }
 

diff --git a/kernels/common/scene_instance_array.cpp b/kernels/common/scene_instance_array.cpp
@@ -203,7 +203,11 @@ namespace embree
     offset += sizeof(InstanceArray);
 
     const size_t offsetObjects = offset;
-    std::memcpy(data_host + offset, objects, numObjects * sizeof(Accel*));
+    Accel** objects_host = (Accel**)(data_host + offsetObjects);
+    for (size_t i = 0; i < numObjects; ++i) {
+      objects_host[i] = (Accel*)((Scene*)objects[i])->getDevicePointer();
+    }
+
     offset += numObjects * sizeof(Accel*);
     iarray->objects = (Accel**)(data_device + offsetObjects);
 

diff --git a/kernels/sycl/scene_sycl.cpp b/kernels/sycl/scene_sycl.cpp
@@ -180,6 +180,12 @@ void Scene::syncWithDevice(sycl::queue* queue_in)
     device->free(offsets);
 
   } // run
+
+  if (scene_device) {
+    device->free(scene_device);
+  }
+  scene_device = (Scene*) device->malloc(sizeof(Scene), 16, EmbreeMemoryType::DEVICE);
+  queue.memcpy(scene_device, (void*)this, sizeof(Scene));
 
   if (!queue_in)
     queue.wait_and_throw();

diff --git a/tutorials/curve_geometry/curve_geometry_device.cpp b/tutorials/curve_geometry/curve_geometry_device.cpp
@@ -253,6 +253,7 @@ extern "C" void renderFrameStandard (int* pixels,
 {
 #if defined(EMBREE_SYCL_TUTORIAL) && !defined(EMBREE_SYCL_RT_SIMULATION)
   TutorialData ldata = data;
+  ldata.g_scene = rtcGetSceneDevicePointer(data.g_scene);
   sycl::event event = global_gpu_queue->submit([=](sycl::handler& cgh){
     const sycl::nd_range<2> nd_range = make_nd_range(height,width);
     cgh.parallel_for(nd_range,[=](sycl::nd_item<2> item) {

diff --git a/tutorials/debug_device_memory/debug_device_memory_device.cpp b/tutorials/debug_device_memory/debug_device_memory_device.cpp
@@ -411,6 +411,7 @@ extern "C" void renderFrameStandard (int* pixels,
 {
 #if defined(EMBREE_SYCL_TUTORIAL) && !defined(EMBREE_SYCL_RT_SIMULATION) && !defined(EMBREE_SYCL_RT_SIMULATION)
   TutorialData ldata = data;
+  ldata.g_scene = rtcGetSceneDevicePointer(data.g_scene);
   sycl::event event = global_gpu_queue->submit([=](sycl::handler& cgh){
     const sycl::nd_range<2> nd_range = make_nd_range(height,width);
     cgh.parallel_for(nd_range,[=](sycl::nd_item<2> item) {

diff --git a/tutorials/dynamic_scene/dynamic_scene_device.cpp b/tutorials/dynamic_scene/dynamic_scene_device.cpp
@@ -279,6 +279,7 @@ extern "C" void renderFrameStandard (int* pixels,
   /* render all pixels */
 #if defined(EMBREE_SYCL_TUTORIAL) && !defined(EMBREE_SYCL_RT_SIMULATION)
   TutorialData ldata = data;
+  ldata.g_scene = rtcGetSceneDevicePointer(data.g_scene);
   sycl::event event = global_gpu_queue->submit([=](sycl::handler& cgh){
     const sycl::nd_range<2> nd_range = make_nd_range(height,width);
     cgh.parallel_for(nd_range,[=](sycl::nd_item<2> item) {

diff --git a/tutorials/forest/forest_device.cpp b/tutorials/forest/forest_device.cpp
@@ -151,7 +151,7 @@ void update_trees(float time)
   RTCBounds bounds;
   rtcGetSceneBounds(scene_terrain, &bounds);
 
-  RTCScene scene = scene_terrain;
+  RTCScene scene = rtcGetSceneDevicePointer(scene_terrain);
   TutorialData ldata = data;
   unsigned int lnum_trees_sqrt = num_trees_sqrt;
 
@@ -430,6 +430,7 @@ extern "C" void renderFrameStandard (int* pixels,
 {
 #if defined(EMBREE_SYCL_TUTORIAL) && !defined(EMBREE_SYCL_RT_SIMULATION)
   TutorialData ldata = data;
+  ldata.g_scene = rtcGetSceneDevicePointer(data.g_scene);
   sycl::event event = global_gpu_queue->submit([=](sycl::handler& cgh){
     const sycl::nd_range<2> nd_range = make_nd_range(height,width);
     cgh.parallel_for(nd_range,[=](sycl::nd_item<2> item) {

diff --git a/tutorials/grid_geometry/grid_geometry_device.cpp b/tutorials/grid_geometry/grid_geometry_device.cpp
@@ -670,6 +670,7 @@ extern "C" void renderFrameStandard (int* pixels,
   /* render image */
 #if defined(EMBREE_SYCL_TUTORIAL) && !defined(EMBREE_SYCL_RT_SIMULATION)
   TutorialData ldata = data;
+  ldata.g_scene = rtcGetSceneDevicePointer(data.g_scene);
   sycl::event event = global_gpu_queue->submit([=](sycl::handler& cgh){
     const sycl::nd_range<2> nd_range = make_nd_range(height,width);
     cgh.parallel_for(nd_range,[=](sycl::nd_item<2> item) {

diff --git a/tutorials/hair_geometry/hair_geometry_device.cpp b/tutorials/hair_geometry/hair_geometry_device.cpp
@@ -429,6 +429,7 @@ extern "C" void renderFrameStandard (int* pixels,
   /* render frame */
 #if defined(EMBREE_SYCL_TUTORIAL) && !defined(EMBREE_SYCL_RT_SIMULATION)
   TutorialData ldata = data;
+  ldata.scene = rtcGetSceneDevicePointer(data.scene);
   sycl::event event = global_gpu_queue->submit([=](sycl::handler& cgh){
     const sycl::nd_range<2> nd_range = make_nd_range(height,width);
     cgh.parallel_for(nd_range,[=](sycl::nd_item<2> item) {

diff --git a/tutorials/instanced_geometry/instanced_geometry_device.cpp b/tutorials/instanced_geometry/instanced_geometry_device.cpp
@@ -282,6 +282,7 @@ extern "C" void renderFrameStandard (int* pixels,
   /* render all pixels */
 #if defined(EMBREE_SYCL_TUTORIAL) && !defined(EMBREE_SYCL_RT_SIMULATION)
   TutorialData ldata = data;
+  ldata.g_scene = rtcGetSceneDevicePointer(data.g_scene);
   sycl::event event = global_gpu_queue->submit([=](sycl::handler& cgh){
     const sycl::nd_range<2> nd_range = make_nd_range(height,width);
     cgh.parallel_for(nd_range,[=](sycl::nd_item<2> item) {

diff --git a/tutorials/intersection_filter/intersection_filter_device.cpp b/tutorials/intersection_filter/intersection_filter_device.cpp
@@ -402,6 +402,7 @@ extern "C" void renderFrameStandard (int* pixels,
 {
 #if defined(EMBREE_SYCL_TUTORIAL) && !defined(EMBREE_SYCL_RT_SIMULATION)
   TutorialData ldata = data;
+  ldata.g_scene = rtcGetSceneDevicePointer(data.g_scene);
   sycl::event event = global_gpu_queue->submit([=](sycl::handler& cgh){
     const sycl::nd_range<2> nd_range = make_nd_range(height,width);
     cgh.parallel_for(nd_range,[=](sycl::nd_item<2> item) {

diff --git a/tutorials/minimal/minimal_sycl.cpp b/tutorials/minimal/minimal_sycl.cpp
@@ -203,10 +203,12 @@ RTCScene initializeScene(RTCDevice device, const sycl::queue& queue)
  * (dx, dy, dz).
  */
 
-void castRay(sycl::queue& queue, const RTCScene scene, 
+void castRay(sycl::queue& queue, const RTCScene scene_in, 
              float ox, float oy, float oz,
              float dx, float dy, float dz, Result* result)
 {
+  RTCScene scene = rtcGetSceneDevicePointer(scene_in);
+
   queue.submit([=](sycl::handler& cgh)
   {
     cgh.set_specialization_constant<feature_mask>(required_features);

diff --git a/tutorials/motion_blur_geometry/motion_blur_geometry_device.cpp b/tutorials/motion_blur_geometry/motion_blur_geometry_device.cpp
@@ -738,6 +738,7 @@ extern "C" void renderFrameStandard (int* pixels,
   /* render next frame */
 #if defined(EMBREE_SYCL_TUTORIAL) && !defined(EMBREE_SYCL_RT_SIMULATION)
   TutorialData ldata = data;
+  ldata.g_scene = rtcGetSceneDevicePointer(data.g_scene);
   sycl::event event = global_gpu_queue->submit([=](sycl::handler& cgh){
     const sycl::nd_range<2> nd_range = make_nd_range(height,width);
     cgh.parallel_for(nd_range,[=](sycl::nd_item<2> item) {

diff --git a/tutorials/multiscene_geometry/multiscene_geometry_device.cpp b/tutorials/multiscene_geometry/multiscene_geometry_device.cpp
@@ -300,6 +300,7 @@ namespace embree {
   {
 #if defined(EMBREE_SYCL_TUTORIAL) && !defined(EMBREE_SYCL_RT_SIMULATION)
   TutorialData ldata = data;
+  ldata.g_curr_scene = rtcGetSceneDevicePointer(data.g_curr_scene);
   sycl::event event = global_gpu_queue->submit([=](sycl::handler& cgh){
     const sycl::nd_range<2> nd_range = make_nd_range(height,width);
     cgh.parallel_for(nd_range,[=](sycl::nd_item<2> item) {

diff --git a/tutorials/next_hit/next_hit_device.cpp b/tutorials/next_hit/next_hit_device.cpp
@@ -467,6 +467,7 @@ extern "C" void renderFrameStandard (int* pixels,
 {
 #if defined(EMBREE_SYCL_TUTORIAL) && !defined(EMBREE_SYCL_RT_SIMULATION)
   TutorialData ldata = data;
+  ldata.scene = rtcGetSceneDevicePointer(data.scene);
 
 #if defined(USE_SPECIALIZATION_CONSTANTS)
   sycl::event event = global_gpu_queue->submit([=](sycl::handler& cgh) {

diff --git a/tutorials/pathtracer/pathtracer_device.cpp b/tutorials/pathtracer/pathtracer_device.cpp
@@ -1746,6 +1746,7 @@ extern "C" void renderFrameStandard (int* pixels,
 /* render image */
 #if defined(EMBREE_SYCL_TUTORIAL) && !defined(EMBREE_SYCL_RT_SIMULATION)
   TutorialData ldata = data;
+  ldata.scene = rtcGetSceneDevicePointer(data.scene);
 
 #if defined(USE_SPECIALIZATION_CONSTANTS)
   sycl::event event = global_gpu_queue->submit([=](sycl::handler& cgh){

diff --git a/tutorials/point_geometry/point_geometry_device.cpp b/tutorials/point_geometry/point_geometry_device.cpp
@@ -190,6 +190,7 @@ extern "C" void renderFrameStandard (int* pixels,
 {
 #if defined(EMBREE_SYCL_TUTORIAL) && !defined(EMBREE_SYCL_RT_SIMULATION)
   TutorialData ldata = data;
+  ldata.g_scene = rtcGetSceneDevicePointer(data.g_scene);
   sycl::event event = global_gpu_queue->submit([=](sycl::handler& cgh){
     const sycl::nd_range<2> nd_range = make_nd_range(height,width);
     cgh.parallel_for(nd_range,[=](sycl::nd_item<2> item) {

diff --git a/tutorials/quaternion_motion_blur/quaternion_motion_blur_device.cpp b/tutorials/quaternion_motion_blur/quaternion_motion_blur_device.cpp
@@ -389,6 +389,7 @@ extern "C" void renderFrameStandard (int* pixels,
 {
 #if defined(EMBREE_SYCL_TUTORIAL) && !defined(EMBREE_SYCL_RT_SIMULATION)
   TutorialData ldata = data;
+  ldata.g_scene = rtcGetSceneDevicePointer(data.g_scene);
   sycl::event event = global_gpu_queue->submit([=](sycl::handler& cgh){
     const sycl::nd_range<2> nd_range = make_nd_range(height,width);
     cgh.parallel_for(nd_range,[=](sycl::nd_item<2> item) {

diff --git a/tutorials/ray_mask/ray_mask_device.cpp b/tutorials/ray_mask/ray_mask_device.cpp
@@ -205,6 +205,7 @@ extern "C" void renderFrameStandard (int* pixels,
 {
 #if defined(EMBREE_SYCL_TUTORIAL) && !defined(EMBREE_SYCL_RT_SIMULATION)
   TutorialData ldata = data;
+  ldata.g_scene = rtcGetSceneDevicePointer(data.g_scene);
   sycl::event event = global_gpu_queue->submit([=](sycl::handler& cgh){
     const sycl::nd_range<2> nd_range = make_nd_range(height,width);
     cgh.parallel_for(nd_range,[=](sycl::nd_item<2> item) {

diff --git a/tutorials/triangle_geometry/triangle_geometry_device.cpp b/tutorials/triangle_geometry/triangle_geometry_device.cpp
@@ -192,6 +192,7 @@ extern "C" void renderFrameStandard (int* pixels,
 {
 #if defined(EMBREE_SYCL_TUTORIAL) && !defined(EMBREE_SYCL_RT_SIMULATION) && !defined(EMBREE_SYCL_RT_SIMULATION)
   TutorialData ldata = data;
+  ldata.g_scene = rtcGetSceneDevicePointer(data.g_scene);
   sycl::event event = global_gpu_queue->submit([=](sycl::handler& cgh){
     const sycl::nd_range<2> nd_range = make_nd_range(height,width);
     cgh.parallel_for(nd_range,[=](sycl::nd_item<2> item) {

diff --git a/tutorials/user_geometry/user_geometry_device.cpp b/tutorials/user_geometry/user_geometry_device.cpp
@@ -208,7 +208,7 @@ Instance* createInstance (RTCScene scene, RTCScene object, int geomID, const Vec
 #if !ENABLE_NATIVE_INSTANCING
   Instance* instance = (Instance*) alignedUSMMalloc(sizeof(Instance),16);
   instance->type = USER_GEOMETRY_INSTANCE;
-  instance->object = object;
+  instance->object = rtcGetSceneDevicePointer(object);
   instance->lower = lower;
   instance->upper = upper;
   instance->local2world.l.vx = Vec3fa(1,0,0);
@@ -872,6 +872,7 @@ extern "C" void renderFrameStandard (int* pixels,
 {
 #if defined(EMBREE_SYCL_TUTORIAL) && !defined(EMBREE_SYCL_RT_SIMULATION)
   TutorialData ldata = data;
+  ldata.g_scene = rtcGetSceneDevicePointer(data.g_scene);
   sycl::event event = global_gpu_queue->submit([=](sycl::handler& cgh){
     const sycl::nd_range<2> nd_range = make_nd_range(height,width);
     cgh.parallel_for(nd_range,[=](sycl::nd_item<2> item) {

diff --git a/tutorials/viewer/viewer_device.cpp b/tutorials/viewer/viewer_device.cpp
@@ -320,6 +320,7 @@ extern "C" void renderFrameStandard (int* pixels,
 {
 #if defined(EMBREE_SYCL_TUTORIAL) && !defined(EMBREE_SYCL_RT_SIMULATION)
   TutorialData ldata = data;
+  ldata.scene = rtcGetSceneDevicePointer(data.scene);
 
 #if defined(USE_SPECIALIZATION_CONSTANTS)
   sycl::event event = global_gpu_queue->submit([=](sycl::handler& cgh) {

diff --git a/tutorials/viewer/viewer_device_debug.cpp b/tutorials/viewer/viewer_device_debug.cpp
@@ -39,7 +39,7 @@ struct DebugShaderData
 
 void DebugShaderData_Constructor(DebugShaderData* This)
 {
-  This->scene = g_scene;
+  This->scene = rtcGetSceneDevicePointer(g_scene);
   This->ispc_scene = g_ispc_scene;
   This->scale = scale;
   This->debug = g_debug;