Add kernel launch support

oneapi-rs-sys/build.rs

@@ -20,6 +20,8 @@ fn main() {
         "src/queue-sys.rs",
         "src/usm-sys.rs",
         "src/event-sys.rs",
+        "src/context-sys.rs",
+        "src/kernel-bundle-sys.rs",
     ];
 
     let cpp_sources = [
@@ -28,6 +30,8 @@ fn main() {
         "src/queue.cpp",
         "src/usm.cpp",
         "src/event.cpp",
+        "src/context.cpp",
+        "src/kernel-bundle.cpp",
     ];
 
     let cpp_headers = [
@@ -37,6 +41,8 @@ fn main() {
         "include/queue.hpp",
         "include/usm.hpp",
         "include/event.hpp",
+        "include/context.hpp",
+        "include/kernel-bundle.hpp",
     ];
 
     cxx_build::bridges(&rust_sources)
@@ -47,6 +53,7 @@ fn main() {
         .compile("oneapi-shim");
 
     println!("cargo::rustc-link-lib=sycl");
+    println!("cargo::rustc-link-lib=ze_loader");
     println!("cargo::rustc-link-lib=intlc");
 
     for source in cpp_sources {

oneapi-rs-sys/include/context.hpp

@@ -0,0 +1,24 @@
+//
+// Copyright (C) 2026 Intel Corporation
+//
+// Under the MIT License or the Apache License v2.0.
+// See LICENSE-MIT and LICENSE-APACHE for license information.
+// SPDX-License-Identifier: MIT OR Apache-2.0
+//
+
+#pragma once
+
+#include <memory>
+
+#include <sycl/sycl.hpp>
+
+#include "oneapi-rs-sys/include/types.hpp"
+#include "rust/cxx.h"
+
+namespace sycl_shims {
+struct DevicePtr;
+} // namespace sycl_shims
+
+namespace sycl_shims::context {
+std::unique_ptr<Context> new_context(rust::Vec<DevicePtr>);
+} // namespace sycl_shims::context

oneapi-rs-sys/include/device.hpp

@@ -25,4 +25,5 @@ DeviceType get_device_type(Device const &);
 rust::String get_version(Device const &);
 rust::String get_name(Device const &);
 std::unique_ptr<Platform> get_platform(Device const &);
+std::unique_ptr<Device> clone(Device const &);
 } // namespace sycl_shims::device

oneapi-rs-sys/include/kernel-bundle.hpp

@@ -0,0 +1,25 @@
+//
+// Copyright (C) 2026 Intel Corporation
+//
+// Under the MIT License or the Apache License v2.0.
+// See LICENSE-MIT and LICENSE-APACHE for license information.
+// SPDX-License-Identifier: MIT OR Apache-2.0
+//
+
+#pragma once
+
+#include <memory>
+
+#include <sycl/sycl.hpp>
+
+#include "oneapi-rs-sys/include/types.hpp"
+#include "rust/cxx.h"
+
+namespace sycl_shims::kernel_bundle {
+std::unique_ptr<SourceKernelBundle>
+create_kernel_bundle_from_source(Context const &ctxt, rust::Str source);
+std::unique_ptr<ExecutableKernelBundle>
+build(std::unique_ptr<SourceKernelBundle> &source);
+std::unique_ptr<Kernel> get_kernel(std::unique_ptr<ExecutableKernelBundle> &,
+                                   rust::Str);
+} // namespace sycl_shims::kernel_bundle

oneapi-rs-sys/include/queue.hpp

@@ -15,16 +15,35 @@
 
 namespace sycl_shims {
 struct EventPtr;
+struct Range1;
+struct Range2;
+struct Range3;
 } // namespace sycl_shims
 
 namespace sycl_shims::queue {
 std::unique_ptr<Queue> new_queue();
 std::unique_ptr<Queue> new_queue_immediate();
 std::unique_ptr<Queue> new_queue_from_device(Device const &);
+std::unique_ptr<Context> get_context(Queue const &);
 std::unique_ptr<Queue> clone(Queue const &);
 std::unique_ptr<Event> memset(std::unique_ptr<Queue> &, std::uint8_t *ptr,
                               int value, std::size_t num_bytes,
                               rust::Vec<EventPtr>);
 std::unique_ptr<Event> barrier(std::unique_ptr<Queue> &, rust::Vec<EventPtr>);
 void wait(std::unique_ptr<Queue> &);
+
+std::unique_ptr<Event>
+launch_1d(std::unique_ptr<Queue> &, Range1 global_size, Range1 local_size,
+          Kernel const &,
+          rust::Slice<rust::slice<std::uint8_t const> const> args);
+
+std::unique_ptr<Event>
+launch_2d(std::unique_ptr<Queue> &, Range2 global_size, Range2 local_size,
+          Kernel const &,
+          rust::Slice<rust::slice<std::uint8_t const> const> args);
+
+std::unique_ptr<Event>
+launch_3d(std::unique_ptr<Queue> &, Range3 global_size, Range3 local_size,
+          Kernel const &,
+          rust::Slice<rust::slice<std::uint8_t const> const> args);
 } // namespace sycl_shims::queue

oneapi-rs-sys/include/types.hpp

@@ -15,4 +15,10 @@ using Device = sycl::device;
 using Platform = sycl::platform;
 using Queue = sycl::queue;
 using Event = sycl::event;
+using Context = sycl::context;
+using Kernel = sycl::kernel;
+using SourceKernelBundle =
+    sycl::kernel_bundle<sycl::bundle_state::ext_oneapi_source>;
+using ExecutableKernelBundle =
+    sycl::kernel_bundle<sycl::bundle_state::executable>;
 } // namespace sycl_shims

oneapi-rs-sys/src/context-sys.rs

@@ -0,0 +1,25 @@
+//
+// Copyright (C) 2026 Intel Corporation
+//
+// Under the MIT License or the Apache License v2.0.
+// See LICENSE-MIT and LICENSE-APACHE for license information.
+// SPDX-License-Identifier: MIT OR Apache-2.0
+//
+
+#[cxx::bridge(namespace = "sycl_shims::context")]
+pub mod ffi {
+    #[namespace = "sycl_shims"]
+    extern "C++" {
+        include!("oneapi-rs-sys/src/types-sys.rs.h");
+        type DevicePtr = crate::types::ffi::DevicePtr;
+    }
+
+    unsafe extern "C++" {
+        include!("oneapi-rs-sys/include/context.hpp");
+
+        #[namespace = "sycl_shims"]
+        type Context = crate::types::ffi::Context;
+
+        fn new_context(devices: Vec<DevicePtr>) -> UniquePtr<Context>;
+    }
+}

oneapi-rs-sys/src/context.cpp

@@ -0,0 +1,19 @@
+//
+// Copyright (C) 2026 Intel Corporation
+//
+// Under the MIT License or the Apache License v2.0.
+// See LICENSE-MIT and LICENSE-APACHE for license information.
+// SPDX-License-Identifier: MIT OR Apache-2.0
+//
+
+#include "oneapi-rs-sys/include/context.hpp"
+#include "oneapi-rs-sys/src/context-sys.rs.h"
+
+namespace sycl_shims::context {
+std::unique_ptr<Context> new_context(rust::Vec<DevicePtr> devices) {
+  std::vector<sycl::device> raw_devices;
+  for (auto &&d : devices)
+    raw_devices.push_back(std::move(*d.ptr.release()));
+  return std::make_unique<Context>(raw_devices);
+}
+} // namespace sycl_shims::context

oneapi-rs-sys/src/device-sys.rs

@@ -31,5 +31,6 @@ pub mod ffi {
         fn get_device_type(device: &Device) -> DeviceType;
         fn get_version(device: &Device) -> String;
         fn get_name(device: &Device) -> String;
+        fn clone(device: &Device) -> UniquePtr<Device>;
     }
 }

oneapi-rs-sys/src/device.cpp

@@ -53,4 +53,8 @@ rust::String get_name(Device const &device) {
 std::unique_ptr<Platform> get_platform(Device const &device) {
   return std::make_unique<Platform>(device.get_platform());
 }
+
+std::unique_ptr<Device> clone(Device const &device) {
+  return std::make_unique<Device>(sycl::device(device));
+}
 } // namespace sycl_shims::device

oneapi-rs-sys/src/kernel-bundle-sys.rs

@@ -0,0 +1,36 @@
+//
+// Copyright (C) 2026 Intel Corporation
+//
+// Under the MIT License or the Apache License v2.0.
+// See LICENSE-MIT and LICENSE-APACHE for license information.
+// SPDX-License-Identifier: MIT OR Apache-2.0
+//
+
+#[cxx::bridge(namespace = "sycl_shims::kernel_bundle")]
+pub mod ffi {
+    unsafe extern "C++" {
+        include!("oneapi-rs-sys/include/kernel-bundle.hpp");
+
+        #[namespace = "sycl_shims"]
+        type Context = crate::types::ffi::Context;
+
+        #[namespace = "sycl_shims"]
+        type SourceKernelBundle = crate::types::ffi::SourceKernelBundle;
+
+        #[namespace = "sycl_shims"]
+        type ExecutableKernelBundle = crate::types::ffi::ExecutableKernelBundle;
+
+        #[namespace = "sycl_shims"]
+        type Kernel = crate::types::ffi::Kernel;
+
+        fn create_kernel_bundle_from_source(
+            ctxt: &Context,
+            source: &str,
+        ) -> UniquePtr<SourceKernelBundle>;
+        fn build(source: &mut UniquePtr<SourceKernelBundle>) -> UniquePtr<ExecutableKernelBundle>;
+        fn get_kernel(
+            bundle: &mut UniquePtr<ExecutableKernelBundle>,
+            name: &str,
+        ) -> UniquePtr<Kernel>;
+    }
+}

oneapi-rs-sys/src/kernel-bundle.cpp

@@ -0,0 +1,32 @@
+//
+// Copyright (C) 2026 Intel Corporation
+//
+// Under the MIT License or the Apache License v2.0.
+// See LICENSE-MIT and LICENSE-APACHE for license information.
+// SPDX-License-Identifier: MIT OR Apache-2.0
+//
+
+#include "oneapi-rs-sys/include/kernel-bundle.hpp"
+#include "oneapi-rs-sys/src/kernel-bundle-sys.rs.h"
+
+namespace syclexp = sycl::ext::oneapi::experimental;
+
+namespace sycl_shims::kernel_bundle {
+std::unique_ptr<SourceKernelBundle>
+create_kernel_bundle_from_source(Context const &ctxt, rust::Str source) {
+  return std::make_unique<SourceKernelBundle>(
+      syclexp::create_kernel_bundle_from_source(
+          ctxt, syclexp::source_language::sycl, std::string(source)));
+}
+
+std::unique_ptr<ExecutableKernelBundle>
+build(std::unique_ptr<SourceKernelBundle> &source) {
+  return std::make_unique<ExecutableKernelBundle>(syclexp::build(*source));
+}
+
+std::unique_ptr<Kernel>
+get_kernel(std::unique_ptr<ExecutableKernelBundle> &bundle, rust::Str name) {
+  return std::make_unique<Kernel>(
+      bundle->ext_oneapi_get_kernel(static_cast<std::string const &>(name)));
+}
+} // namespace sycl_shims::kernel_bundle

oneapi-rs-sys/src/lib.rs

@@ -23,3 +23,9 @@ pub mod usm;
 
 #[path = "event-sys.rs"]
 pub mod event;
+
+#[path = "context-sys.rs"]
+pub mod context;
+
+#[path = "kernel-bundle-sys.rs"]
+pub mod kernel_bundle;

oneapi-rs-sys/src/queue-sys.rs

@@ -12,6 +12,9 @@ pub mod ffi {
     extern "C++" {
         include!("oneapi-rs-sys/src/types-sys.rs.h");
         type EventPtr = crate::types::ffi::EventPtr;
+        type Range1 = crate::types::ffi::Range1;
+        type Range2 = crate::types::ffi::Range2;
+        type Range3 = crate::types::ffi::Range3;
     }
 
     unsafe extern "C++" {
@@ -22,20 +25,51 @@ pub mod ffi {
         #[namespace = "sycl_shims"]
         type Device = crate::types::ffi::Device;
         #[namespace = "sycl_shims"]
+        type Context = crate::types::ffi::Context;
+        #[namespace = "sycl_shims"]
         type Event = crate::types::ffi::Event;
+        #[namespace = "sycl_shims"]
+        type Kernel = crate::types::ffi::Kernel;
 
         fn new_queue() -> UniquePtr<Queue>;
         fn new_queue_immediate() -> UniquePtr<Queue>;
         fn new_queue_from_device(device: &Device) -> UniquePtr<Queue>;
+        fn get_context(queue: &Queue) -> UniquePtr<Context>;
         fn clone(queue: &Queue) -> UniquePtr<Queue>;
+
         unsafe fn memset(
             queue: &mut UniquePtr<Queue>,
             ptr: *mut u8,
             value: i32,
             num_bytes: usize,
             dep_events: Vec<EventPtr>,
         ) -> UniquePtr<Event>;
+
         fn barrier(queue: &mut UniquePtr<Queue>, dep_events: Vec<EventPtr>) -> UniquePtr<Event>;
         fn wait(queue: &mut UniquePtr<Queue>);
+
+        unsafe fn launch_1d(
+            queue: &mut UniquePtr<Queue>,
+            global_size: Range1,
+            local_size: Range1,
+            kernel: &Kernel,
+            args: &[&[u8]],
+        ) -> UniquePtr<Event>;
+
+        unsafe fn launch_2d(
+            queue: &mut UniquePtr<Queue>,
+            global_size: Range2,
+            local_size: Range2,
+            kernel: &Kernel,
+            args: &[&[u8]],
+        ) -> UniquePtr<Event>;
+
+        unsafe fn launch_3d(
+            queue: &mut UniquePtr<Queue>,
+            global_size: Range3,
+            local_size: Range3,
+            kernel: &Kernel,
+            args: &[&[u8]],
+        ) -> UniquePtr<Event>;
     }
 }