diff --git a/RenderingToolkit/GettingStarted/03_openvkl_gsg/CMakeLists.txt b/RenderingToolkit/GettingStarted/03_openvkl_gsg/CMakeLists.txt
deleted file mode 100644
index 862e41d23d..0000000000
--- a/RenderingToolkit/GettingStarted/03_openvkl_gsg/CMakeLists.txt
+++ /dev/null
@@ -1,35 +0,0 @@
-cmake_minimum_required(VERSION 3.16)
-project(OPENVKL_GSG LANGUAGES C CXX)
-
-set(ONEAPI_ROOT "")
-if(DEFINED ENV{ONEAPI_ROOT})
-  set(ONEAPI_ROOT "$ENV{ONEAPI_ROOT}")
-  message(STATUS "ONEAPI_ROOT FROM ENVIRONMENT: ${ONEAPI_ROOT}")
-else()
-  if(WIN32)
-    set(ONEAPI_ROOT "C:/Program Files (x86)/Intel/oneAPI")
-  else()
-    set(ONEAPI_ROOT /opt/intel/oneapi)
-  endif()
-  message(STATUS "ONEAPI_ROOT DEFAULT: ${ONEAPI_ROOT}")
-endif(DEFINED ENV{ONEAPI_ROOT})
-
-set(OPENVKL_ROOT ${ONEAPI_ROOT}/openvkl/latest)
-set(OPENVKL_INCLUDE_DIR ${OPENVKL_ROOT}/include)
-message(STATUS "OPENVKL_INCLUDE_DIR: ${OPENVKL_INCLUDE_DIR}")
-find_package(openvkl REQUIRED PATHS ${ONEAPI_ROOT})
-
-if(MSVC)
-  set(CMAKE_CXX_STANDARD 11)
-  set(CMAKE_CXX_STANDARD_REQUIRED ON)
-  set(CMAKE_CXX_EXTENSIONS OFF)
-else()
-  set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++11")
-endif(MSVC)
-
-include_directories(${OPENVKL_INCLUDE_DIR})
-link_directories(${OPENVKL_ROOT}/lib)
-
-add_executable(vklTutorial src/vklTutorial.c)
-target_link_libraries(vklTutorial PRIVATE openvkl)
-install(TARGETS vklTutorial RUNTIME DESTINATION ${CMAKE_INSTALL_BINDIR})
diff --git a/RenderingToolkit/GettingStarted/03_openvkl_gsg/README.md b/RenderingToolkit/GettingStarted/03_openvkl_gsg/README.md
index 5b22f353a1..8aa6b400b0 100644
--- a/RenderingToolkit/GettingStarted/03_openvkl_gsg/README.md
+++ b/RenderingToolkit/GettingStarted/03_openvkl_gsg/README.md
@@ -5,63 +5,12 @@ high-performance volume computation kernels. Improve performance of volume
 rendering applications by using performance optimized volume traversal and
 sampling functionality for a variety of data formats.
 
-| Minimum Requirements              | Description
-|:---                               |:---
-| OS                                | Linux* Ubuntu* 18.04 <br>CentOS 8 (or compatible) <br> Windows* 10 <br>macOS* 10.15+
-| Hardware                          | Intel 64 Penryn or newer with SSE4.1 extensions, ARM64 with NEON extensions <br>(Optimized requirements: Intel 64 Skylake or newer with AVX512 extentions, ARM64 with NEON extensions) 
-| Compiler Toolchain                | Windows OS: MSVS 2019 installed with Windows* SDK and CMake*; Other platforms: C++11 compiler, a C99 compiler (for example. gcc/c++/clang), and CMake*
-| Libraries                         | Install Intel&reg; oneAPI Rendering Toolkit (Render Kit),  including Intel&reg; Embree and Intel&reg; Open VKL
+## Versions
 
-| Objective                         | Description
-|:---                               |:---
-| What you will learn               | How to build and run a basic rendering program using the Intel&reg; Open VKL API from the Render Kit.
-| Time to complete                  | 5 minutes
-
-## Purpose
-
-This sample program, `vklTutorial`, shows sampling amongst a proceedurally
-generated volume the different volumetric sampling capabilities with Intel&reg;
-Open VKL. Output is written to the console (stdout).
-
-## Key Implementation Details
-
-`vklTutorial` is written in C99 and is constructed to compile with a C++ or C99
-compiler.
-
-## Build and Run
-
-### Windows
-
-1. Run a new **x64 Native Tools Command Prompt for MSVS 2019**.
-
-```
-call <path-to-oneapi-folder>\setvars.bat
-cd <path-to-oneAPI-samples>\RenderingToolkit\GettingStarted\03_openvkl_gsg
-mkdir build
-cd build
-cmake ..
-cmake --build . --config Release
-cd Release
-vklTutorial.exe
-```
-
-2. Review the terminal output (stdout).
-
-
-### Linux and macOS
-
-1. Start a new Terminal session.
-```
-source <path-to-oneapi-folder>/setvars.sh
-cd <path-to-oneAPI-samples>/RenderingToolkit/GettingStarted/03_openvkl_gsg
-mkdir build
-cd build
-cmake ..
-cmake --build .
-./vklTutorial
-```
-
-2. Review the terminal output (stdout).
+1. [CPU](./cpu/) - for Intel64 (x86-64) Host
+- This version of the `vklTutorialCPU` program uses a C++11 (or C99) system compiler to target the host processor.
+2. [GPU](./gpu/) - for Intel&reg; Arc&trade; Graphics, Intel&reg; Data Center Flex Series, or Intel&reg; Data Center Max Series or higher (Xe-HPG, DG2-128, DG2-512 or higher)
+- This `vklTutorialGPU` program uses Intel&reg; oneAPI DPC/C++ Compiler and SYCL* Runtimes to target the GPU.
 
 
 ## License
@@ -70,4 +19,4 @@ This code sample is licensed under the Apache 2.0 license. See
 [LICENSE.txt](LICENSE.txt) for details.
 
 Third party program Licenses can be found here:
