From 052e892cd5efb1de2dadc77be82bdb9ba1281daf Mon Sep 17 00:00:00 2001 From: Zhiwei Liang Date: Sat, 31 Oct 2026 07:10:22 +0000 Subject: [PATCH 1/3] Carry a kind for each kernel argument through the launch bridge The launch functions in sycl-rs-sys took each kernel argument as plain bytes and always handed them to raw_kernel_arg, so the C++ shim could not tell a USM pointer from a value. A backend may take pointers through a different entry point than plain bytes: the OpenCL CPU device rejects USM pointers passed through raw_kernel_arg. Pass each argument as a KernelArg with a kind. Bytes still go through raw_kernel_arg, and pointers go through the typed handler::set_arg overload. sycl-rs passes every argument as Bytes for now, so behaviour is unchanged. Signed-off-by: Zhiwei Liang --- sycl/sycl-rs-sys/include/queue.hpp | 22 ++++++------ sycl/sycl-rs-sys/src/queue-sys.rs | 13 +++---- sycl/sycl-rs-sys/src/queue.cpp | 56 ++++++++++++++++++++++-------- sycl/sycl-rs-sys/src/types-sys.rs | 16 +++++++++ sycl/sycl-rs/src/range.rs | 15 ++++++-- 5 files changed, 86 insertions(+), 36 deletions(-) diff --git a/sycl/sycl-rs-sys/include/queue.hpp b/sycl/sycl-rs-sys/include/queue.hpp index 0fe5bbe..9f0c08e 100644 --- a/sycl/sycl-rs-sys/include/queue.hpp +++ b/sycl/sycl-rs-sys/include/queue.hpp @@ -18,6 +18,7 @@ struct EventPtr; struct Range1; struct Range2; struct Range3; +struct KernelArg; } // namespace sycl_shims namespace sycl_shims::queue { @@ -37,20 +38,17 @@ std::unique_ptr memset(std::unique_ptr &, std::uint8_t *ptr, 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_1d(std::unique_ptr &, Range1 global_size, + Range1 local_size, Kernel const &, + rust::Slice 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_2d(std::unique_ptr &, Range2 global_size, + Range2 local_size, Kernel const &, + rust::Slice args); -std::unique_ptr -launch_3d(std::unique_ptr &, Range3 global_size, Range3 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 args); std::unique_ptr memcpy(std::unique_ptr &, std::uint8_t *dest, std::uint8_t const *src, std::size_t num_bytes, diff --git a/sycl/sycl-rs-sys/src/queue-sys.rs b/sycl/sycl-rs-sys/src/queue-sys.rs index 2310585..707fd0e 100644 --- a/sycl/sycl-rs-sys/src/queue-sys.rs +++ b/sycl/sycl-rs-sys/src/queue-sys.rs @@ -17,6 +17,7 @@ pub mod ffi { type Range1 = crate::types::ffi::Range1; type Range2 = crate::types::ffi::Range2; type Range3 = crate::types::ffi::Range3; + type KernelArg<'a> = crate::types::ffi::KernelArg<'a>; } unsafe extern "C++" { @@ -82,36 +83,36 @@ pub mod ffi { /// /// Each entry in `args` must match the corresponding kernel parameter's size, layout, /// and alignment. - unsafe fn launch_1d( + unsafe fn launch_1d<'a>( queue: &mut UniquePtr, global_size: Range1, local_size: Range1, kernel: &Kernel, - args: &[&[u8]], + args: &[KernelArg<'a>], ) -> Result>; /// # Safety /// /// Each entry in `args` must match the corresponding kernel parameter's size, layout, /// and alignment. - unsafe fn launch_2d( + unsafe fn launch_2d<'a>( queue: &mut UniquePtr, global_size: Range2, local_size: Range2, kernel: &Kernel, - args: &[&[u8]], + args: &[KernelArg<'a>], ) -> Result>; /// # Safety /// /// Each entry in `args` must match the corresponding kernel parameter's size, layout, /// and alignment. - unsafe fn launch_3d( + unsafe fn launch_3d<'a>( queue: &mut UniquePtr, global_size: Range3, local_size: Range3, kernel: &Kernel, - args: &[&[u8]], + args: &[KernelArg<'a>], ) -> Result>; } } diff --git a/sycl/sycl-rs-sys/src/queue.cpp b/sycl/sycl-rs-sys/src/queue.cpp index 5698ffc..9e097a4 100644 --- a/sycl/sycl-rs-sys/src/queue.cpp +++ b/sycl/sycl-rs-sys/src/queue.cpp @@ -10,6 +10,9 @@ #include "sycl-rs-sys/include/utils.hpp" #include "sycl-rs-sys/src/queue-sys.rs.h" +#include +#include + using sycl::ext::intel::property::queue::immediate_command_list; using sycl::property::queue::in_order; using sycl_shims::utils::vec_to_vector; @@ -62,42 +65,65 @@ std::unique_ptr barrier(std::unique_ptr &queue, void wait(std::unique_ptr &queue) { queue->wait(); } +// Pointers go through the typed set_arg overload rather than raw_kernel_arg: a +// backend may take them through a different entry point than plain bytes (the +// OpenCL CPU device rejects USM pointers passed as bytes), and raw_kernel_arg +// only takes non-pointer values. +static void set_kernel_arg(sycl::handler &cgh, std::size_t index, + KernelArg const &arg) { + switch (arg.kind) { + case KernelArgKind::Bytes: + cgh.set_arg(index, + syclexp::raw_kernel_arg(arg.bytes.data(), arg.bytes.size())); + return; + case KernelArgKind::Pointer: { + void *pointer; + if (arg.bytes.size() != sizeof(pointer)) + throw std::invalid_argument( + "pointer kernel argument is not pointer-sized"); + std::memcpy(&pointer, arg.bytes.data(), sizeof(pointer)); + cgh.set_arg(index, pointer); + return; + } + } + throw std::invalid_argument("unknown kernel argument kind"); +} + template std::unique_ptr launch(std::unique_ptr &queue, sycl::nd_range nd_range, - Kernel const &kernel, - rust::Slice const> args) { + Kernel const &kernel, rust::Slice 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())); + set_kernel_arg(cgh, i, args[i]); 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) { +std::unique_ptr launch_1d(std::unique_ptr &queue, + Range1 global_size, Range1 local_size, + Kernel const &kernel, + rust::Slice 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) { +std::unique_ptr launch_2d(std::unique_ptr &queue, + Range2 global_size, Range2 local_size, + Kernel const &kernel, + rust::Slice 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) { +std::unique_ptr launch_3d(std::unique_ptr &queue, + Range3 global_size, Range3 local_size, + Kernel const &kernel, + rust::Slice args) { return launch( queue, sycl::nd_range<3>{ diff --git a/sycl/sycl-rs-sys/src/types-sys.rs b/sycl/sycl-rs-sys/src/types-sys.rs index 3136ca6..2dfd99c 100644 --- a/sycl/sycl-rs-sys/src/types-sys.rs +++ b/sycl/sycl-rs-sys/src/types-sys.rs @@ -106,6 +106,22 @@ pub mod ffi { data: [u64; 3], } + /// How the bytes of a kernel argument are handed to the backend. A backend may take + /// pointers through a different entry point than plain bytes, and cannot tell the two + /// apart from the bytes alone. + #[derive(Debug, Hash)] + enum KernelArgKind { + /// The bytes of a trivially copyable value that is not a pointer. + Bytes, + /// The bytes of a pointer to USM memory. + Pointer, + } + + struct KernelArg<'a> { + kind: KernelArgKind, + bytes: &'a [u8], + } + impl UniquePtr {} impl UniquePtr {} impl UniquePtr {} diff --git a/sycl/sycl-rs/src/range.rs b/sycl/sycl-rs/src/range.rs index c8b4059..3fd2aa0 100644 --- a/sycl/sycl-rs/src/range.rs +++ b/sycl/sycl-rs/src/range.rs @@ -72,7 +72,10 @@ impl ValidDimension for NdRange<1> { data: self.local_size, }, &kernel.0, - &args.as_raw_arg_list(), + &args.as_raw_arg_list().map(|bytes| types::ffi::KernelArg { + kind: types::ffi::KernelArgKind::Bytes, + bytes, + }), ) } .map(Into::into) @@ -97,7 +100,10 @@ impl ValidDimension for NdRange<2> { data: self.local_size, }, &kernel.0, - &args.as_raw_arg_list(), + &args.as_raw_arg_list().map(|bytes| types::ffi::KernelArg { + kind: types::ffi::KernelArgKind::Bytes, + bytes, + }), ) } .map(Into::into) @@ -122,7 +128,10 @@ impl ValidDimension for NdRange<3> { data: self.local_size, }, &kernel.0, - &args.as_raw_arg_list(), + &args.as_raw_arg_list().map(|bytes| types::ffi::KernelArg { + kind: types::ffi::KernelArgKind::Bytes, + bytes, + }), ) } .map(Into::into) From 922039508c0ab47942f2d4c15bfc6dccc5db85b9 Mon Sep 17 00:00:00 2001 From: Zhiwei Liang Date: Sat, 31 Oct 2026 07:11:17 +0000 Subject: [PATCH 2/3] Return RawKernelArg from KernelArgument KernelArgument::as_raw_arg and KernelArgumentList::as_raw_arg_list returned plain byte slices. Return a RawKernelArg instead, so that a later change can mark USM pointers. Every argument is still passed as bytes, so behaviour is unchanged. RawKernelArg is non-exhaustive so that more kinds can be added without a breaking change. Signed-off-by: Zhiwei Liang --- sycl/sycl-rs-derive/src/lib.rs | 4 ++-- sycl/sycl-rs/src/kernel.rs | 34 ++++++++++++++++++++++++++-------- sycl/sycl-rs/src/prelude.rs | 2 +- sycl/sycl-rs/src/range.rs | 17 ++++------------- sycl/sycl-rs/src/usmbox.rs | 14 +++++++------- 5 files changed, 40 insertions(+), 31 deletions(-) diff --git a/sycl/sycl-rs-derive/src/lib.rs b/sycl/sycl-rs-derive/src/lib.rs index bc37a49..150c5cd 100644 --- a/sycl/sycl-rs-derive/src/lib.rs +++ b/sycl/sycl-rs-derive/src/lib.rs @@ -43,7 +43,7 @@ pub fn derive_kernel_argument_list(input: TokenStream) -> TokenStream { let expanded = quote! { unsafe impl #impl_generics #oneapi::kernel::KernelArgumentList<#argc> for #ident #ty_generics #where_clause { - unsafe fn as_raw_arg_list(&self) -> [&[u8]; #argc] { + unsafe fn as_raw_arg_list(&self) -> [#oneapi::kernel::RawKernelArg<'_>; #argc] { [ #(unsafe { self.#members.as_raw_arg() }),* ] } } @@ -69,7 +69,7 @@ fn get_single_tuple_impl(argc: usize) -> proc_macro2::TokenStream { quote! { unsafe impl<#(#types),*> crate::kernel::KernelArgumentList<#argc> for (#(#types),*) where #(#types: crate::kernel::KernelArgument),* { - unsafe fn as_raw_arg_list(&self) -> [&[u8]; #argc] { + unsafe fn as_raw_arg_list(&self) -> [crate::kernel::RawKernelArg<'_>; #argc] { [ #(unsafe { self.#iter.as_raw_arg() }),* ] } } diff --git a/sycl/sycl-rs/src/kernel.rs b/sycl/sycl-rs/src/kernel.rs index d1ce3a8..fce8293 100644 --- a/sycl/sycl-rs/src/kernel.rs +++ b/sycl/sycl-rs/src/kernel.rs @@ -49,6 +49,24 @@ impl From> for Kernel { } } +/// The raw representation of a SYCL kernel argument. +#[derive(Clone, Copy, Debug)] +#[non_exhaustive] +pub enum RawKernelArg<'a> { + /// The bytes of a trivially copyable value. + Bytes(&'a [u8]), +} + +impl<'a> RawKernelArg<'a> { + pub(crate) fn into_ffi(self) -> types::ffi::KernelArg<'a> { + let Self::Bytes(bytes) = self; + types::ffi::KernelArg { + kind: types::ffi::KernelArgKind::Bytes, + bytes, + } + } +} + /// Types which can be passed as SYCL kernel arguments. /// /// # Safety @@ -56,18 +74,18 @@ impl From> for Kernel { /// A type implementing this trait must mirror the representation and alignment of the /// corresponding SYCL kernel argument structure. pub unsafe trait KernelArgument { - /// Converts self to a raw byte representation. + /// Converts self to its raw representation. /// /// # Safety /// /// This function returns a reference to raw bytes. These bytes will be passed to FFI /// functions. The caller must make sure these functions respect Rust's aliasing rules. - unsafe fn as_raw_arg(&self) -> &[u8]; + unsafe fn as_raw_arg(&self) -> RawKernelArg<'_>; } unsafe impl KernelArgument for T { - unsafe fn as_raw_arg(&self) -> &[u8] { - bytemuck::bytes_of(self) + unsafe fn as_raw_arg(&self) -> RawKernelArg<'_> { + RawKernelArg::Bytes(bytemuck::bytes_of(self)) } } @@ -78,23 +96,23 @@ unsafe impl KernelArgument for T { /// A type implementing this trait must mirror the representation and alignment of each /// corresponding SYCL kernel argument inside the returned array. pub unsafe trait KernelArgumentList { - /// Converts each struct member to a raw byte representation. + /// Converts each struct member to its raw representation. /// /// # Safety /// /// This function returns references to raw bytes. These bytes will be passed to FFI /// functions. The caller must make sure these functions respect Rust's aliasing rules. - unsafe fn as_raw_arg_list(&self) -> [&[u8]; ARGC]; + unsafe fn as_raw_arg_list(&self) -> [RawKernelArg<'_>; ARGC]; } unsafe impl KernelArgumentList<0> for () { - unsafe fn as_raw_arg_list(&self) -> [&[u8]; 0] { + unsafe fn as_raw_arg_list(&self) -> [RawKernelArg<'_>; 0] { [] } } unsafe impl KernelArgumentList<1> for T { - unsafe fn as_raw_arg_list(&self) -> [&[u8]; 1] { + unsafe fn as_raw_arg_list(&self) -> [RawKernelArg<'_>; 1] { [unsafe { self.as_raw_arg() }] } } diff --git a/sycl/sycl-rs/src/prelude.rs b/sycl/sycl-rs/src/prelude.rs index 2ab3549..2bee2aa 100644 --- a/sycl/sycl-rs/src/prelude.rs +++ b/sycl/sycl-rs/src/prelude.rs @@ -10,7 +10,7 @@ pub use crate::{ context::Context, device::{Aspect, Device}, info::{self, InfoTarget}, - kernel::{Kernel, KernelArgument, KernelArgumentList}, + kernel::{Kernel, KernelArgument, KernelArgumentList, RawKernelArg}, platform::Platform, queue::Queue, range::{NdRange, Range}, diff --git a/sycl/sycl-rs/src/range.rs b/sycl/sycl-rs/src/range.rs index 3fd2aa0..b82dfbe 100644 --- a/sycl/sycl-rs/src/range.rs +++ b/sycl/sycl-rs/src/range.rs @@ -11,7 +11,7 @@ use sycl_rs_sys::types; use crate::{ event::Event, - kernel::{Kernel, KernelArgumentList}, + kernel::{Kernel, KernelArgumentList, RawKernelArg}, private::Sealed, queue::Queue, }; @@ -72,10 +72,7 @@ impl ValidDimension for NdRange<1> { data: self.local_size, }, &kernel.0, - &args.as_raw_arg_list().map(|bytes| types::ffi::KernelArg { - kind: types::ffi::KernelArgKind::Bytes, - bytes, - }), + &args.as_raw_arg_list().map(RawKernelArg::into_ffi), ) } .map(Into::into) @@ -100,10 +97,7 @@ impl ValidDimension for NdRange<2> { data: self.local_size, }, &kernel.0, - &args.as_raw_arg_list().map(|bytes| types::ffi::KernelArg { - kind: types::ffi::KernelArgKind::Bytes, - bytes, - }), + &args.as_raw_arg_list().map(RawKernelArg::into_ffi), ) } .map(Into::into) @@ -128,10 +122,7 @@ impl ValidDimension for NdRange<3> { data: self.local_size, }, &kernel.0, - &args.as_raw_arg_list().map(|bytes| types::ffi::KernelArg { - kind: types::ffi::KernelArgKind::Bytes, - bytes, - }), + &args.as_raw_arg_list().map(RawKernelArg::into_ffi), ) } .map(Into::into) diff --git a/sycl/sycl-rs/src/usmbox.rs b/sycl/sycl-rs/src/usmbox.rs index b9ba1ab..e706447 100644 --- a/sycl/sycl-rs/src/usmbox.rs +++ b/sycl/sycl-rs/src/usmbox.rs @@ -21,7 +21,7 @@ use pin_project::pin_project; use crate::{ Result, event::{Event, EventFuture}, - kernel::KernelArgument, + kernel::{KernelArgument, RawKernelArg}, usm::{ DeviceAllocator, HostAccessible, HostAllocator, SharedAllocator, UsmAlloc, UsmAllocator, }, @@ -169,19 +169,19 @@ pub type SharedUsmBoxFuture = UsmBoxFuture>; pub type DeviceUsmBoxFuture = UsmBoxFuture>; unsafe impl KernelArgument for UsmBox { - unsafe fn as_raw_arg(&self) -> &[u8] { - unsafe { self.as_raw_arg_impl() } + unsafe fn as_raw_arg(&self) -> RawKernelArg<'_> { + RawKernelArg::Bytes(unsafe { self.as_raw_arg_impl() }) } } unsafe impl KernelArgument for &UsmBox { - unsafe fn as_raw_arg(&self) -> &[u8] { - unsafe { self.as_raw_arg_impl() } + unsafe fn as_raw_arg(&self) -> RawKernelArg<'_> { + RawKernelArg::Bytes(unsafe { self.as_raw_arg_impl() }) } } unsafe impl KernelArgument for &mut UsmBox { - unsafe fn as_raw_arg(&self) -> &[u8] { - unsafe { self.as_raw_arg_impl() } + unsafe fn as_raw_arg(&self) -> RawKernelArg<'_> { + RawKernelArg::Bytes(unsafe { self.as_raw_arg_impl() }) } } From c8c5b587ea7897d4dd89d91ab35e05760eaa1945 Mon Sep 17 00:00:00 2001 From: Zhiwei Liang Date: Sun, 1 Nov 2026 01:40:59 +0000 Subject: [PATCH 3/3] Mark USM pointer kernel arguments UsmBox kernel arguments were passed as plain bytes. On the OpenCL CPU device, kernel launches with them failed with UR_RESULT_ERROR_INVALID_MEM_OBJECT, so no kernel could run on a CPU. Add RawKernelArg::Pointer, return it for UsmBox arguments, and require KernelArgument and KernelArgumentList implementations to mark USM pointers this way. The launch bridge then passes them through handler::set_arg. Since the pointer bytes now matter, take their size from the NonNull field instead of an unrelated *const u8. Add a regression test, which fails without this change on the OpenCL CPU device. It prefers the CPU device, because the default device is a GPU when there is one, and GPUs accepted pointers passed as plain bytes. Signed-off-by: Zhiwei Liang --- CHANGELOG.md | 3 ++ sycl/sycl-rs-sys/src/queue-sys.rs | 6 ++-- sycl/sycl-rs/src/kernel.rs | 33 ++++++++++------- sycl/sycl-rs/src/usmbox.rs | 12 +++---- sycl/sycl-rs/tests/launch.rs | 59 +++++++++++++++++++++++++++++++ 5 files changed, 92 insertions(+), 21 deletions(-) create mode 100644 sycl/sycl-rs/tests/launch.rs diff --git a/CHANGELOG.md b/CHANGELOG.md index 446be5c..1cd03df 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -27,3 +27,6 @@ can be found at: - Runtime kernel bundle compilation and one-, two-, and three-dimensional kernel launch support. - Derive support for typed kernel argument lists. +- `RawKernelArg`, which marks USM pointer kernel arguments apart from plain + bytes, so kernels also launch on backends that take pointers through a + separate entry point, such as the OpenCL CPU device. diff --git a/sycl/sycl-rs-sys/src/queue-sys.rs b/sycl/sycl-rs-sys/src/queue-sys.rs index 707fd0e..3d34107 100644 --- a/sycl/sycl-rs-sys/src/queue-sys.rs +++ b/sycl/sycl-rs-sys/src/queue-sys.rs @@ -82,7 +82,7 @@ pub mod ffi { /// # Safety /// /// Each entry in `args` must match the corresponding kernel parameter's size, layout, - /// and alignment. + /// and alignment, and must be marked as a pointer exactly when it holds a USM pointer. unsafe fn launch_1d<'a>( queue: &mut UniquePtr, global_size: Range1, @@ -94,7 +94,7 @@ pub mod ffi { /// # Safety /// /// Each entry in `args` must match the corresponding kernel parameter's size, layout, - /// and alignment. + /// and alignment, and must be marked as a pointer exactly when it holds a USM pointer. unsafe fn launch_2d<'a>( queue: &mut UniquePtr, global_size: Range2, @@ -106,7 +106,7 @@ pub mod ffi { /// # Safety /// /// Each entry in `args` must match the corresponding kernel parameter's size, layout, - /// and alignment. + /// and alignment, and must be marked as a pointer exactly when it holds a USM pointer. unsafe fn launch_3d<'a>( queue: &mut UniquePtr, global_size: Range3, diff --git a/sycl/sycl-rs/src/kernel.rs b/sycl/sycl-rs/src/kernel.rs index fce8293..dce18d0 100644 --- a/sycl/sycl-rs/src/kernel.rs +++ b/sycl/sycl-rs/src/kernel.rs @@ -50,20 +50,25 @@ impl From> for Kernel { } /// The raw representation of a SYCL kernel argument. +/// +/// A backend may take pointers through a different entry point than plain bytes, and cannot +/// tell the two apart from the bytes alone, so pointers are marked as such. #[derive(Clone, Copy, Debug)] #[non_exhaustive] pub enum RawKernelArg<'a> { - /// The bytes of a trivially copyable value. + /// The bytes of a trivially copyable value that is not a pointer. Bytes(&'a [u8]), + /// The bytes of a pointer to USM memory. + Pointer(&'a [u8]), } impl<'a> RawKernelArg<'a> { pub(crate) fn into_ffi(self) -> types::ffi::KernelArg<'a> { - let Self::Bytes(bytes) = self; - types::ffi::KernelArg { - kind: types::ffi::KernelArgKind::Bytes, - bytes, - } + let (kind, bytes) = match self { + Self::Bytes(bytes) => (types::ffi::KernelArgKind::Bytes, bytes), + Self::Pointer(bytes) => (types::ffi::KernelArgKind::Pointer, bytes), + }; + types::ffi::KernelArg { kind, bytes } } } @@ -72,14 +77,16 @@ impl<'a> RawKernelArg<'a> { /// # Safety /// /// A type implementing this trait must mirror the representation and alignment of the -/// corresponding SYCL kernel argument structure. +/// corresponding SYCL kernel argument structure, and must return [`RawKernelArg::Pointer`] +/// exactly when it represents a pointer to USM memory. pub unsafe trait KernelArgument { /// Converts self to its raw representation. /// /// # Safety /// - /// This function returns a reference to raw bytes. These bytes will be passed to FFI - /// functions. The caller must make sure these functions respect Rust's aliasing rules. + /// This function returns a reference to raw bytes, marked as a pointer or plain bytes. + /// These bytes will be passed to FFI functions. The caller must make sure these functions + /// respect Rust's aliasing rules. unsafe fn as_raw_arg(&self) -> RawKernelArg<'_>; } @@ -94,14 +101,16 @@ unsafe impl KernelArgument for T { /// # Safety /// /// A type implementing this trait must mirror the representation and alignment of each -/// corresponding SYCL kernel argument inside the returned array. +/// corresponding SYCL kernel argument inside the returned array, and each entry must be +/// [`RawKernelArg::Pointer`] exactly when it represents a pointer to USM memory. pub unsafe trait KernelArgumentList { /// Converts each struct member to its raw representation. /// /// # Safety /// - /// This function returns references to raw bytes. These bytes will be passed to FFI - /// functions. The caller must make sure these functions respect Rust's aliasing rules. + /// This function returns references to raw bytes, each marked as a pointer or plain bytes. + /// These bytes will be passed to FFI functions. The caller must make sure these functions + /// respect Rust's aliasing rules. unsafe fn as_raw_arg_list(&self) -> [RawKernelArg<'_>; ARGC]; } diff --git a/sycl/sycl-rs/src/usmbox.rs b/sycl/sycl-rs/src/usmbox.rs index e706447..da21d31 100644 --- a/sycl/sycl-rs/src/usmbox.rs +++ b/sycl/sycl-rs/src/usmbox.rs @@ -76,10 +76,10 @@ impl UsmBox { self.len } + /// Returns the bytes of the USM pointer itself, which is what a kernel receives. unsafe fn as_raw_arg_impl(&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_val(&cast_ptr)) } + let data: *const NonNull = &self.data; + unsafe { slice::from_raw_parts(data.cast::(), std::mem::size_of::>()) } } } @@ -170,18 +170,18 @@ pub type DeviceUsmBoxFuture = UsmBoxFuture>; unsafe impl KernelArgument for UsmBox { unsafe fn as_raw_arg(&self) -> RawKernelArg<'_> { - RawKernelArg::Bytes(unsafe { self.as_raw_arg_impl() }) + RawKernelArg::Pointer(unsafe { self.as_raw_arg_impl() }) } } unsafe impl KernelArgument for &UsmBox { unsafe fn as_raw_arg(&self) -> RawKernelArg<'_> { - RawKernelArg::Bytes(unsafe { self.as_raw_arg_impl() }) + RawKernelArg::Pointer(unsafe { self.as_raw_arg_impl() }) } } unsafe impl KernelArgument for &mut UsmBox { unsafe fn as_raw_arg(&self) -> RawKernelArg<'_> { - RawKernelArg::Bytes(unsafe { self.as_raw_arg_impl() }) + RawKernelArg::Pointer(unsafe { self.as_raw_arg_impl() }) } } diff --git a/sycl/sycl-rs/tests/launch.rs b/sycl/sycl-rs/tests/launch.rs new file mode 100644 index 0000000..e27b681 --- /dev/null +++ b/sycl/sycl-rs/tests/launch.rs @@ -0,0 +1,59 @@ +// +// 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 sycl_rs::prelude::*; + +static SAXPY_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 saxpy(float a, const float *x, float *y) { + size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); + y[id] = a * x[id] + y[id]; +} +"#; + +// Mixes a plain value with USM pointers, which some backends (such as the OpenCL CPU device) +// only accept when the pointers are passed through their own entry point. +#[test] +fn launch_with_value_and_pointer_arguments() -> sycl_rs::Result<()> { + const LEN: usize = 64; + // Queue::new() picks a GPU when there is one. Prefer the CPU, whose OpenCL runtime rejects + // pointers passed as plain bytes. + let mut queue = Device::get_devices() + .iter() + .find(|device| device.get_info::() == info::DeviceType::Cpu) + .map_or_else(Queue::new, Queue::from); + let mut x = queue.alloc_shared::(LEN)?.wait()?; + let mut y = queue.alloc_shared::(LEN)?.wait()?; + for (index, (x, y)) in x.iter_mut().zip(y.iter_mut()).enumerate() { + *x = index as f32; + *y = 1.0; + } + + let kernel = queue + .get_context() + .create_kernel_bundle_from_source(SAXPY_SRC)? + .build()? + .get_kernel("saxpy")?; + unsafe { + queue.launch( + NdRange::new([LEN as u64], [16]), + &kernel, + (2.0f32, &x, &mut y), + ) + }? + .wait()?; + + let expected: Vec = (0..LEN).map(|index| 2.0 * index as f32 + 1.0).collect(); + assert_eq!(y.as_ref(), expected.as_slice()); + Ok(()) +}