diff --git a/oneapi-rs-sys/build.rs b/oneapi-rs-sys/build.rs index 853dffc..121cc68 100644 --- a/oneapi-rs-sys/build.rs +++ b/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 { diff --git a/oneapi-rs-sys/include/context.hpp b/oneapi-rs-sys/include/context.hpp new file mode 100644 index 0000000..ffc53b8 --- /dev/null +++ b/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 + +#include + +#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 new_context(rust::Vec); +} // namespace sycl_shims::context diff --git a/oneapi-rs-sys/include/device.hpp b/oneapi-rs-sys/include/device.hpp index 0cbbf1a..bbed562 100644 --- a/oneapi-rs-sys/include/device.hpp +++ b/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 get_platform(Device const &); +std::unique_ptr clone(Device const &); } // namespace sycl_shims::device diff --git a/oneapi-rs-sys/include/kernel-bundle.hpp b/oneapi-rs-sys/include/kernel-bundle.hpp new file mode 100644 index 0000000..cc7dbe7 --- /dev/null +++ b/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 + +#include + +#include "oneapi-rs-sys/include/types.hpp" +#include "rust/cxx.h" + +namespace sycl_shims::kernel_bundle { +std::unique_ptr +create_kernel_bundle_from_source(Context const &ctxt, rust::Str source); +std::unique_ptr +build(std::unique_ptr &source); +std::unique_ptr get_kernel(std::unique_ptr &, + rust::Str); +} // namespace sycl_shims::kernel_bundle diff --git a/oneapi-rs-sys/include/queue.hpp b/oneapi-rs-sys/include/queue.hpp index c8fa997..8ac204a 100644 --- a/oneapi-rs-sys/include/queue.hpp +++ b/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 new_queue(); std::unique_ptr new_queue_immediate(); std::unique_ptr new_queue_from_device(Device const &); +std::unique_ptr get_context(Queue const &); std::unique_ptr clone(Queue const &); std::unique_ptr memset(std::unique_ptr &, std::uint8_t *ptr, int value, std::size_t num_bytes, rust::Vec); std::unique_ptr barrier(std::unique_ptr &, rust::Vec); void wait(std::unique_ptr &); + +std::unique_ptr +launch_1d(std::unique_ptr &, Range1 global_size, Range1 local_size, + Kernel const &, + rust::Slice const> args); + +std::unique_ptr +launch_2d(std::unique_ptr &, Range2 global_size, Range2 local_size, + Kernel const &, + rust::Slice const> args); + +std::unique_ptr +launch_3d(std::unique_ptr &, Range3 global_size, Range3 local_size, + Kernel const &, + rust::Slice const> args); } // namespace sycl_shims::queue diff --git a/oneapi-rs-sys/include/types.hpp b/oneapi-rs-sys/include/types.hpp index db147b1..567da69 100644 --- a/oneapi-rs-sys/include/types.hpp +++ b/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; +using ExecutableKernelBundle = + sycl::kernel_bundle; } // namespace sycl_shims diff --git a/oneapi-rs-sys/src/context-sys.rs b/oneapi-rs-sys/src/context-sys.rs new file mode 100644 index 0000000..e5f9093 --- /dev/null +++ b/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) -> UniquePtr; + } +} diff --git a/oneapi-rs-sys/src/context.cpp b/oneapi-rs-sys/src/context.cpp new file mode 100644 index 0000000..dd91592 --- /dev/null +++ b/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 new_context(rust::Vec devices) { + std::vector raw_devices; + for (auto &&d : devices) + raw_devices.push_back(std::move(*d.ptr.release())); + return std::make_unique(raw_devices); +} +} // namespace sycl_shims::context diff --git a/oneapi-rs-sys/src/device-sys.rs b/oneapi-rs-sys/src/device-sys.rs index f272172..9c460ea 100644 --- a/oneapi-rs-sys/src/device-sys.rs +++ b/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; } } diff --git a/oneapi-rs-sys/src/device.cpp b/oneapi-rs-sys/src/device.cpp index 6d87359..1063f0f 100644 --- a/oneapi-rs-sys/src/device.cpp +++ b/oneapi-rs-sys/src/device.cpp @@ -53,4 +53,8 @@ rust::String get_name(Device const &device) { std::unique_ptr get_platform(Device const &device) { return std::make_unique(device.get_platform()); } + +std::unique_ptr clone(Device const &device) { + return std::make_unique(sycl::device(device)); +} } // namespace sycl_shims::device diff --git a/oneapi-rs-sys/src/kernel-bundle-sys.rs b/oneapi-rs-sys/src/kernel-bundle-sys.rs new file mode 100644 index 0000000..2e53ab4 --- /dev/null +++ b/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; + fn build(source: &mut UniquePtr) -> UniquePtr; + fn get_kernel( + bundle: &mut UniquePtr, + name: &str, + ) -> UniquePtr; + } +} diff --git a/oneapi-rs-sys/src/kernel-bundle.cpp b/oneapi-rs-sys/src/kernel-bundle.cpp new file mode 100644 index 0000000..27fe928 --- /dev/null +++ b/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 +create_kernel_bundle_from_source(Context const &ctxt, rust::Str source) { + return std::make_unique( + syclexp::create_kernel_bundle_from_source( + ctxt, syclexp::source_language::sycl, std::string(source))); +} + +std::unique_ptr +build(std::unique_ptr &source) { + return std::make_unique(syclexp::build(*source)); +} + +std::unique_ptr +get_kernel(std::unique_ptr &bundle, rust::Str name) { + return std::make_unique( + bundle->ext_oneapi_get_kernel(static_cast(name))); +} +} // namespace sycl_shims::kernel_bundle diff --git a/oneapi-rs-sys/src/lib.rs b/oneapi-rs-sys/src/lib.rs index f2f723d..514e0d8 100644 --- a/oneapi-rs-sys/src/lib.rs +++ b/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; diff --git a/oneapi-rs-sys/src/queue-sys.rs b/oneapi-rs-sys/src/queue-sys.rs index f79c6f5..d66e620 100644 --- a/oneapi-rs-sys/src/queue-sys.rs +++ b/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,12 +25,18 @@ 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; fn new_queue_immediate() -> UniquePtr; fn new_queue_from_device(device: &Device) -> UniquePtr; + fn get_context(queue: &Queue) -> UniquePtr; fn clone(queue: &Queue) -> UniquePtr; + unsafe fn memset( queue: &mut UniquePtr, ptr: *mut u8, @@ -35,7 +44,32 @@ pub mod ffi { num_bytes: usize, dep_events: Vec, ) -> UniquePtr; + fn barrier(queue: &mut UniquePtr, dep_events: Vec) -> UniquePtr; fn wait(queue: &mut UniquePtr); + + unsafe fn launch_1d( + queue: &mut UniquePtr, + global_size: Range1, + local_size: Range1, + kernel: &Kernel, + args: &[&[u8]], + ) -> UniquePtr; + + unsafe fn launch_2d( + queue: &mut UniquePtr, + global_size: Range2, + local_size: Range2, + kernel: &Kernel, + args: &[&[u8]], + ) -> UniquePtr; + + unsafe fn launch_3d( + queue: &mut UniquePtr, + global_size: Range3, + local_size: Range3, + kernel: &Kernel, + args: &[&[u8]], + ) -> UniquePtr; } } diff --git a/oneapi-rs-sys/src/queue.cpp b/oneapi-rs-sys/src/queue.cpp index 3209cd1..afef17b 100644 --- a/oneapi-rs-sys/src/queue.cpp +++ b/oneapi-rs-sys/src/queue.cpp @@ -12,6 +12,8 @@ using sycl::ext::intel::property::queue::immediate_command_list; using sycl::property::queue::in_order; +namespace syclexp = sycl::ext::oneapi::experimental; + namespace sycl_shims::queue { std::unique_ptr new_queue() { return std::make_unique(sycl::queue({in_order()})); @@ -26,6 +28,10 @@ std::unique_ptr new_queue_from_device(Device const &device) { return std::make_unique(sycl::queue(device, {in_order()})); } +std::unique_ptr get_context(Queue const &queue) { + return std::make_unique(queue.get_context()); +} + std::unique_ptr clone(Queue const &queue) { return std::make_unique(sycl::queue(queue)); } @@ -48,4 +54,48 @@ std::unique_ptr barrier(std::unique_ptr &queue, } void wait(std::unique_ptr &queue) { queue->wait(); } + +template +std::unique_ptr +launch(std::unique_ptr &queue, sycl::nd_range nd_range, + Kernel const &kernel, + rust::Slice const> args) { + return std::make_unique(queue->submit([&](sycl::handler &cgh) { + for (std::size_t i = 0; i < args.size(); ++i) + cgh.set_arg(i, syclexp::raw_kernel_arg(args[i].data(), args[i].size())); + + cgh.parallel_for(nd_range, kernel); + })); +} + +std::unique_ptr +launch_1d(std::unique_ptr &queue, Range1 global_size, Range1 local_size, + Kernel const &kernel, + rust::Slice const> args) { + return launch(queue, + sycl::nd_range<1>{{global_size.data[0]}, {local_size.data[0]}}, + kernel, args); +} + +std::unique_ptr +launch_2d(std::unique_ptr &queue, Range2 global_size, Range2 local_size, + Kernel const &kernel, + rust::Slice const> args) { + return launch(queue, + sycl::nd_range<2>{{global_size.data[0], global_size.data[1]}, + {local_size.data[0], local_size.data[1]}}, + kernel, args); +} + +std::unique_ptr +launch_3d(std::unique_ptr &queue, Range3 global_size, Range3 local_size, + Kernel const &kernel, + rust::Slice const> args) { + return launch( + queue, + sycl::nd_range<3>{ + {global_size.data[0], global_size.data[1], global_size.data[2]}, + {local_size.data[0], local_size.data[1], local_size.data[2]}}, + kernel, args); +} } // namespace sycl_shims::queue diff --git a/oneapi-rs-sys/src/types-sys.rs b/oneapi-rs-sys/src/types-sys.rs index 2c9a49e..c675165 100644 --- a/oneapi-rs-sys/src/types-sys.rs +++ b/oneapi-rs-sys/src/types-sys.rs @@ -37,6 +37,10 @@ pub mod ffi { type Platform; type Queue; type Event; + type Context; + type Kernel; + type SourceKernelBundle; + type ExecutableKernelBundle; } // This is a workaround - cxx currently doesn't support passing @@ -74,10 +78,27 @@ pub mod ffi { Unknown, } + // cxx doesn't support const generic parameters + struct Range1 { + data: [u64; 1], + } + + struct Range2 { + data: [u64; 2], + } + + struct Range3 { + data: [u64; 3], + } + impl UniquePtr {} impl UniquePtr {} impl UniquePtr {} impl UniquePtr {} + impl UniquePtr {} + impl UniquePtr {} + impl UniquePtr {} + impl UniquePtr {} impl Vec {} impl Vec {} diff --git a/oneapi-rs/examples/kernel_launch.rs b/oneapi-rs/examples/kernel_launch.rs new file mode 100644 index 0000000..6224b27 --- /dev/null +++ b/oneapi-rs/examples/kernel_launch.rs @@ -0,0 +1,69 @@ +// +// 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 +// + +use oneapi_rs::{ + buffer::Buffer, + kernel::{KernelArgument, KernelArgumentList}, + queue::Queue, + range::NdRange, + usm::{SharedAllocator, UsmAllocator}, +}; + +static IOTA_SRC: &str = r#" +#include +namespace syclext = sycl::ext::oneapi; +namespace syclexp = sycl::ext::oneapi::experimental; + +extern "C" +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +void iota(float start, float *ptr) { + size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); + ptr[id] = start + static_cast(id); +} +"#; + +struct IotaArgs<'a> { + start: f32, + buffer: &'a mut Buffer>, +} + +unsafe impl<'a> KernelArgumentList<2> for IotaArgs<'a> { + unsafe fn as_raw_arg_list(&self) -> [&[u8]; 2] { + return [unsafe { self.start.as_raw_arg() }, unsafe { + self.buffer.as_raw_arg() + }]; + } +} + +fn main() { + let mut queue = Queue::new(); + let mut buffer = queue.alloc_shared::(1024).wait(); + + let kernel = queue + .get_context() + .create_kernel_bundle_from_source(IOTA_SRC) + .build() + .get_kernel("iota"); + + unsafe { + queue.launch( + NdRange::new([1024], [16]), + &kernel, + IotaArgs { + start: 3.14, + buffer: &mut buffer, + }, + ) + } + .wait(); + + for e in buffer.iter() { + print!("{e} "); + } + println!(); +} diff --git a/oneapi-rs/src/buffer.rs b/oneapi-rs/src/buffer.rs index 32d2a44..4ef71c2 100644 --- a/oneapi-rs/src/buffer.rs +++ b/oneapi-rs/src/buffer.rs @@ -15,10 +15,12 @@ use std::{ task::{Context, Poll}, }; +use bytemuck::Pod; use pin_project::pin_project; use crate::{ event::{Event, EventFuture}, + kernel::KernelArgument, usm::UsmAlloc, }; @@ -138,3 +140,11 @@ impl IntoFuture for EnqueuedBuffer { } } } + +unsafe impl KernelArgument for Buffer { + unsafe fn as_raw_arg(&self) -> &[u8] { + let data_ptr: *const NonNull<_> = &self.data; + let cast_ptr = data_ptr as *const u8; + unsafe { slice::from_raw_parts(cast_ptr, std::mem::size_of::<*mut u8>()) } + } +} diff --git a/oneapi-rs/src/context.rs b/oneapi-rs/src/context.rs new file mode 100644 index 0000000..0dfff64 --- /dev/null +++ b/oneapi-rs/src/context.rs @@ -0,0 +1,38 @@ +// +// 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 +// + +use oneapi_rs_sys::{context::ffi, kernel_bundle, types::ffi::DevicePtr}; + +use crate::{device::Device, kernel::SourceKernelBundle}; + +/// A context represents the runtime data structures and state required by a SYCL backend API +/// to interact with a group of devices associated with a platform. +pub struct Context(pub(crate) cxx::UniquePtr); + +impl From> for Context { + fn from(value: cxx::UniquePtr) -> Self { + Self(value) + } +} + +impl Context { + pub fn new(devices: &[&Device]) -> Self { + let devices = devices + .iter() + .map(|d| DevicePtr { + ptr: (*d).clone().0, + }) + .collect::>(); + + ffi::new_context(devices).into() + } + + pub fn create_kernel_bundle_from_source(&self, source: &str) -> SourceKernelBundle { + kernel_bundle::ffi::create_kernel_bundle_from_source(&self.0, source).into() + } +} diff --git a/oneapi-rs/src/device.rs b/oneapi-rs/src/device.rs index cee1aac..8560062 100644 --- a/oneapi-rs/src/device.rs +++ b/oneapi-rs/src/device.rs @@ -15,6 +15,12 @@ use crate::{info::device::DeviceInfo, platform::Platform}; /// The `Device` struct provides the common reference semantics. pub struct Device(pub(crate) cxx::UniquePtr); +impl From> for Device { + fn from(value: cxx::UniquePtr) -> Self { + Self(value) + } +} + impl Device { /// Returns a [`Vec`] containing all the root devices from all SYCL backends /// available in the system which have the device type encapsulated by [`DeviceType`](crate::info::DeviceType). @@ -39,3 +45,9 @@ impl Device { Platform(raw_platform) } } + +impl Clone for Device { + fn clone(&self) -> Self { + ffi::clone(&self.0).into() + } +} diff --git a/oneapi-rs/src/kernel.rs b/oneapi-rs/src/kernel.rs new file mode 100644 index 0000000..0e5d6a4 --- /dev/null +++ b/oneapi-rs/src/kernel.rs @@ -0,0 +1,65 @@ +// +// 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 +// + +use bytemuck::Pod; +use oneapi_rs_sys::{kernel_bundle::ffi, types}; + +/// A kernel bundle which stores loaded SYCL source code. +pub struct SourceKernelBundle(pub(crate) cxx::UniquePtr); + +impl From> for SourceKernelBundle { + fn from(value: cxx::UniquePtr) -> Self { + Self(value) + } +} + +impl SourceKernelBundle { + pub fn build(&mut self) -> ExecutableKernelBundle { + ffi::build(&mut self.0).into() + } +} + +/// A kernel bundle which stores compiled SYCL kernels. +pub struct ExecutableKernelBundle(pub(crate) cxx::UniquePtr); + +impl From> for ExecutableKernelBundle { + fn from(value: cxx::UniquePtr) -> Self { + Self(value) + } +} + +impl ExecutableKernelBundle { + pub fn get_kernel(&mut self, name: &str) -> Kernel { + ffi::get_kernel(&mut self.0, name).into() + } +} + +/// An executable SYCL kernel. +pub struct Kernel(pub(crate) cxx::UniquePtr); + +impl From> for Kernel { + fn from(value: cxx::UniquePtr) -> Self { + Self(value) + } +} + +/// Types which can be passed as SYCL kernel arguments. +pub unsafe trait KernelArgument { + unsafe fn as_raw_arg(&self) -> &[u8]; +} + +unsafe impl KernelArgument for T { + unsafe fn as_raw_arg(&self) -> &[u8] { + bytemuck::bytes_of(self) + } +} + +/// Types which describe an argument list for a SYCL kernel. +pub unsafe trait KernelArgumentList { + unsafe fn as_raw_arg_list(&self) -> [&[u8]; ARGC]; +} diff --git a/oneapi-rs/src/lib.rs b/oneapi-rs/src/lib.rs index 922b460..3e274b2 100644 --- a/oneapi-rs/src/lib.rs +++ b/oneapi-rs/src/lib.rs @@ -7,9 +7,12 @@ // pub mod buffer; +pub mod context; pub mod device; pub mod event; pub mod info; +pub mod kernel; pub mod platform; pub mod queue; +pub mod range; pub mod usm; diff --git a/oneapi-rs/src/queue.rs b/oneapi-rs/src/queue.rs index 70ba4fc..2f8e351 100644 --- a/oneapi-rs/src/queue.rs +++ b/oneapi-rs/src/queue.rs @@ -11,8 +11,11 @@ use oneapi_rs_sys::{queue::ffi, types::ffi::EventPtr}; use crate::{ buffer::{Buffer, EnqueuedBuffer}, + context::Context, device::Device, event::Event, + kernel::{Kernel, KernelArgumentList}, + range::{NdRange, ValidDimension}, usm::{HostAllocator, SharedAllocator, UsmAlloc, UsmAllocator}, }; @@ -32,6 +35,11 @@ impl Queue { Self(ffi::new_queue_immediate()) } + /// Returns the SYCL queue’s context. + pub fn get_context(&self) -> Context { + ffi::get_context(&self.0).into() + } + /// Allocates zeroed memory and creates a host-side [`Buffer`] that can store an array of T. pub fn alloc_host( &mut self, @@ -125,6 +133,20 @@ impl Queue { pub fn wait(&mut self) { ffi::wait(&mut self.0); } + + /// Enqueues a kernel object to the queue an an ND-range kernel, using the number of work-items + /// specified by the [`NdRange`] nd_range. + pub unsafe fn launch( + &mut self, + nd_range: NdRange, + kernel: &Kernel, + args: impl KernelArgumentList, + ) -> Event + where + NdRange: ValidDimension, + { + unsafe { nd_range.launch(self, kernel, args) } + } } impl From<&Device> for Queue { diff --git a/oneapi-rs/src/range.rs b/oneapi-rs/src/range.rs new file mode 100644 index 0000000..8c11aca --- /dev/null +++ b/oneapi-rs/src/range.rs @@ -0,0 +1,120 @@ +// +// 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 +// + +use oneapi_rs_sys::types; + +use crate::{ + event::Event, + kernel::{Kernel, KernelArgumentList}, + queue::Queue, +}; + +/// `Range` is a 1D, 2D or 3D vector that defines the iteration domain of either a single work-group +/// in a parallel dispatch, or the overall dimensions of the dispatch. +pub type Range = [u64; DIMENSIONS]; + +/// The `NdRange` struct defines the iteration domain of both the work-groups and the overall +/// dispatch. +/// +/// An `NdRange` comprises two [`Range`] parameters: the whole range over which the kernel is to be +/// executed and the range of each work group. +pub struct NdRange { + pub group_size: Range, + pub local_size: Range, +} + +impl NdRange { + pub fn new(group_size: Range, local_size: Range) -> Self { + Self { + group_size, + local_size, + } + } +} + +/// [`NdRange`] types which are limited to 1, 2 or 3 dimensions. +pub trait ValidDimension { + unsafe fn launch( + &self, + queue: &mut Queue, + kernel: &Kernel, + args: impl KernelArgumentList, + ) -> Event; +} + +impl ValidDimension for NdRange<1> { + unsafe fn launch( + &self, + queue: &mut Queue, + kernel: &Kernel, + args: impl KernelArgumentList, + ) -> Event { + unsafe { + oneapi_rs_sys::queue::ffi::launch_1d( + &mut queue.0, + types::ffi::Range1 { + data: self.group_size, + }, + types::ffi::Range1 { + data: self.local_size, + }, + &kernel.0, + &args.as_raw_arg_list(), + ) + } + .into() + } +} + +impl ValidDimension for NdRange<2> { + unsafe fn launch( + &self, + queue: &mut Queue, + kernel: &Kernel, + args: impl KernelArgumentList, + ) -> Event { + unsafe { + oneapi_rs_sys::queue::ffi::launch_2d( + &mut queue.0, + types::ffi::Range2 { + data: self.group_size, + }, + types::ffi::Range2 { + data: self.local_size, + }, + &kernel.0, + &args.as_raw_arg_list(), + ) + } + .into() + } +} + +impl ValidDimension for NdRange<3> { + unsafe fn launch( + &self, + queue: &mut Queue, + kernel: &Kernel, + args: impl KernelArgumentList, + ) -> Event { + unsafe { + oneapi_rs_sys::queue::ffi::launch_3d( + &mut queue.0, + types::ffi::Range3 { + data: self.group_size, + }, + types::ffi::Range3 { + data: self.local_size, + }, + &kernel.0, + &args.as_raw_arg_list(), + ) + } + .into() + } +}