-[third-party-programs.txt](https://github.com/oneapi-src/oneAPI-samples/blob/master/third-party-programs.txt).
\ No newline at end of file
+[third-party-programs.txt](https://github.com/oneapi-src/oneAPI-samples/blob/master/third-party-programs.txt).
diff --git a/RenderingToolkit/GettingStarted/03_openvkl_gsg/cpu/CMakeLists.txt b/RenderingToolkit/GettingStarted/03_openvkl_gsg/cpu/CMakeLists.txt
new file mode 100644
index 0000000000..676a8956e3
--- /dev/null
+++ b/RenderingToolkit/GettingStarted/03_openvkl_gsg/cpu/CMakeLists.txt
@@ -0,0 +1,31 @@
+cmake_minimum_required(VERSION 3.16)
+project(OPENVKL_CPU_GSG LANGUAGES C CXX)
+
+set(ONEAPI_ROOT "")
+if(DEFINED ENV{ONEAPI_ROOT})
+  set(ONEAPI_ROOT "$ENV{ONEAPI_ROOT}")
+  message(STATUS "ONEAPI_ROOT FROM ENVIRONMENT: ${ONEAPI_ROOT}")
+else()
+  message(FATAL_ERROR "ONEAPI_ROOT DEFAULT: ${ONEAPI_ROOT}")
+endif(DEFINED ENV{ONEAPI_ROOT})
+
+set(OPENVKL_BASE_DIR "")
+if(EXISTS ${ONEAPI_ROOT}/oneapi-vars.sh OR EXISTS ${ONEAPI_ROOT}/oneapi-vars.bat)
+  set(OPENVKL_BASE_DIR ${ONEAPI_ROOT})
+else()
+  set(OPENVKL_BASE_DIR ${ONEAPI_ROOT}/openvkl/latest)
+endif(EXISTS ${ONEAPI_ROOT}/oneapi-vars.sh OR EXISTS ${ONEAPI_ROOT}/oneapi-vars.bat)
+
+find_package(openvkl REQUIRED PATHS ${ONEAPI_ROOT})
+
+if(MSVC)
+  set(CMAKE_CXX_STANDARD 11)
+  set(CMAKE_CXX_STANDARD_REQUIRED ON)
+  set(CMAKE_CXX_EXTENSIONS OFF)
+else()
+  set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++11")
+endif(MSVC)
+
+add_executable(vklTutorialCPU src/vklTutorialCPU.c)
+target_link_libraries(vklTutorialCPU PRIVATE openvkl::openvkl openvkl::openvkl_module_cpu_device)
+install(TARGETS vklTutorialCPU RUNTIME DESTINATION ${CMAKE_INSTALL_BINDIR})
diff --git a/RenderingToolkit/GettingStarted/03_openvkl_gsg/cpu/README.md b/RenderingToolkit/GettingStarted/03_openvkl_gsg/cpu/README.md
new file mode 100644
index 0000000000..05e42385b9
--- /dev/null
+++ b/RenderingToolkit/GettingStarted/03_openvkl_gsg/cpu/README.md
@@ -0,0 +1,75 @@
+# Getting Started Sample for Intel&reg; Rendering Toolkit (Render Kit): Intel&reg; Open Volume Kernel Library (Intel&reg; Open VKL) on CPU
+
+Intel&reg; Open Volume Kernel Library (Intel&reg; Open VKL) is a collection of
+high-performance volume computation kernels. Improve performance of volume
+rendering applications by using performance optimized volume traversal and
+sampling functionality for a variety of data formats.
+
+| Minimum Requirements              | Description
+|:---                               |:---
+| OS                                | Linux* Ubuntu* 22.04 <br>CentOS 8 (or compatible) <br> Windows* 10 or 11<br>macOS* 10.15+
+| Hardware                          | Intel 64 Penryn or newer with SSE4.1 extensions, ARM64 with NEON extensions <br>(Optimized requirements: Intel 64 Skylake or newer with AVX512 extentions, ARM64 with NEON extensions) 
+| Compiler Toolchain                | Windows OS: MSVS 2022 (or 2019) installed with Windows* SDK and CMake*; Other platforms: C++11 compiler, a C99 compiler (for example. gcc/c++/clang), and CMake*
+| Libraries                         | Install Intel&reg; Rendering Toolkit (Render Kit),  including Intel&reg; Embree and Intel&reg; Open VKL
+
+| Objective                         | Description
+|:---                               |:---
+| What you will learn               | How to build and run a basic rendering program using the Intel&reg; Open VKL API from the Render Kit.
+| Time to complete                  | 5 minutes
+
+## Purpose
+
+This sample program, `vklTutorialCPU`, shows sampling amongst a proceedurally
+generated volume the different volumetric sampling capabilities with Intel&reg;
+Open VKL. Output is written to the console (stdout).
+
+## Key Implementation Details
+
+`vklTutorialCPU` is written in C99 and is constructed to compile with a C++ or C99
+compiler.
+
+## Build and Run
+
+### Windows
+
+1. Run a new **x64 Native Tools Command Prompt for MSVS 2022**.
+
+```
+call <path-to-oneapi-folder>\setvars.bat
+cd <path-to-oneAPI-samples>\RenderingToolkit\GettingStarted\03_openvkl_gsg
+mkdir build
+cd build
+cmake ..
+cmake --build . --config Release
+cd Release
+vklTutorialCPU.exe
+```
+
+Note: MSVS 2019 should use an **x64 Native Tools Command Prompt for MSVS 2019**
+
+2. Review the terminal output (stdout).
+
+
+### Linux and macOS
+
+1. Start a new Terminal session.
+```
+source <path-to-oneapi-folder>/setvars.sh
+cd <path-to-oneAPI-samples>/RenderingToolkit/GettingStarted/03_openvkl_gsg
+mkdir build
+cd build
+cmake -DCMAKE_BUILD_TYPE=Release ..
+cmake --build .
+./vklTutorialCPU
+```
+
+2. Review the terminal output (stdout).
+
+
+## License
+
+This code sample is licensed under the Apache 2.0 license. See
+[LICENSE.txt](LICENSE.txt) for details.
+
+Third party program Licenses can be found here:
+[third-party-programs.txt](https://github.com/oneapi-src/oneAPI-samples/blob/master/third-party-programs.txt).
diff --git a/RenderingToolkit/GettingStarted/03_openvkl_gsg/src/vklTutorial.c b/RenderingToolkit/GettingStarted/03_openvkl_gsg/cpu/src/vklTutorialCPU.c
similarity index 89%
rename from RenderingToolkit/GettingStarted/03_openvkl_gsg/src/vklTutorial.c
rename to RenderingToolkit/GettingStarted/03_openvkl_gsg/cpu/src/vklTutorialCPU.c
index 438ece4447..91212f31ec 100644
--- a/RenderingToolkit/GettingStarted/03_openvkl_gsg/src/vklTutorial.c
+++ b/RenderingToolkit/GettingStarted/03_openvkl_gsg/cpu/src/vklTutorialCPU.c
@@ -1,7 +1,8 @@
-// Copyright 2019-2021 Intel Corporation
+// Copyright 2019 Intel Corporation
 // SPDX-License-Identifier: Apache-2.0
 
 #include <openvkl/openvkl.h>
+#include <openvkl/device/openvkl.h>
 #include <stdio.h>
 
 #if defined(_MSC_VER)
@@ -39,8 +40,8 @@ void demoScalarAPI(VKLDevice device, VKLVolume volume) {
   // sample, gradient (first attribute)
   unsigned int attributeIndex = 0;
   float time = 0.f;
-  float sample = vklComputeSample(sampler, &coord, attributeIndex, time);
-  vkl_vec3f grad = vklComputeGradient(sampler, &coord, attributeIndex, time);
+  float sample = vklComputeSample(&sampler, &coord, attributeIndex, time);
+  vkl_vec3f grad = vklComputeGradient(&sampler, &coord, attributeIndex, time);
   printf("\tsampling and gradient computation (first attribute)\n");
   printf("\t\tsample = %f\n", sample);
   printf("\t\tgrad   = %f %f %f\n\n", grad.x, grad.y, grad.z);
@@ -49,7 +50,7 @@ void demoScalarAPI(VKLDevice device, VKLVolume volume) {
   unsigned int M = 3;
   unsigned int attributeIndices[] = {0, 1, 2};
   float samples[3];
-  vklComputeSampleM(sampler, &coord, samples, M, attributeIndices, time);
+  vklComputeSampleM(&sampler, &coord, samples, M, attributeIndices, time);
   printf("\tsampling (multiple attributes)\n");
   printf("\t\tsamples = %f %f %f\n\n", samples[0], samples[1], samples[2]);
 
@@ -99,12 +100,12 @@ void demoScalarAPI(VKLDevice device, VKLVolume volume) {
 #if defined(_MSC_VER)
     // MSVC does not support variable length arrays, but provides a
     // safer version of alloca.
-    char *buffer = _malloca(vklGetIntervalIteratorSize(intervalContext));
+    char *buffer = _malloca(vklGetIntervalIteratorSize(&intervalContext));
 #else
-    char buffer[vklGetIntervalIteratorSize(intervalContext)];
+    char buffer[vklGetIntervalIteratorSize(&intervalContext)];
 #endif
     VKLIntervalIterator intervalIterator = vklInitIntervalIterator(
-        intervalContext, &rayOrigin, &rayDirection, &rayTRange, time, buffer);
+        &intervalContext, &rayOrigin, &rayDirection, &rayTRange, time, buffer);
 
     printf("\n\tinterval iterator for value ranges {%f %f} {%f %f}\n",
            ranges[0].lower, ranges[0].upper, ranges[1].lower, ranges[1].upper);
@@ -130,12 +131,12 @@ void demoScalarAPI(VKLDevice device, VKLVolume volume) {
 #if defined(_MSC_VER)
     // MSVC does not support variable length arrays, but provides a
     // safer version of alloca.
-    char *buffer = _malloca(vklGetHitIteratorSize(hitContext));
+    char *buffer = _malloca(vklGetHitIteratorSize(&hitContext));
 #else
-    char buffer[vklGetHitIteratorSize(hitContext)];
+    char buffer[vklGetHitIteratorSize(&hitContext)];
 #endif
     VKLHitIterator hitIterator = vklInitHitIterator(
-        hitContext, &rayOrigin, &rayDirection, &rayTRange, time, buffer);
+        &hitContext, &rayOrigin, &rayDirection, &rayTRange, time, buffer);
 
     printf("\thit iterator for values %f %f\n", values[0], values[1]);
 
@@ -182,8 +183,8 @@ void demoVectorAPI(VKLVolume volume) {
   float time4[4] = {0.f};
   float sample4[4];
   vkl_vvec3f4 grad4;
-  vklComputeSample4(valid, sampler, &coord4, sample4, attributeIndex, time4);
-  vklComputeGradient4(valid, sampler, &coord4, &grad4, attributeIndex, time4);
+  vklComputeSample4(valid, &sampler, &coord4, sample4, attributeIndex, time4);
+  vklComputeGradient4(valid, &sampler, &coord4, &grad4, attributeIndex, time4);
 
   printf("\n\tsampling and gradient computation (first attribute)\n");
 
@@ -197,7 +198,7 @@ void demoVectorAPI(VKLVolume volume) {
   unsigned int M = 3;
   unsigned int attributeIndices[] = {0, 1, 2};
   float samples[3 * 4];
-  vklComputeSampleM4(valid, sampler, &coord4, samples, M, attributeIndices,
+  vklComputeSampleM4(valid, &sampler, &coord4, samples, M, attributeIndices,
                      time4);
 
   printf("\n\tsampling (multiple attributes)\n");
@@ -240,8 +241,8 @@ void demoStreamAPI(VKLVolume volume) {
   float time[5] = {0.f};
   float sample[5];
   vkl_vec3f grad[5];
-  vklComputeSampleN(sampler, 5, coord, sample, attributeIndex, time);
-  vklComputeGradientN(sampler, 5, coord, grad, attributeIndex, time);
+  vklComputeSampleN(&sampler, 5, coord, sample, attributeIndex, time);
+  vklComputeGradientN(&sampler, 5, coord, grad, attributeIndex, time);
 
   for (int i = 0; i < 5; i++) {
     printf("\t\tsample[%d] = %f\n", i, sample[i]);
@@ -252,7 +253,7 @@ void demoStreamAPI(VKLVolume volume) {
   unsigned int M = 3;
   unsigned int attributeIndices[] = {0, 1, 2};
   float samples[3 * 5];
-  vklComputeSampleMN(sampler, 5, coord, samples, M, attributeIndices, time);
+  vklComputeSampleMN(&sampler, 5, coord, samples, M, attributeIndices, time);
 
   printf("\n\tsampling (multiple attributes)\n");
 
@@ -271,7 +272,7 @@ void demoStreamAPI(VKLVolume volume) {
 }
 
 int main() {
-  vklLoadModule("cpu_device");
+  vklInit();
 
   VKLDevice device = vklNewDevice("cpu");
   vklCommitDevice(device);
diff --git a/RenderingToolkit/GettingStarted/03_openvkl_gsg/gpu/CMakeLists.txt b/RenderingToolkit/GettingStarted/03_openvkl_gsg/gpu/CMakeLists.txt
new file mode 100644
index 0000000000..189de83e09
--- /dev/null
+++ b/RenderingToolkit/GettingStarted/03_openvkl_gsg/gpu/CMakeLists.txt
@@ -0,0 +1,77 @@
+cmake_minimum_required(VERSION 3.25)
+project(OPENVKL_GPU_GSG LANGUAGES CXX)
+
+set(ONEAPI_ROOT "")
+if(DEFINED ENV{ONEAPI_ROOT})
+  set(ONEAPI_ROOT "$ENV{ONEAPI_ROOT}")
+  message(STATUS "ONEAPI_ROOT FROM ENVIRONMENT: ${ONEAPI_ROOT}")
+else()
+  message(FATAL_ERROR "ONEAPI_ROOT NOT set. Please source environment variables.")
+endif(DEFINED ENV{ONEAPI_ROOT})
+
+#lots of helper variables to get access to other components
+set(OPENVKL_BASE_DIR "")
+set(COMPILER_SYCL_INCLUDE "")
+set(COMPILER_CL_INCLUDE "")
+set(COMPILER_PATH "")
+set(RKCOMMON_BASE_DIR "")
+set(ISPCRT_BASE_DIR "")
+set(TBB_BASE_DIR "")
+set(EMBREE_BASE_DIR "")
+if(EXISTS ${ONEAPI_ROOT}/oneapi-vars.sh OR EXISTS ${ONEAPI_ROOT}/oneapi-bars.bat)
+  set(OPENVKL_BASE_DIR ${ONEAPI_ROOT})
+  set(RKCOMMON_BASE_DIR ${ONEAPI_ROOT})
+  set(ISPCRT_BASE_DIR ${ONEAPI_ROOT})
+  set(TBB_BASE_DIR ${ONEAPI_ROOT})
+  set(EMBREE_BASE_DIR ${ONEAPI_ROOT})
+  if(MSVC)
+
+  set(COMPILER_SYCL_INCLUDE ${ONEAPI_ROOT}/include)
+  set(COMPILER_CL_INCLUDE ${ONEAPI_ROOT}/include/sycl)
+  #  set(ONEAPI_INCLUDE_DIR ${ONEAPI_ROOT}/include 
+  set(COMPILER_PATH ${ONEAPI_ROOT}/bin/compiler)
+  else()
+  set(COMPILER_SYCL_INCLUDE ${ONEAPI_ROOT}/opt/compiler/include/sycl)
+  set(COMPILER_CL_INCLUDE ${ONEAPI_ROOT}/opt/compiler/include/sycl/sycl)
+  #  set(ONEAPI_INCLUDE_DIR ${ONEAPI_ROOT}/include 
+  set(COMPILER_PATH ${ONEAPI_ROOT}/bin/compiler)
+  endif(MSVC)
+
+else()
+  set(OPENVKL_BASE_DIR ${ONEAPI_ROOT}/openvkl/latest)
+  set(RKCOMMON_BASE_DIR ${ONEAPI_ROOT}/rkcommon/latest)
+  set(ISPCRT_BASE_DIR ${ONEAPI_ROOT}/ispc/latest)
+  set(TBB_BASE_DIR ${ONEAPI_ROOT}/tbb/latest)
+  set(EMBREE_BASE_DIR ${ONEAPI_ROOT}/embree/latest)
+  if(MSVC)
+  set(COMPILER_SYCL_INCLUDE ${ONEAPI_ROOT}/compiler/latest/include)
+  set(COMPILER_CL_INCLUDE ${ONEAPI_ROOT}/compiler/latest/include/sycl)
+  set(COMPILER_PATH ${ONEAPI_ROOT}/compiler/latest/bin/compiler)
+  else()
+  set(COMPILER_SYCL_INCLUDE ${ONEAPI_ROOT}/compiler/latest/linux/include)
+  set(COMPILER_CL_INCLUDE ${ONEAPI_ROOT}/compiler/latest/linux/include/sycl)
+  set(COMPILER_PATH ${ONEAPI_ROOT}/compiler/latest/bin/compiler)
+  endif(MSVC)
+endif(EXISTS ${ONEAPI_ROOT}/oneapi-vars.sh OR EXISTS ${ONEAPI_ROOT}/oneapi-bars.bat)
+
+if(NOT DEFINED CMAKE_CXX_COMPILER)
+	message(FATAL_ERROR "CMAKE_CXX_COMPILER not set use icpx")
+endif(NOT DEFINED CMAKE_CXX_COMPILER)
+
+set(OPENVKL_INCLUDE_DIR ${OPENVKL_BASE_DIR}/include)
+#target headers
+set(OPENVKL_GPU_INCLUDE_DIR ${OPENVKL_INCLUDE_DIR}/openvkl/devices/gpu)
+
+find_package(openvkl 2.0 REQUIRED PATHS ${OPENVKL_BASE_DIR})
+
+find_package(rkcommon REQUIRED PATHS ${RKCOMMON_BASE_DIR})
+
+find_package(TBB REQUIRED PATHS ${TBB_BASE_DIR}/lib/cmake/tbb)
+
+find_package(embree 4.0 REQUIRED PATHS ${EMBREE_BASE_DIR})
+
+link_directories(${OPENVKL_BASE_DIR}/lib ${RKCOMMON_BASE_DIR}/lib ${ISPCRT_BASE_DIR}/lib)
+
+add_executable(vklTutorialGPU src/vklTutorialGPU.cpp)
+target_link_libraries(vklTutorialGPU PRIVATE openvkl::openvkl openvkl::openvkl_module_gpu_device rkcommon ispcrt tbb embree)
+install(TARGETS vklTutorialGPU RUNTIME DESTINATION ${CMAKE_INSTALL_BINDIR})
diff --git a/RenderingToolkit/GettingStarted/03_openvkl_gsg/gpu/README.md b/RenderingToolkit/GettingStarted/03_openvkl_gsg/gpu/README.md
new file mode 100644
index 0000000000..876cf02976
--- /dev/null
+++ b/RenderingToolkit/GettingStarted/03_openvkl_gsg/gpu/README.md
@@ -0,0 +1,76 @@
+# Getting Started Sample for Intel&reg; Rendering Toolkit (Render Kit): Intel&reg; Open Volume Kernel Library (Intel&reg; Open VKL) for GPU
+
+Intel&reg; Open Volume Kernel Library (Intel&reg; Open VKL) is a collection of
+high-performance volume computation kernels. Improve performance of volume
+rendering applications by using performance optimized volume traversal and
+sampling functionality for a variety of data formats.
+
+| Minimum Requirements              | Description
+|:---                               |:---
+| OS                                | Linux* Ubuntu* 22.04 <br>CentOS 8 (or compatible) <br> Windows* 10 or 11
+| Hardware                          | Intel&reg; Arc Graphics (Xe-HPG architecture, DG2-128, DG2-512) or higher
+| Compiler Toolchain                | Windows OS: MSVS 2022 (or 2019) installed with Windows* SDK and CMake*; All Platforms: Intel&reg; oneAPI DPC++ Compiler from the Intel&reg; oneAPI Base Toolkit
+| SYCL Compiler                     | oneAPI DPC++ 2024.0.0 compiler or higher
+| Libraries                         | Install Intel&reg; Rendering Toolkit (Render Kit), including Intel&reg; Embree and Intel&reg; Open VKL
+| Knowledge                         | First, build and run the [CPU](../cpu) get started program `vklTutorialCPU`
+
+| Objective                         | Description
+|:---                               |:---
+| What you will learn               | How to build and run a basic rendering program using the Intel&reg; Open VKL API from the Render Kit targeting GPU
+| Time to complete                  | 5 minutes
+
+## Purpose
+
+This sample program, `vklTutorialGPU`, shows sampling amongst a proceedurally
+generated volume the different volumetric sampling capabilities with Intel&reg;
+Open VKL. Output is written to the console (stdout).
+
+## Key Implementation Details
+
+`vklTutorialGPU` is written for C++17 with SYCL. It is constructed to build with the Intel&reg; oneAPI DPC++ Compiler. On Windows* OS it must also build with MSVS system libraries.
+
+## Build and Run
+
+### Windows
+
+1. Run a new **x64 Native Tools Command Prompt for MSVS 2022**.
+
+```
+call <path-to-oneapi-folder>\setvars.bat
+cd <path-to-oneAPI-samples>\RenderingToolkit\GettingStarted\03_openvkl_gsg\gpu
+mkdir build
+cd build
+cmake -G"NMake Makefiles" -DCMAKE_BUILD_TYPE=Release -DCMAKE_CXX_COMPILER=icx-cl ..
+cmake --build . --config Release
+cd Release
+.\vklTutorialGPU.exe
+```
+
+Note: MSVS 2019 users should use a **x64 Native Tools Command Prompt for MSVS 2019**
+
+2. Review the terminal output (stdout).
+
+
+### Linux
+
+1. Start a new Terminal session.
+```
+source <path-to-oneapi-folder>/setvars.sh
+cd <path-to-oneAPI-samples>/RenderingToolkit/GettingStarted/03_openvkl_gsg/gpu
+mkdir build
+cd build
+cmake -DCMAKE_BUILD_TYPE=Release -DCMAKE_CXX_COMPILER=icpx ..
+cmake --build .
+./vklTutorialGPU
+```
+
+2. Review the terminal output (stdout).
+
+
+## License
+
+This code sample is licensed under the Apache 2.0 license. See
+[LICENSE.txt](LICENSE.txt) for details.
+
+Third party program Licenses can be found here:
+[third-party-programs.txt](https://github.com/oneapi-src/oneAPI-samples/blob/master/third-party-programs.txt).
diff --git a/RenderingToolkit/GettingStarted/03_openvkl_gsg/gpu/src/vklTutorialGPU.cpp b/RenderingToolkit/GettingStarted/03_openvkl_gsg/gpu/src/vklTutorialGPU.cpp
new file mode 100644
index 0000000000..25583d00a0
--- /dev/null
+++ b/RenderingToolkit/GettingStarted/03_openvkl_gsg/gpu/src/vklTutorialGPU.cpp
@@ -0,0 +1,379 @@
+// Copyright 2022 Intel Corporation
+// SPDX-License-Identifier: Apache-2.0
+
+
+#include <openvkl/openvkl.h>
+#include <openvkl/device/openvkl.h>
+
+#include <iomanip>
+#include <iostream>
+
+// setup specialization constant for feature flags
+static_assert(std::is_trivially_copyable<VKLFeatureFlags>::value);
+
+constexpr static sycl::specialization_id<VKLFeatureFlags> samplerSpecId{
+    VKL_FEATURE_FLAGS_DEFAULT};
+
+void demoGpuAPI(sycl::queue &syclQueue, VKLDevice device, VKLVolume volume) {
+  std::cout << "demo of GPU API" << std::endl;
+
+  std::cout << std::fixed << std::setprecision(6);
+
+  VKLSampler sampler = vklNewSampler(volume);
+  vklCommit(sampler);
+
+  // feature flags improve performance on GPU, as well as JIT times
+  const VKLFeatureFlags requiredFeatures = vklGetFeatureFlags(sampler);
+
+  // bounding box
+  vkl_box3f bbox = vklGetBoundingBox(volume);
+  std::cout << "\tbounding box" << std::endl;
+  std::cout << "\t\tlower = " << bbox.lower.x << " " << bbox.lower.y << " "
+            << bbox.lower.z << std::endl;
+  std::cout << "\t\tupper = " << bbox.upper.x << " " << bbox.upper.y << " "
+            << bbox.upper.z << std::endl
+            << std::endl;
+
+  // number of attributes
+  unsigned int numAttributes = vklGetNumAttributes(volume);
+  std::cout << "\tnum attributes = " << numAttributes << std::endl;
+
+  // value range for all attributes
+  for (unsigned int i = 0; i < numAttributes; i++) {
+    vkl_range1f valueRange = vklGetValueRange(volume, i);
+    std::cout << "\tvalue range (attribute " << i << ") = (" << valueRange.lower
+              << " " << valueRange.upper << ")" << std::endl;
+  }
+
+  std::cout << std::endl << "\tsampling" << std::endl;
+
+  // coordinate for sampling / gradients
+  vkl_vec3f coord = {1.f, 2.f, 3.f};
+  std::cout << "\n\tcoord = " << coord.x << " " << coord.y << " " << coord.z
+            << std::endl
+            << std::endl;
+
+  // sample, gradient (first attribute)
+  const unsigned int attributeIndex = 0;
+  const float time = 0.f;
+
+  // USM shared allocations, required when we want to pass results back from GPU
+  float *sample = sycl::malloc_shared<float>(1, syclQueue);
+  vkl_vec3f *grad = sycl::malloc_shared<vkl_vec3f>(1, syclQueue);
+
+  syclQueue
+      .submit([=](sycl::handler &cgh) {
+        cgh.set_specialization_constant<samplerSpecId>(requiredFeatures);
+
+        cgh.single_task([=](sycl::kernel_handler kh) {
+          const VKLFeatureFlags featureFlags =
+              kh.get_specialization_constant<samplerSpecId>();
+
+          *sample = vklComputeSample(&sampler, &coord, attributeIndex, time,
+                                     featureFlags);
+          *grad = vklComputeGradient(&sampler, &coord, attributeIndex, time,
+                                     featureFlags);
+        });
+      })
+      .wait();
+
+  std::cout << "\tsampling and gradient computation (first attribute)"
+            << std::endl;
+  std::cout << "\t\tsample = " << *sample << std::endl;
+  std::cout << "\t\tgrad   = " << grad->x << " " << grad->y << " " << grad->z
+            << std::endl
+            << std::endl;
+
+  sycl::free(sample, syclQueue);
+  sycl::free(grad, syclQueue);
+
+  // sample (multiple attributes)
+  const unsigned int M = 3;
+  const unsigned int attributeIndices[] = {0, 1, 2};
+
+  float *samples = sycl::malloc_shared<float>(M, syclQueue);
+
+  syclQueue
+      .submit([=](sycl::handler &cgh) {
+        cgh.set_specialization_constant<samplerSpecId>(requiredFeatures);
+
+        cgh.single_task([=](sycl::kernel_handler kh) {
+          const VKLFeatureFlags featureFlags =
+              kh.get_specialization_constant<samplerSpecId>();
+
+          vklComputeSampleM(&sampler, &coord, samples, M, attributeIndices,
+                            time, featureFlags);
+        });
+      })
+      .wait();
+
+  std::cout << "\tsampling (multiple attributes)" << std::endl;
+  std::cout << "\t\tsamples = " << samples[0] << " " << samples[1] << " "
+            << samples[2] << std::endl;
+
+  sycl::free(samples, syclQueue);
+
+  // interval iterator context setup
+  std::cout << std::endl << "\tinterval iteration" << std::endl << std::endl;
+
+  std::vector<vkl_range1f> ranges{{10, 20}, {50, 75}};
+  VKLData rangesData =
+      vklNewData(device, ranges.size(), VKL_BOX1F, ranges.data());
+
+  VKLIntervalIteratorContext intervalContext =
+      vklNewIntervalIteratorContext(sampler);
+
+  vklSetInt(intervalContext, "attributeIndex", 0);
+
+  vklSetData(intervalContext, "valueRanges", rangesData);
+  vklRelease(rangesData);
+
+  vklCommit(intervalContext);
+
+  // ray definition for iterators
+  vkl_vec3f rayOrigin{0.f, 1.f, 1.f};
+  vkl_vec3f rayDirection{1.f, 0.f, 0.f};
+  vkl_range1f rayTRange{0.f, 200.f};
+  std::cout << "\trayOrigin = " << rayOrigin.x << " " << rayOrigin.y << " "
+            << rayOrigin.z << std::endl;
+  std::cout << "\trayDirection = " << rayDirection.x << " " << rayDirection.y
+            << " " << rayDirection.z << std::endl;
+  std::cout << "\trayTRange = " << rayTRange.lower << " " << rayTRange.upper
+            << std::endl
+            << std::endl;
+
+  // interval iteration
+  char *iteratorBuffer = sycl::malloc_device<char>(
+      vklGetIntervalIteratorSize(&intervalContext), syclQueue);
+
+  int *numIntervals = sycl::malloc_shared<int>(1, syclQueue);
+  *numIntervals = 0;
+
+  const size_t maxNumIntervals = 999;
+
+  VKLInterval *intervalsBuffer =
+      sycl::malloc_shared<VKLInterval>(maxNumIntervals, syclQueue);
+  memset(intervalsBuffer, 0, maxNumIntervals * sizeof(VKLInterval));
+
+  std::cout << "\tinterval iterator for value ranges";
+
+  for (const auto &r : ranges) {
+    std::cout << " {" << r.lower << " " << r.upper << "}";
+  }
+  std::cout << std::endl << std::endl;
+
+  syclQueue
+      .submit([=](sycl::handler &cgh) {
+        cgh.set_specialization_constant<samplerSpecId>(requiredFeatures);
+
+        cgh.single_task([=](sycl::kernel_handler kh) {
+          const VKLFeatureFlags featureFlags =
+              kh.get_specialization_constant<samplerSpecId>();
+
+          VKLIntervalIterator intervalIterator = vklInitIntervalIterator(
+              &intervalContext, &rayOrigin, &rayDirection, &rayTRange, time,
+              (void *)iteratorBuffer, featureFlags);
+
+          for (;;) {
+            VKLInterval interval;
+            int result =
+                vklIterateInterval(intervalIterator, &interval, featureFlags);
+            if (!result) {
+              break;
+            }
+            intervalsBuffer[*numIntervals] = interval;
+
+            *numIntervals = *numIntervals + 1;
+            if (*numIntervals >= maxNumIntervals) break;
+          }
+        });
+      })
+      .wait();
+
+  for (int i = 0; i < *numIntervals; ++i) {
+    std::cout << "\t\ttRange (" << intervalsBuffer[i].tRange.lower << " "
+              << intervalsBuffer[i].tRange.upper << ")" << std::endl;
+    std::cout << "\t\tvalueRange (" << intervalsBuffer[i].valueRange.lower
+              << " " << intervalsBuffer[i].valueRange.upper << ")" << std::endl;
+    std::cout << "\t\tnominalDeltaT " << intervalsBuffer[i].nominalDeltaT
+              << std::endl
+              << std::endl;
+  }
+
+  sycl::free(iteratorBuffer, syclQueue);
+  sycl::free(numIntervals, syclQueue);
+  sycl::free(intervalsBuffer, syclQueue);
+
+  vklRelease(intervalContext);
+
+  // hit iteration
+  std::cout << std::endl << "\thit iteration" << std::endl << std::endl;
+
+  // hit iterator context setup
+  float values[2] = {32.f, 96.f};
+  int num_values = 2;
+  VKLData valuesData = vklNewData(device, num_values, VKL_FLOAT, values);
+
+  VKLHitIteratorContext hitContext = vklNewHitIteratorContext(sampler);
+
+  vklSetInt(hitContext, "attributeIndex", 0);
+
+  vklSetData(hitContext, "values", valuesData);
+  vklRelease(valuesData);
+
+  vklCommit(hitContext);
+
+  // ray definition for iterators
+  // see rayOrigin, Direction and TRange above
+
+  char *hitIteratorBuffer =
+      sycl::malloc_device<char>(vklGetHitIteratorSize(&hitContext), syclQueue);
+
+  int *numHits = sycl::malloc_shared<int>(1, syclQueue);
+  *numHits = 0;
+
+  const size_t maxNumHits = 999;
+
+  VKLHit *hitBuffer = sycl::malloc_shared<VKLHit>(maxNumHits, syclQueue);
+  memset(hitBuffer, 0, maxNumHits * sizeof(VKLHit));
+
+  std::cout << "\thit iterator for values";
+
+  for (const auto &r : values) {
+    std::cout << " " << r << " ";
+  }
+  std::cout << std::endl << std::endl;
+
+  syclQueue
+      .submit([=](sycl::handler &cgh) {
+        cgh.set_specialization_constant<samplerSpecId>(requiredFeatures);
+
+        cgh.single_task([=](sycl::kernel_handler kh) {
+          const VKLFeatureFlags featureFlags =
+              kh.get_specialization_constant<samplerSpecId>();
+
+          VKLHitIterator hitIterator = vklInitHitIterator(
+              &hitContext, &rayOrigin, &rayDirection, &rayTRange, time,
+              (void *)hitIteratorBuffer, featureFlags);
+
+          for (;;) {
+            VKLHit hit;
+            int result = vklIterateHit(hitIterator, &hit, featureFlags);
+            if (!result) {
+              break;
+            }
+            hitBuffer[*numHits] = hit;
+
+            *numHits = *numHits + 1;
+            if (*numHits >= maxNumHits) break;
+          }
+        });
+      })
+      .wait();
+
+  for (int i = 0; i < *numHits; ++i) {
+    std::cout << "\t\tt " << hitBuffer[i].t << std::endl;
+    std::cout << "\t\tsample " << hitBuffer[i].sample << std::endl;
+    std::cout << "\t\tepsilon " << hitBuffer[i].epsilon << std::endl
+              << std::endl;
+  }
+
+  sycl::free(hitIteratorBuffer, syclQueue);
+  sycl::free(numHits, syclQueue);
+  sycl::free(hitBuffer, syclQueue);
+
+  vklRelease(hitContext);
+
+  vklRelease(sampler);
+}
+
+int main() {
+  auto IntelGPUDeviceSelector = [](const sycl::device &device) {
+    using namespace sycl::info;
+    const std::string deviceName = device.get_info<device::name>();
+    bool match = device.is_gpu() &&
+                 device.get_info<sycl::info::device::vendor_id>() == 0x8086 &&
+                 device.get_backend() == sycl::backend::ext_oneapi_level_zero;
+    return match ? 1 : -1;
+  };
+
+  sycl::queue syclQueue(IntelGPUDeviceSelector);
+
+  sycl::context syclContext = syclQueue.get_context();
+
+  std::cout << "Target SYCL device: "
+            << syclQueue.get_device().get_info<sycl::info::device::name>()
+            << std::endl
+            << std::endl;
+
+  vklInit();
+
+  VKLDevice device = vklNewDevice("gpu");
+  vklDeviceSetVoidPtr(device, "syclContext", static_cast<void *>(&syclContext));
+  vklCommitDevice(device);
+
+  const int dimensions[] = {128, 128, 128};
+
+  const int numVoxels = dimensions[0] * dimensions[1] * dimensions[2];
+
+  const int numAttributes = 3;
+
+  VKLVolume volume = vklNewVolume(device, "structuredRegular");
+  vklSetVec3i(volume, "dimensions", dimensions[0], dimensions[1],
+              dimensions[2]);
+  vklSetVec3f(volume, "gridOrigin", 0, 0, 0);
+  vklSetVec3f(volume, "gridSpacing", 1, 1, 1);
+
+  std::vector<float> voxels(numVoxels);
+
+  // volume attribute 0: x-grad
+  for (int k = 0; k < dimensions[2]; k++)
+    for (int j = 0; j < dimensions[1]; j++)
+      for (int i = 0; i < dimensions[0]; i++)
+        voxels[k * dimensions[0] * dimensions[1] + j * dimensions[2] + i] =
+            (float)i;
+
+  VKLData data0 = vklNewData(device, numVoxels, VKL_FLOAT, voxels.data());
+
+  // volume attribute 1: y-grad
+  for (int k = 0; k < dimensions[2]; k++)
+    for (int j = 0; j < dimensions[1]; j++)
+      for (int i = 0; i < dimensions[0]; i++)
+        voxels[k * dimensions[0] * dimensions[1] + j * dimensions[2] + i] =
+            (float)j;
+
+  VKLData data1 = vklNewData(device, numVoxels, VKL_FLOAT, voxels.data());
+
+  // volume attribute 2: z-grad
+  for (int k = 0; k < dimensions[2]; k++)
+    for (int j = 0; j < dimensions[1]; j++)
+      for (int i = 0; i < dimensions[0]; i++)
+        voxels[k * dimensions[0] * dimensions[1] + j * dimensions[2] + i] =
+            (float)k;
+
+  VKLData data2 = vklNewData(device, numVoxels, VKL_FLOAT, voxels.data());
+
+  VKLData attributes[] = {data0, data1, data2};
+
+  VKLData attributesData =
+      vklNewData(device, numAttributes, VKL_DATA, attributes);
+
+  vklRelease(data0);
+  vklRelease(data1);
+  vklRelease(data2);
+
+  vklSetData(volume, "data", attributesData);
+  vklRelease(attributesData);
+
+  vklCommit(volume);
+
+  demoGpuAPI(syclQueue, device, volume);
+
+  vklRelease(volume);
+
+  vklReleaseDevice(device);
+
+  std::cout << "complete." << std::endl;
+
+  return 0;
+}
diff --git a/RenderingToolkit/GettingStarted/03_openvkl_gsg/sample.json b/RenderingToolkit/GettingStarted/03_openvkl_gsg/sample.json
index bff4cae380..f93a4b22a1 100644
--- a/RenderingToolkit/GettingStarted/03_openvkl_gsg/sample.json
+++ b/RenderingToolkit/GettingStarted/03_openvkl_gsg/sample.json
@@ -4,20 +4,27 @@
   "categories": ["Toolkit/oneAPI Libraries/Open VKL"],
   "description": "This introductory hello rendering toolkit sample program demonstrates how to sample into volumes with Intel Open VKL",
   "builder": ["cli"],
-  "languages": [{"c":{}}],
+  "languages": [{"cpp":{}}],
   "dependencies": ["tbb","openvkl","embree","rkcommon"],
   "os":["linux", "windows", "darwin"],
-  "targetDevice": ["CPU"],
+  "targetDevice": ["CPU","GPU"],
   "ciTests": {
         "linux": [
                 {
                 "id": "Intel_OpenVKL_vklTutorial_lin",
                 "steps": [
+                    "cd cpu",
                     "mkdir build",
                     "cd build",
                     "cmake ..",
                     "cmake --build . ",
-                    "./vklTutorial"
+                    "./vklTutorialCPU",
+		    "cd ../gpu",
+		    "mkdir build",
+		    "cd build",
+		    "cmake -DCMAKE_CXX_COMPILER=icpx ..",
+		    "cmake --build .",
+		    "./vklTutorialGPU"
                  ]
                  }
         ],
@@ -25,12 +32,20 @@
                 {
                 "id": "Intel_OpenVKL_vklTutorial_win",
                 "steps":[
+                    "cd cpu",
                     "mkdir build",
                     "cd build",
                     "cmake ..",
                     "cmake --build . --config Release",
                     "cd Release",
-                    ".\\vklTutorial.exe"
+                    ".\\vklTutorialCPU.exe",
+		    "cd ..",
+		    "cd gpu",
+		    "mkdir build",
+		    "cd build",
+		    "cmake -G\"NMake Makefiles\" -DCMAKE_CXX_COMPILER=icx-cl ..",
+		    "cmake --build . --config Release",
+		    ".\\vklTutorialGPU.exe"
                 ]
 
                 }
@@ -39,11 +54,13 @@
                 {
                 "id": "Intel_OpenVKL_vklTutorial_mac",
                 "steps": [
+                    "cd cpu",
                     "mkdir build",
                     "cd build",
                     "cmake ..",
                     "cmake --build . ",
-                    "export DYLD_LIBRARY_PATH=${ONEAPI_ROOT}/openvkl/latest/lib:${ONEAPI_ROOT}/rkcommon/latest/lib:${ONEAPI_ROOT}/tbb/latest/lib:${ONEAPI_ROOT}/embree/latest/lib && ./vklTutorial"
+                    "export DYLD_LIBRARY_PATH=${ONEAPI_ROOT}/lib && ./vklTutorialCPU",
+                    "export DYLD_LIBRARY_PATH=${ONEAPI_ROOT}/openvkl/latest/lib:${ONEAPI_ROOT}/rkcommon/latest/lib:${ONEAPI_ROOT}/tbb/latest/lib:${ONEAPI_ROOT}/embree/latest/lib && ./vklTutorialCPU"
                  ]
                  }
         ]