fix issues with host/device memory API on iGPU

RenderKit · Dec 13, 2024 · 76f8c1f · 76f8c1f
1 parent 9950968
commit 76f8c1f
Show file tree

Hide file tree

Showing 3 changed files with 31 additions and 5 deletions.
diff --git a/kernels/common/buffer.h b/kernels/common/buffer.h
@@ -61,8 +61,16 @@ namespace embree
       else if (device->is_gpu() && device->has_unified_memory())
       {
         ptr = alloc(ptr_in, shared, EmbreeMemoryType::USM_SHARED);
-        dshared = true;
-        dptr = ptr;
+
+        if (device->get_memory_type(ptr) != EmbreeMemoryType::USM_SHARED)
+        {
+          dptr = alloc(dptr_in, dshared, EmbreeMemoryType::USM_DEVICE);
+        }
+        else
+        {
+          dshared = true;
+          dptr = ptr;
+        }
       }
       else
 #endif
@@ -92,6 +100,11 @@ namespace embree
     virtual void free()
     {
       if (!shared && ptr) {
+#if defined(EMBREE_SYCL_SUPPORT)
+        if (dptr == ptr) {
+          dptr = nullptr;
+        }
+#endif
         device->free(ptr);
         device->memoryMonitor(-ssize_t(this->bytes()), true);
         ptr = nullptr;

diff --git a/kernels/common/device.h b/kernels/common/device.h
@@ -132,6 +132,8 @@ namespace embree
     /*! returns true if device and host have shared memory system (e.g., integrated GPU) */
     virtual bool has_unified_memory() const { return true; }
 
+    virtual EmbreeMemoryType get_memory_type(void* ptr) const { return EmbreeMemoryType::MALLOC; }
+
   private:
 
     /*! initializes the tasking system */
@@ -192,6 +194,15 @@ namespace embree
     /*! returns true if device and host have shared memory system (e.g., integrated GPU) */
     virtual bool has_unified_memory() const override;
 
+    virtual EmbreeMemoryType get_memory_type(void* ptr) const override {
+      switch(sycl::get_pointer_type(ptr, gpu_context)) {
+        case sycl::usm::alloc::host: return EmbreeMemoryType::USM_HOST;
+        case sycl::usm::alloc::device: return EmbreeMemoryType::USM_DEVICE;
+        case sycl::usm::alloc::shared: return EmbreeMemoryType::USM_SHARED;
+        default: return EmbreeMemoryType::MALLOC;
+      }
+    }
+
   private:
     sycl::context gpu_context;
     sycl::device  gpu_device;

diff --git a/tutorials/host_device_memory/host_device_memory_device.cpp b/tutorials/host_device_memory/host_device_memory_device.cpp
@@ -141,7 +141,6 @@ unsigned int addCubeShared (RTCScene scene, Vec3fa d)
     triangles[11].v0 = 3; triangles[11].v1 = 5; triangles[11].v2 = 7;
 #if defined(EMBREE_SYCL_TUTORIAL)
   });
-  global_gpu_queue->wait_and_throw();
 #endif
 
   rtcSetSharedGeometryBuffer(mesh, RTC_BUFFER_TYPE_INDEX, 0, RTC_FORMAT_UINT3, triangles, 0, sizeof(Triangle), 12);
@@ -275,7 +274,6 @@ unsigned int addCubeBufferShared (RTCScene scene, Vec3fa d)
     triangles[11].v0 = 3; triangles[11].v1 = 5; triangles[11].v2 = 7;
 #if defined(EMBREE_SYCL_TUTORIAL)
   });
-  global_gpu_queue->wait_and_throw();
 #endif
 
   RTCBuffer triangleBuffer = rtcNewSharedBuffer(g_device, triangles, 12 * sizeof(Triangle));
@@ -356,8 +354,12 @@ extern "C" void device_init (char* cfg)
   /* add ground plane */
   addGroundPlane(data.g_scene);
 
-  /* commit changes to scene */
 #if defined(EMBREE_SYCL_SUPPORT) && defined(EMBREE_SYCL_TUTORIAL)
+  // we fill data on GPU and copy to host. we have to ensure
+  // that the data is on host before building the scene
+  // with rtcCommitScene
+  global_gpu_queue->wait_and_throw();
+
   rtcCommitSceneWithQueue (data.g_scene, *global_gpu_queue);
 #else
   rtcCommitScene (data.g_scene);