diff --git a/include/embree4/rtcore_scene.h b/include/embree4/rtcore_scene.h
index b02244cdc9..92222b1c96 100644
--- 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
index 04c0715e3b..8be5519f43 100644
--- 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
index 4c987b1aed..51b4d18047 100644
--- 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),
diff --git a/kernels/common/scene.h b/kernels/common/scene.h
index ae1911ff50..1d73e74c7a 100644
--- 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,17 @@ namespace embree
     
 #if defined(EMBREE_SYCL_SUPPORT)
     void syncWithDevice(sycl::queue* queue);
+  
+  public:
+    Scene* getDevicePointer() {
+      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
index 6a77c7ba6c..9340af33c2 100644
--- 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/sycl/scene_sycl.cpp b/kernels/sycl/scene_sycl.cpp
index 9bc89da8d1..917fdf2543 100644
--- a/kernels/sycl/scene_sycl.cpp
+++ b/kernels/sycl/scene_sycl.cpp
@@ -180,6 +180,9 @@ void Scene::syncWithDevice(sycl::queue* queue_in)
     device->free(offsets);
 
   } // run
+  
+  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
index d6822ae916..d9b106cd9e 100644
--- 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
index ac66c9fd2d..b2c081b49a 100644
--- 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
index 76b9717794..b8c28e830f 100644
--- 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
index 17f651f596..d94e841bbc 100644
--- a/tutorials/forest/forest_device.cpp
+++ b/tutorials/forest/forest_device.cpp
@@ -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
index 81d8dd618e..0d75278581 100644
--- 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
index adb1a2418a..42351777a5 100644
--- 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
index 0ee86f6d5f..b1a15b1c4d 100644
--- 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
index 1b893b1beb..77e91e49d7 100644
--- 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
index e180d43bbf..a5e44b011a 100644
--- 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
index 14e9b8fcf1..bc1c943592 100644
--- 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/next_hit/next_hit_device.cpp b/tutorials/next_hit/next_hit_device.cpp
index 23ddf97dea..1dea00b254 100644
--- 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/point_geometry/point_geometry_device.cpp b/tutorials/point_geometry/point_geometry_device.cpp
index f207617749..8bfee3502b 100644
--- 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/ray_mask/ray_mask_device.cpp b/tutorials/ray_mask/ray_mask_device.cpp
index b81b637cd7..b187cd72a4 100644
--- 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
index f7b257ff12..b21fba64bd 100644
--- 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
index 804b996cc5..f18a590ac3 100644
--- a/tutorials/user_geometry/user_geometry_device.cpp
+++ b/tutorials/user_geometry/user_geometry_device.cpp
@@ -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) {