diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc
index a01b9504a5f28..c03fbc7b814bd 100644
--- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc
+++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc
@@ -41,6 +41,12 @@ This extension is written against the SYCL 2020 revision 7 specification.  All
 references below to the "core SYCL specification" or to section numbers in the
 SYCL specification refer to that revision.
 
+The `nd_launch` overloads that take the kernel arguments as a `sycl::span`
+depend on
+version 2 of
+link:../experimental/sycl_ext_oneapi_raw_kernel_arg.asciidoc[sycl_ext_oneapi_raw_kernel_arg],
+which is what lets an element of that sequence represent a pointer argument.
+
 
 == Status
 
@@ -165,8 +171,13 @@ implementation supports.
 |Description
 
 |1
-|The APIs of this experimental extension are not versioned, so the
- feature-test macro always has this value.
+|Initial version of this extension.
+
+|2
+|Adds the `nd_launch` overloads that take the arguments of a `sycl::kernel` as
+ a `sycl::span` of `raw_kernel_arg`, and specifies that a parameter pack
+ overload handed such a sequence as its only argument passes it on as the
+ argument list.
 |===
 
 
@@ -525,6 +536,51 @@ a!
 ----
 namespace sycl::ext::oneapi::experimental {
 
+template <int Dimensions>
+void nd_launch(sycl::queue q, sycl::nd_range<Dimensions> r,
+               const sycl::kernel& k, sycl::span<const raw_kernel_arg> args);
+
+template <int Dimensions>
+void nd_launch(sycl::handler &h, sycl::nd_range<Dimensions> r,
+               const sycl::kernel& k, sycl::span<const raw_kernel_arg> args);
+
+}
+----
+!====
+_Effects_: Enqueues a kernel object to the `sycl::queue` or `sycl::handler`
+as a basic kernel, using the number of work-items specified by a
+`sycl::nd_range`. The element `args[i]` is passed to the kernel as the argument
+at index `i`, as if it had been passed to `set_arg`. The sequence referenced by
+`args` need only remain valid until the function returns.
+
+If one of the parameter pack overloads above is called with a single argument
+whose type is convertible to `sycl::span<const raw_kernel_arg>`, that argument
+is the kernel argument list and the effects are those of the overload taking a
+`sycl::span`.
+
+[_Note:_ These overloads exist for an application whose kernel argument list is
+only known at run time, which would otherwise have to instantiate a parameter
+pack overload for each number of arguments it may encounter. An application
+holding its arguments in a `std::vector<raw_kernel_arg>`, in a `std::array` of
+them or in a `std::span` of them can therefore pass the container itself, since
+all of those convert to a `sycl::span`. A single `raw_kernel_arg` still selects
+the parameter pack overload, as it is one argument rather than a sequence of
+them. _end note_]
+
+[_Note:_ An element that represents a pointer argument has to be constructed
+with the `pointer_arg` overload of `raw_kernel_arg`, as described in
+link:../experimental/sycl_ext_oneapi_raw_kernel_arg.asciidoc[sycl_ext_oneapi_raw_kernel_arg],
+since a `raw_kernel_arg` holding the bytes of a pointer is only guaranteed to
+bind that pointer on the Level Zero backend. _end note_]
+
+a|
+[frame=all,grid=none]
+!====
+a!
+[source,c++]
+----
+namespace sycl::ext::oneapi::experimental {
+
 template <int Dimensions, typename Properties, typename... Args>
 void nd_launch(sycl::queue q,
                launch_config<sycl::nd_range<Dimensions>, Properties> c,
diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_raw_kernel_arg.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_raw_kernel_arg.asciidoc
index 6226aa1ec8be9..6826c2c9b719f 100644
--- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_raw_kernel_arg.asciidoc
+++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_raw_kernel_arg.asciidoc
@@ -54,11 +54,17 @@ specification.*
 
 == Backend support status
 
-This extension is currently implemented in {dpcpp} only for GPU devices and
-only when using the Level Zero backend.  Attempting to use this extension in
-kernels that run on other devices or backends may result in undefined
-behavior.  Be aware that the compiler is not able to issue a diagnostic to
-warn you if this happens.
+This extension is currently implemented in {dpcpp} for the Level Zero, OpenCL,
+CUDA and Native CPU backends.  Attempting to use this extension on other
+backends may result in undefined behavior.  Be aware that the compiler is not
+able to issue a diagnostic to warn you if this happens.
+
+A kernel argument that is a pointer must be constructed with the `pointer_arg`
+overload described below.  Passing the byte representation of a pointer to the
+byte overload is only guaranteed to bind that pointer on the Level Zero backend,
+because a backend may take a pointer argument through a different entry point
+than the one that takes plain bytes, and it cannot tell the two apart from the
+bytes alone.
 
 
 == Overview
@@ -85,6 +91,19 @@ h.set_arg(1, sycl::ext::oneapi::experimental::raw_kernel_arg(opaque_type, nbytes
 h.parallel_for(range, kernel);
 ----
 
+An argument that is a pointer says so, since a backend may bind a pointer
+through a different entry point than a sequence of bytes:
+
+[source,c++]
+----
+namespace syclex = sycl::ext::oneapi::experimental;
+
+int* ptr = sycl::malloc_device<int>(n, q);
+...
+h.set_arg(0, syclex::raw_kernel_arg(&ptr, syclex::pointer_arg));
+h.parallel_for(range, kernel);
+----
+
 
 == Specification
 
@@ -104,8 +123,10 @@ implementation supports.
 |Description
 
 |1
-|The APIs of this experimental extension are not versioned, so the
- feature-test macro always has this value.
+|Initial version of this extension.
+
+|2
+|Adds the `raw_kernel_arg` constructor that takes a pointer argument.
 |===
 
 === The `raw_kernel_arg` class
@@ -117,9 +138,15 @@ kernel arguments via a raw byte representation.
 ----
 namespace sycl::ext::oneapi::experimental {
 
+  struct pointer_arg_t {};
+  inline constexpr pointer_arg_t pointer_arg{};
+
   class raw_kernel_arg {
   public:
     raw_kernel_arg(const void* bytes, size_t count);
+
+    template <typename T>
+    raw_kernel_arg(T* const* pointer_location, pointer_arg_t tag);
   };
 
 } // namespace sycl::ext::oneapi::experimental
@@ -131,13 +158,34 @@ raw_kernel_arg(const void* bytes, size_t count);
 ----
 _Preconditions_: `bytes` must point to an array of at least `count` bytes,
 which is the byte representation of a kernel argument that is trivially
-copyable.
+copyable.  If the argument is a pointer, only the Level Zero backend is
+guaranteed to bind it; see _Backend support status_ above.
 
 _Effects_: Constructs a `raw_kernel_arg` representing a view of the `count`
 bytes starting at the address specified by `bytes`. Since the `raw_kernel_arg`
 object is only a view, the caller must ensure that the lifetime of the `bytes`
 memory lasts at least as long as the lifetime of the `raw_kernel_arg` object.
 
+[source,c++]
+----
+template <typename T>
+raw_kernel_arg(T* const* pointer_location, pointer_arg_t tag);
+----
+_Preconditions_: `pointer_location` must point to a pointer that is a valid
+kernel argument, such as a pointer to memory allocated by one of the USM
+allocation functions.
+
+_Effects_: Constructs a `raw_kernel_arg` representing a view of the pointer
+stored at `pointer_location`, to be bound as a pointer argument rather than as
+the `sizeof(T*)` bytes it is made of. Since the `raw_kernel_arg` object is only
+a view, the caller must ensure that the lifetime of the pointer object at
+`pointer_location` lasts at least as long as the lifetime of the
+`raw_kernel_arg` object.
+
+[_Note:_ The constructor takes the address of the pointer, in keeping with the
+byte overload taking the address of the bytes, so that passing the pointer
+itself does not compile. _{endnote}_]
+
 === Using a raw kernel argument
 
 Instances of `raw_kernel_arg` are passed to kernels via the existing `set_arg`
@@ -154,7 +202,8 @@ argument in `args` was passed to `set_arg` ", adding a new overload of
 void set_arg(int argIndex, sycl::ext::oneapi::experimental::raw_kernel_arg&& arg);
 ----
 _Effects_: Sets the kernel argument associated with index `argIndex` using the
-bytes represented by `arg`.
+bytes represented by `arg`, or as a pointer argument if `arg` was constructed
+with the `pointer_arg` overload.
 
 
 == Issues
diff --git a/sycl/include/sycl/detail/kernel_arg_view.hpp b/sycl/include/sycl/detail/kernel_arg_view.hpp
new file mode 100644
index 0000000000000..6062e6ed2ec1d
--- /dev/null
+++ b/sycl/include/sycl/detail/kernel_arg_view.hpp
@@ -0,0 +1,35 @@
+//==---- kernel_arg_view.hpp --- SYCL kernel argument as bytes and kind ----==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <sycl/detail/kernel_desc.hpp> // for kernel_param_kind_t
+
+#include <stddef.h> // for size_t
+
+namespace sycl {
+inline namespace _V1 {
+namespace detail {
+
+inline namespace kernel_arg_view_v1 {
+
+// A kernel argument reduced to what the runtime needs in order to bind it. Used
+// by the enqueue functions that launch a `sycl::kernel` without building a
+// command group, where the arguments are only known as bytes plus a kind. This
+// is being passed across the ABI boundary, hence the versioned namespace.
+struct KernelArgView {
+  const void *MPtr;
+  size_t MSize;
+  kernel_param_kind_t MKind;
+};
+
+} // namespace kernel_arg_view_v1
+
+} // namespace detail
+} // namespace _V1
+} // namespace sycl
diff --git a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp
index efc355875b6a5..c3520ef60795b 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp
+++ b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp
@@ -15,6 +15,7 @@
 #include <sycl/ext/oneapi/experimental/enqueue_types.hpp>
 #include <sycl/ext/oneapi/experimental/free_function_traits.hpp>
 #include <sycl/ext/oneapi/experimental/graph.hpp>
+#include <sycl/ext/oneapi/experimental/raw_kernel_arg.hpp>
 #include <sycl/ext/oneapi/properties.hpp>
 #include <sycl/handler.hpp>
 #include <sycl/nd_range.hpp>
@@ -98,6 +99,52 @@ template <typename LCRangeT, typename LCPropertiesT> struct LaunchConfigAccess {
   }
 };
 
+// The argument type as the kernel sees it. Deliberately not `std::decay_t`,
+// which turns an array into a pointer: an array has to keep being bound as the
+// bytes it is, which is what `handler::setArgHelper` does with it.
+template <typename T>
+using plain_arg_t = std::remove_cv_t<std::remove_reference_t<T>>;
+
+// An argument that can be bound as plain bytes, i.e. one that carries no
+// requirement for the scheduler to track. Accessors, local accessors, streams
+// and work group memory are deliberately excluded and keep using the command
+// group path; `HasSpecialCaptures` in the runtime draws the same line.
+template <typename T>
+inline constexpr bool is_plain_kernel_arg_v =
+    std::is_arithmetic_v<plain_arg_t<T>> || std::is_enum_v<plain_arg_t<T>> ||
+    std::is_pointer_v<plain_arg_t<T>> ||
+    std::is_same_v<plain_arg_t<T>, raw_kernel_arg>;
+
+// A pointer has to keep its kind. The runtime binds a pointer argument as
+// UR_EXP_KERNEL_ARG_TYPE_POINTER, which the OpenCL adapter passes to
+// clSetKernelArgMemPointerINTEL rather than to clSetKernelArg, so plain bytes
+// are not a substitute. The Native CPU adapter draws the same distinction: it
+// puts a pointer argument straight into the argument slot, whereas a value
+// argument lands there as the address of the adapter's own copy.
+template <typename T>
+sycl::detail::KernelArgView makeKernelArgView(const T &Arg) {
+  using sycl::detail::kernel_param_kind_t;
+  if constexpr (std::is_same_v<plain_arg_t<T>, raw_kernel_arg>)
+    return {RawKernelArgAccess::getData(Arg), RawKernelArgAccess::getSize(Arg),
+            RawKernelArgAccess::isPointer(Arg)
+                ? kernel_param_kind_t::kind_pointer
+                : kernel_param_kind_t::kind_std_layout};
+  else
+    return {&Arg, sizeof(plain_arg_t<T>),
+            std::is_pointer_v<plain_arg_t<T>>
+                ? kernel_param_kind_t::kind_pointer
+                : kernel_param_kind_t::kind_std_layout};
+}
+
+// True when a parameter pack overload was handed the argument list itself, as a
+// container that converts to the span the sibling overload takes. A pack is an
+// exact match and wins overload resolution, so such a call has to be forwarded
+// rather than bound as one argument.
+template <typename... ArgsT>
+inline constexpr bool is_arg_list_container_v =
+    sizeof...(ArgsT) == 1 &&
+    (std::is_convertible_v<ArgsT, sycl::span<const raw_kernel_arg>> && ...);
+
 template <typename CommandGroupFunc, typename PropertiesT>
 void submit_impl(const queue &Q, PropertiesT Props, CommandGroupFunc &&CGF,
                  const sycl::detail::code_location &CodeLoc) {
@@ -398,19 +445,82 @@ void nd_launch(queue Q, launch_config<nd_range<Dimensions>, Properties> Config,
   }
 }
 
+template <int Dimensions>
+void nd_launch(handler &CGH, nd_range<Dimensions> Range,
+               const kernel &KernelObj, sycl::span<const raw_kernel_arg> Args);
+
 template <int Dimensions, typename... ArgsT>
 void nd_launch(handler &CGH, nd_range<Dimensions> Range,
                const kernel &KernelObj, ArgsT &&...Args) {
-  CGH.set_args<ArgsT...>(std::forward<ArgsT>(Args)...);
-  CGH.parallel_for(Range, KernelObj);
+  if constexpr (detail::is_arg_list_container_v<ArgsT...>) {
+    nd_launch(CGH, Range, KernelObj,
+              sycl::span<const raw_kernel_arg>{std::forward<ArgsT>(Args)...});
+  } else {
+    CGH.set_args<ArgsT...>(std::forward<ArgsT>(Args)...);
+    CGH.parallel_for(Range, KernelObj);
+  }
 }
 
+template <int Dimensions>
+void nd_launch(queue Q, nd_range<Dimensions> Range, const kernel &KernelObj,
+               sycl::span<const raw_kernel_arg> Args,
+               const sycl::detail::code_location &CodeLoc =
+                   sycl::detail::code_location::current());
+
 template <int Dimensions, typename... ArgsT>
 void nd_launch(queue Q, nd_range<Dimensions> Range, const kernel &KernelObj,
                ArgsT &&...Args) {
-  submit(std::move(Q), [&](handler &CGH) {
-    nd_launch(CGH, Range, KernelObj, std::forward<ArgsT>(Args)...);
-  });
+  // A container of raw_kernel_arg is the argument list, not one argument, and
+  // the pack is what overload resolution picks for it, so hand it over to the
+  // overload that takes a span.
+  if constexpr (detail::is_arg_list_container_v<ArgsT...>) {
+    nd_launch(std::move(Q), Range, KernelObj,
+              sycl::span<const raw_kernel_arg>{std::forward<ArgsT>(Args)...});
+  } else if constexpr ((detail::is_plain_kernel_arg_v<ArgsT> && ...)) {
+    // Bind the arguments straight from this call, so that neither a handler nor
+    // a command group object has to be created. The array is one element longer
+    // than the pack so that a zero-argument kernel stays well formed.
+    const sycl::detail::KernelArgView ArgViews[sizeof...(ArgsT) + 1] = {
+        detail::makeKernelArgView(Args)...};
+    // An overload that ends in a parameter pack cannot take a trailing
+    // code_location parameter, so the location is the one this header sees,
+    // as it was when this overload went through submit(). Seed the TLS slot
+    // rather than leaving it default-constructed, which the instrumentation
+    // reads as a null file and function name.
+    sycl::detail::tls_code_loc_t TlsCodeLocCapture{
+        sycl::detail::code_location::current()};
+    sycl::submit_kernel_obj_direct_without_event_impl(
+        Q, sycl::detail::nd_range_view(Range), KernelObj,
+        {ArgViews, sizeof...(ArgsT)}, TlsCodeLocCapture.query(),
+        TlsCodeLocCapture.isToplevel());
+  } else {
+    submit(std::move(Q), [&](handler &CGH) {
+      nd_launch(CGH, Range, KernelObj, std::forward<ArgsT>(Args)...);
+    });
+  }
+}
+
+template <int Dimensions>
+void nd_launch(handler &CGH, nd_range<Dimensions> Range,
+               const kernel &KernelObj, sycl::span<const raw_kernel_arg> Args) {
+  // set_arg only takes an rvalue raw_kernel_arg; an lvalue would select the
+  // generic overload and bind the object itself as the argument.
+  for (size_t I = 0; I < Args.size(); ++I)
+    CGH.set_arg(static_cast<int>(I), raw_kernel_arg{Args[I]});
+  CGH.parallel_for(Range, KernelObj);
+}
+
+// Takes the kernel arguments as a contiguous sequence instead of a parameter
+// pack, for a caller that only learns its argument list at run time and would
+// otherwise need one instantiation of the pack overload per argument count.
+template <int Dimensions>
+void nd_launch(queue Q, nd_range<Dimensions> Range, const kernel &KernelObj,
+               sycl::span<const raw_kernel_arg> Args,
+               const sycl::detail::code_location &CodeLoc) {
+  sycl::detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
+  sycl::submit_kernel_obj_direct_without_event_impl(
+      Q, sycl::detail::nd_range_view(Range), KernelObj, Args,
+      TlsCodeLocCapture.query(), TlsCodeLocCapture.isToplevel());
 }
 
 template <int Dimensions, typename Properties, typename... ArgsT>
diff --git a/sycl/include/sycl/ext/oneapi/experimental/raw_kernel_arg.hpp b/sycl/include/sycl/ext/oneapi/experimental/raw_kernel_arg.hpp
index 4ca41f315a43f..933a3550a1fa3 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/raw_kernel_arg.hpp
+++ b/sycl/include/sycl/ext/oneapi/experimental/raw_kernel_arg.hpp
@@ -18,21 +18,51 @@ namespace ext::oneapi::experimental {
 
 namespace detail {
 class dynamic_parameter_impl;
+struct RawKernelArgAccess;
 } // namespace detail
 
+// Tells the raw_kernel_arg constructor below that the argument it is given is a
+// pointer rather than a sequence of bytes to copy.
+struct pointer_arg_t {};
+inline constexpr pointer_arg_t pointer_arg{};
+
 class raw_kernel_arg {
 public:
   raw_kernel_arg(const void *bytes, size_t count)
       : MArgData(bytes), MArgSize(count) {}
 
+  // A pointer argument is not interchangeable with the bytes it is made of: a
+  // backend may have to be told that an argument is a pointer to bind it at
+  // all, as OpenCL does, where a USM pointer goes to
+  // clSetKernelArgMemPointerINTEL rather than to clSetKernelArg. Takes the
+  // address of the pointer, like the byte form takes the address of the bytes,
+  // so that passing the pointer itself does not compile.
+  template <typename T>
+  raw_kernel_arg(T *const *pointer_location, pointer_arg_t)
+      : MArgData(pointer_location), MArgSize(sizeof(T *)), MIsPointer(true) {}
+
 private:
   const void *MArgData;
   size_t MArgSize;
+  // Appended last, so that the offsets of the members above, which the library
+  // reads on the paths that predate this one, stay where they were.
+  bool MIsPointer = false;
 
   friend class sycl::handler;
   // For sycl_ext_oneapi_graph integration
   friend class detail::dynamic_parameter_impl;
+  // For the enqueue paths that bind arguments without a handler
+  friend struct detail::RawKernelArgAccess;
+};
+
+namespace detail {
+// Helper for accessing the members of raw_kernel_arg.
+struct RawKernelArgAccess {
+  static const void *getData(const raw_kernel_arg &Arg) { return Arg.MArgData; }
+  static size_t getSize(const raw_kernel_arg &Arg) { return Arg.MArgSize; }
+  static bool isPointer(const raw_kernel_arg &Arg) { return Arg.MIsPointer; }
 };
+} // namespace detail
 
 } // namespace ext::oneapi::experimental
 } // namespace _V1
diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp
index 8354570fd9054..845a3ed0d0948 100644
--- a/sycl/include/sycl/handler.hpp
+++ b/sycl/include/sycl/handler.hpp
@@ -603,8 +603,11 @@ class __SYCL_EXPORT handler {
   void setArgHelper(int ArgIndex,
                     sycl::ext::oneapi::experimental::raw_kernel_arg &&Arg) {
     auto StoredArg = storeRawArg(Arg);
-    addArg(detail::kernel_param_kind_t::kind_std_layout, StoredArg,
-           Arg.MArgSize, ArgIndex);
+    // A pointer argument has to be bound as one: a backend may reach for it
+    // through a different entry point than the one that takes plain bytes.
+    addArg(Arg.MIsPointer ? detail::kernel_param_kind_t::kind_pointer
+                          : detail::kernel_param_kind_t::kind_std_layout,
+           StoredArg, Arg.MArgSize, ArgIndex);
   }
 
   /// Registers a dynamic parameter with the handler for later association with
diff --git a/sycl/include/sycl/queue.hpp b/sycl/include/sycl/queue.hpp
index 629b73e8258f3..5dff032bbdcb8 100644
--- a/sycl/include/sycl/queue.hpp
+++ b/sycl/include/sycl/queue.hpp
@@ -18,6 +18,7 @@
 #include <sycl/detail/common.hpp>             // for code_location
 #include <sycl/detail/defines_elementary.hpp> // for __SYCL2020_DEP...
 #include <sycl/detail/export.hpp>             // for __SYCL_EXPORT
+#include <sycl/detail/kernel_arg_view.hpp>    // for KernelArgView
 #include <sycl/detail/kernel_desc.hpp>        // for KernelInfo
 #include <sycl/detail/nd_range_view.hpp>
 #include <sycl/detail/optional.hpp>
@@ -59,12 +60,34 @@ inline namespace _V1 {
 class context;
 class device;
 class event;
+class kernel;
 class queue;
 
 template <backend BackendName, class SyclObjectT>
 auto get_native(const SyclObjectT &Obj)
     -> backend_return_t<BackendName, SyclObjectT>;
 
+namespace ext::oneapi::experimental {
+class raw_kernel_arg;
+} // namespace ext::oneapi::experimental
+
+// Launches an already built `sycl::kernel` with an explicit argument list,
+// bypassing the handler and the command group object the same way
+// submit_kernel_direct_* does for kernel function objects.
+void __SYCL_EXPORT submit_kernel_obj_direct_without_event_impl(
+    const queue &Queue, const detail::nd_range_view &RangeView,
+    const kernel &Kernel, sycl::span<const detail::KernelArgView> Args,
+    const detail::code_location &CodeLoc, bool IsTopCodeLoc);
+
+// As above, for an argument list that is already a contiguous sequence of
+// `raw_kernel_arg`, which carries its own kind per element, so the caller needs
+// no conversion step.
+void __SYCL_EXPORT submit_kernel_obj_direct_without_event_impl(
+    const queue &Queue, const detail::nd_range_view &RangeView,
+    const kernel &Kernel,
+    sycl::span<const ext::oneapi::experimental::raw_kernel_arg> Args,
+    const detail::code_location &CodeLoc, bool IsTopCodeLoc);
+
 event __SYCL_EXPORT submit_kernel_direct_with_event_impl(
     const queue &Queue, const detail::nd_range_view &RangeView,
     detail::HostKernelRefBase &HostKernel,
diff --git a/sycl/source/detail/queue_impl.cpp b/sycl/source/detail/queue_impl.cpp
index a8c51a03b710e..19b76b670f44d 100644
--- a/sycl/source/detail/queue_impl.cpp
+++ b/sycl/source/detail/queue_impl.cpp
@@ -883,6 +883,128 @@ EventImplPtr queue_impl::submit_kernel_direct_impl(
                        /*InsertBarrierForInOrderCommand*/ false);
 }
 
+namespace {
+// The two argument forms differ only in how one element yields the bytes to
+// bind and their kind: a raw_kernel_arg carries plain bytes unless it was built
+// as a pointer argument.
+inline sycl::detail::KernelArgView
+makeKernelArgView(const sycl::detail::KernelArgView &Arg) {
+  return Arg;
+}
+inline sycl::detail::KernelArgView
+makeKernelArgView(const ext::oneapi::experimental::raw_kernel_arg &Arg) {
+  namespace syclex_detail = ext::oneapi::experimental::detail;
+  return {syclex_detail::RawKernelArgAccess::getData(Arg),
+          syclex_detail::RawKernelArgAccess::getSize(Arg),
+          syclex_detail::RawKernelArgAccess::isPointer(Arg)
+              ? sycl::detail::kernel_param_kind_t::kind_pointer
+              : sycl::detail::kernel_param_kind_t::kind_std_layout};
+}
+} // namespace
+
+template <typename ArgT>
+void queue_impl::submit_kernel_obj_direct_impl(
+    const detail::nd_range_view &RangeView,
+    const std::shared_ptr<detail::kernel_impl> &KernelImpl,
+    sycl::span<const ArgT> Args, const detail::code_location &CodeLoc,
+    bool IsTopCodeLoc) {
+
+  KernelData KData;
+  KData.setDeviceKernelInfoPtr(&KernelImpl->getDeviceKernelInfo());
+  KData.setNDRDesc(NDRDescT(RangeView));
+  KData.getArgs().reserve(Args.size());
+
+  // The kernel may have come from a bundle, and the bundle has to travel with
+  // it the way the handler path lets it: it is what lets enqueueImpKernel
+  // initialize the device globals a bundle keeps to itself.
+  std::shared_ptr<detail::kernel_bundle_impl> KernelBundleImpl =
+      KernelImpl->get_kernel_bundle();
+
+  // This overload carries no properties, so a kernel that needs work group
+  // scratch memory can never have been given a size. The handler path reports
+  // that in handler.cpp, so report it here too rather than launching a kernel
+  // whose scratch allocation is missing.
+  if (KData.getDeviceKernelInfoPtr()->getWorkGroupDynamicLocalMem())
+    throw sycl::exception(
+        sycl::make_error_code(sycl::errc::memory_allocation),
+        "Kernel allocates work group scratch memory but an allocation size "
+        "has not been specified through the work_group_scratch_size property!");
+
+  auto SubmitKernelFunc = [&](detail::CG::StorageInitHelper &&CGData)
+      -> std::pair<EventImplPtr, bool> {
+    bool SchedulerBypass =
+        (CGData.MEvents.size() > 0
+             ? detail::Scheduler::areEventsSafeForSchedulerBypass(
+                   CGData.MEvents, getContextImpl())
+             : true) &&
+        !hasCommandGraph();
+
+    // On the bypass path the argument values are read before this call returns,
+    // so they can be bound where the caller keeps them. Otherwise the command
+    // group outlives the call and they have to be copied into its storage.
+    for (size_t I = 0; I < Args.size(); ++I) {
+      const sycl::detail::KernelArgView View = makeKernelArgView(Args[I]);
+      void *Value = const_cast<void *>(View.MPtr);
+      if (!SchedulerBypass) {
+        const char *Bytes = static_cast<const char *>(View.MPtr);
+        CGData.MArgsStorage.emplace_back(Bytes, Bytes + View.MSize);
+        Value = CGData.MArgsStorage.back().data();
+      }
+      KData.addArg(View.MKind, Value, static_cast<int>(View.MSize),
+                   static_cast<int>(I));
+    }
+
+    if (SchedulerBypass)
+      return {submit_kernel_scheduler_bypass(
+                  KData, CGData.MEvents, /*EventNeeded*/ false,
+                  KernelImpl.get(), KernelBundleImpl.get(), CodeLoc,
+                  IsTopCodeLoc),
+              /*SchedulerBypass*/ true};
+
+    auto CommandGroup = std::make_unique<detail::CGExecKernel>(
+        KData.getNDRDesc(), /*HostKernel*/ nullptr, KernelImpl,
+        KernelBundleImpl, std::move(CGData), std::move(KData).getArgs(),
+        *KData.getDeviceKernelInfoPtr(),
+        std::vector<std::shared_ptr<detail::stream_impl>>{},
+        std::vector<std::shared_ptr<const void>>{}, detail::CGType::Kernel,
+        KData.getKernelCacheConfig(), KData.isCooperative(),
+        KData.usesClusterLaunch(), KData.getKernelWorkGroupMemorySize(),
+        CodeLoc);
+    CommandGroup->MIsTopCodeLoc = IsTopCodeLoc;
+
+    if (auto GraphImpl = getCommandGraph(); GraphImpl)
+      return {submit_command_to_graph(*GraphImpl, std::move(CommandGroup),
+                                      detail::CGType::Kernel),
+              /*SchedulerBypass*/ false};
+
+    return {detail::Scheduler::getInstance().addCG(std::move(CommandGroup),
+                                                   *this, true),
+            /*SchedulerBypass*/ false};
+  };
+
+  submit_direct(/*CallerNeedsEvent*/ false, /*DepEvents*/ {}, SubmitKernelFunc,
+                detail::CGType::Kernel,
+                /*InsertBarrierForInOrderCommand*/ false);
+}
+
+void queue_impl::submit_kernel_obj_direct_without_event(
+    const detail::nd_range_view &RangeView,
+    const std::shared_ptr<detail::kernel_impl> &KernelImpl,
+    sycl::span<const sycl::detail::KernelArgView> Args,
+    const detail::code_location &CodeLoc, bool IsTopCodeLoc) {
+  submit_kernel_obj_direct_impl(RangeView, KernelImpl, Args, CodeLoc,
+                                IsTopCodeLoc);
+}
+
+void queue_impl::submit_kernel_obj_direct_without_event(
+    const detail::nd_range_view &RangeView,
+    const std::shared_ptr<detail::kernel_impl> &KernelImpl,
+    sycl::span<const ext::oneapi::experimental::raw_kernel_arg> Args,
+    const detail::code_location &CodeLoc, bool IsTopCodeLoc) {
+  submit_kernel_obj_direct_impl(RangeView, KernelImpl, Args, CodeLoc,
+                                IsTopCodeLoc);
+}
+
 EventImplPtr queue_impl::submit_graph_direct_impl(
     std::shared_ptr<ext::oneapi::experimental::detail::exec_graph_impl>
         ExecGraph,
diff --git a/sycl/source/detail/queue_impl.hpp b/sycl/source/detail/queue_impl.hpp
index 2128ff4be7b0d..e62dc12b9ec76 100644
--- a/sycl/source/detail/queue_impl.hpp
+++ b/sycl/source/detail/queue_impl.hpp
@@ -378,6 +378,26 @@ class queue_impl : public std::enable_shared_from_this<queue_impl> {
                               CodeLoc, IsTopCodeLoc);
   }
 
+  /// Submits an already built kernel with an explicit argument list, without
+  /// creating a handler or a command group object.
+  ///
+  /// \param RangeView is the execution range.
+  /// \param KernelImpl is the kernel to launch.
+  /// \param Args are the kernel arguments, as bytes plus their kind.
+  void submit_kernel_obj_direct_without_event(
+      const detail::nd_range_view &RangeView,
+      const std::shared_ptr<detail::kernel_impl> &KernelImpl,
+      sycl::span<const sycl::detail::KernelArgView> Args,
+      const detail::code_location &CodeLoc, bool IsTopCodeLoc);
+
+  /// As above, for arguments that are already a sequence of `raw_kernel_arg`,
+  /// each of which carries its own kind.
+  void submit_kernel_obj_direct_without_event(
+      const detail::nd_range_view &RangeView,
+      const std::shared_ptr<detail::kernel_impl> &KernelImpl,
+      sycl::span<const ext::oneapi::experimental::raw_kernel_arg> Args,
+      const detail::code_location &CodeLoc, bool IsTopCodeLoc);
+
   event submit_barrier_direct_with_event(sycl::span<const event> DepEvents,
                                          detail::CGType BarrierType,
                                          const detail::code_location &CodeLoc) {
@@ -986,6 +1006,17 @@ class queue_impl : public std::enable_shared_from_this<queue_impl> {
                 SubmitCommandFuncType &SubmitCommandFunc, detail::CGType Type,
                 bool InsertBarrierForInOrderCommand);
 
+  /// Shared implementation of the two submit_kernel_obj_direct_without_event
+  /// overloads. The element type only decides how one argument yields its bytes
+  /// and kind, so the submission itself is written once. Instantiated for
+  /// KernelArgView and for raw_kernel_arg in queue_impl.cpp.
+  template <typename ArgT>
+  void submit_kernel_obj_direct_impl(
+      const detail::nd_range_view &RangeView,
+      const std::shared_ptr<detail::kernel_impl> &KernelImpl,
+      sycl::span<const ArgT> Args, const detail::code_location &CodeLoc,
+      bool IsTopCodeLoc);
+
   /// Performs barrier submission to the queue.
   ///
   /// \param DepEvents is a vector of dependencies of the operation.
diff --git a/sycl/source/feature_test.hpp.in b/sycl/source/feature_test.hpp.in
index 6aebc773b514c..a8d7396ec122d 100644
--- a/sycl/source/feature_test.hpp.in
+++ b/sycl/source/feature_test.hpp.in
@@ -94,8 +94,8 @@ inline namespace _V1 {
 #define SYCL_EXT_ONEAPI_FORWARD_PROGRESS 1
 #define SYCL_EXT_ONEAPI_FREE_FUNCTION_KERNELS 1
 #define SYCL_EXT_ONEAPI_PROD 1
-#define SYCL_EXT_ONEAPI_ENQUEUE_FUNCTIONS 1
-#define SYCL_EXT_ONEAPI_RAW_KERNEL_ARG 1
+#define SYCL_EXT_ONEAPI_ENQUEUE_FUNCTIONS 2
+#define SYCL_EXT_ONEAPI_RAW_KERNEL_ARG 2
 #define SYCL_EXT_ONEAPI_PROFILING_TAG 1
 #define SYCL_EXT_ONEAPI_ENQUEUE_NATIVE_COMMAND 2
 #define SYCL_EXT_ONEAPI_GET_KERNEL_INFO 1
diff --git a/sycl/source/queue.cpp b/sycl/source/queue.cpp
index 91e5aa5f4f82a..2e1bccd24cb45 100644
--- a/sycl/source/queue.cpp
+++ b/sycl/source/queue.cpp
@@ -13,6 +13,7 @@
 #include <sycl/event.hpp>
 #include <sycl/exception_list.hpp>
 #include <sycl/handler.hpp>
+#include <sycl/kernel.hpp>
 #include <sycl/queue.hpp>
 
 #include <algorithm>
@@ -318,6 +319,23 @@ void submit_kernel_direct_without_event_impl(
       IsTopCodeLoc);
 }
 
+void submit_kernel_obj_direct_without_event_impl(
+    const queue &Queue, const detail::nd_range_view &RangeView,
+    const kernel &Kernel, sycl::span<const detail::KernelArgView> Args,
+    const detail::code_location &CodeLoc, bool IsTopCodeLoc) {
+  detail::getSyclObjImpl(Queue)->submit_kernel_obj_direct_without_event(
+      RangeView, detail::getSyclObjImpl(Kernel), Args, CodeLoc, IsTopCodeLoc);
+}
+
+void submit_kernel_obj_direct_without_event_impl(
+    const queue &Queue, const detail::nd_range_view &RangeView,
+    const kernel &Kernel,
+    sycl::span<const ext::oneapi::experimental::raw_kernel_arg> Args,
+    const detail::code_location &CodeLoc, bool IsTopCodeLoc) {
+  detail::getSyclObjImpl(Queue)->submit_kernel_obj_direct_without_event(
+      RangeView, detail::getSyclObjImpl(Kernel), Args, CodeLoc, IsTopCodeLoc);
+}
+
 event submit_graph_direct_with_event_impl(
     const queue &Queue,
     ext::oneapi::experimental::command_graph<
diff --git a/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_array_arg.cpp b/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_array_arg.cpp
new file mode 100644
index 0000000000000..44caf806aab04
--- /dev/null
+++ b/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_array_arg.cpp
@@ -0,0 +1,78 @@
+// REQUIRES: aspect-usm_shared_allocations
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+// Tests that an array argument reaches a `sycl::kernel` the same way through
+// the handler-less nd_launch overload as it does through the handler.
+//
+// An array is bound as the bytes it is, which is what `handler::set_args` does
+// with it, so the kernel below reads those 16 bytes as its parameter. The
+// handler-less overload must not classify the array as a pointer argument: the
+// runtime would then read the first bytes of the array as an address, which
+// binds neither the bytes nor the array itself.
+//
+// Kernel arguments are sticky on the backend side and a bundle hands out the
+// same underlying kernel for the same kernel id, so a correct launch would mask
+// a wrong one. The handler-less path therefore goes first.
+
+#include <sycl/detail/core.hpp>
+#include <sycl/ext/oneapi/experimental/enqueue_functions.hpp>
+#include <sycl/ext/oneapi/free_function_queries.hpp>
+#include <sycl/kernel_bundle.hpp>
+#include <sycl/properties/all_properties.hpp>
+#include <sycl/usm.hpp>
+
+#include "common.hpp"
+
+namespace syclext = sycl::ext::oneapi;
+
+constexpr size_t N = 8;
+
+struct FourInts {
+  int A, B, C, D;
+};
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((oneapiext::nd_range_kernel<1>))
+void addFourInts(FourInts Values, int *Out) {
+  size_t I = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
+  Out[I] = Values.A + Values.B + Values.C + Values.D;
+}
+
+template <auto *Func> sycl::kernel getKernel(sycl::queue &Q) {
+  auto Bundle =
+      oneapiext::get_kernel_bundle<Func, sycl::bundle_state::executable>(
+          Q.get_context());
+  return Bundle.template ext_oneapi_get_kernel<Func>();
+}
+
+int main() {
+  sycl::queue Q{sycl::property::queue::in_order{}};
+
+  int *Out = sycl::malloc_shared<int>(N, Q);
+  sycl::nd_range<1> Ndr{sycl::range<1>{N}, sycl::range<1>{N}};
+  int Values[4] = {1, 2, 3, 4};
+  constexpr int Expected = 1 + 2 + 3 + 4;
+
+  int Failed = 0;
+
+  sycl::kernel Kernel = getKernel<addFourInts>(Q);
+
+  // The handler-less path, which must bind the array as plain bytes.
+  Q.memset(Out, 0, N * sizeof(int));
+  oneapiext::nd_launch(Q, Ndr, Kernel, Values, Out);
+  Q.wait();
+  for (size_t I = 0; I < N; ++I)
+    Failed += Check(Out, Expected, I, "array without a handler");
+
+  // The command group path has to agree with it.
+  Q.memset(Out, 0, N * sizeof(int));
+  Q.submit([&](sycl::handler &CGH) {
+    oneapiext::nd_launch(CGH, Ndr, Kernel, Values, Out);
+  });
+  Q.wait();
+  for (size_t I = 0; I < N; ++I)
+    Failed += Check(Out, Expected, I, "array through the command group path");
+
+  sycl::free(Out, Q);
+  return Failed;
+}
diff --git a/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_direct.cpp b/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_direct.cpp
new file mode 100644
index 0000000000000..299023f94b896
--- /dev/null
+++ b/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_direct.cpp
@@ -0,0 +1,148 @@
+// REQUIRES: aspect-usm_shared_allocations
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+// Tests the nd_launch overloads that take an already built sycl::kernel and its
+// arguments as a parameter pack, which bypass the handler when every argument
+// can be bound as plain bytes. Arguments that the scheduler has to track, an
+// accessor here, must still reach the kernel through the command group path.
+
+#include <sycl/detail/core.hpp>
+#include <sycl/ext/oneapi/experimental/enqueue_functions.hpp>
+#include <sycl/ext/oneapi/experimental/raw_kernel_arg.hpp>
+#include <sycl/ext/oneapi/free_function_queries.hpp>
+#include <sycl/ext/oneapi/work_group_scratch_memory.hpp>
+#include <sycl/kernel_bundle.hpp>
+#include <sycl/properties/all_properties.hpp>
+#include <sycl/usm.hpp>
+
+#include "common.hpp"
+
+namespace syclext = sycl::ext::oneapi;
+
+constexpr size_t N = 1024;
+constexpr size_t WGSize = 8;
+
+enum class Sign : int { Plus = 1 };
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((oneapiext::nd_range_kernel<1>))
+void addScalars(int *Ptr, int A, int B) {
+  size_t I = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
+  Ptr[I] += A + B;
+}
+
+// A mixture of argument kinds, so that a wrong size or a wrong order shows up
+// as a wrong result rather than as a silent pass.
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((oneapiext::nd_range_kernel<1>))
+void addMixed(int *Ptr, int A, unsigned long B, float C, Sign S) {
+  size_t I = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
+  Ptr[I] +=
+      static_cast<int>(S) * (A + static_cast<int>(B) + static_cast<int>(C));
+}
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((oneapiext::nd_range_kernel<1>))
+void addViaAccessor(sycl::accessor<int, 1> Acc, int A) {
+  size_t I = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
+  Acc[I] += A;
+}
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((oneapiext::nd_range_kernel<1>))
+void usesScratch(int *Ptr) {
+  size_t I = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
+  int *Scratch =
+      reinterpret_cast<int *>(oneapiext::get_work_group_scratch_memory());
+  Scratch[I % WGSize] = static_cast<int>(I);
+  Ptr[I] = Scratch[I % WGSize];
+}
+
+template <auto *Func> sycl::kernel getKernel(sycl::queue &Q) {
+  auto Bundle =
+      oneapiext::get_kernel_bundle<Func, sycl::bundle_state::executable>(
+          Q.get_context());
+  return Bundle.template ext_oneapi_get_kernel<Func>();
+}
+
+int main() {
+  sycl::queue Q{sycl::property::queue::in_order{}};
+  int *Memory = sycl::malloc_shared<int>(N, Q);
+  sycl::nd_range<1> Ndr{sycl::range<1>{N}, sycl::range<1>{WGSize}};
+
+  int Failed = 0;
+
+  // Typed arguments: no handler is created for these.
+  sycl::kernel ScalarsKernel = getKernel<addScalars>(Q);
+  Q.memset(Memory, 0, N * sizeof(int));
+  oneapiext::nd_launch(Q, Ndr, ScalarsKernel, Memory, 3, 4);
+  Q.wait();
+  for (size_t I = 0; I < N; ++I)
+    Failed += Check(Memory, 7, I, "typed arguments");
+
+  // The same scalar arguments as raw bytes, which is how a caller that only
+  // knows the signature as sizes has to pass them. Mixing them with a typed
+  // pointer has to work; nd_launch_kernel_obj_direct_raw_ptr.cpp covers an
+  // argument list that is raw throughout.
+  int A = 10, B = 20;
+  Q.memset(Memory, 0, N * sizeof(int));
+  oneapiext::nd_launch(Q, Ndr, ScalarsKernel, Memory,
+                       oneapiext::raw_kernel_arg{&A, sizeof(A)},
+                       oneapiext::raw_kernel_arg{&B, sizeof(B)});
+  Q.wait();
+  for (size_t I = 0; I < N; ++I)
+    Failed += Check(Memory, 30, I, "raw_kernel_arg arguments");
+
+  // Mixed argument kinds, including an enum and a float.
+  Q.memset(Memory, 0, N * sizeof(int));
+  oneapiext::nd_launch(Q, Ndr, getKernel<addMixed>(Q), Memory, 1, 2ul, 3.0f,
+                       Sign::Plus);
+  Q.wait();
+  for (size_t I = 0; I < N; ++I)
+    Failed += Check(Memory, 6, I, "mixed argument kinds");
+
+  // Ordering against the preceding commands on an in-order queue has to hold
+  // for a launch that bypasses the scheduler.
+  Q.memset(Memory, 0, N * sizeof(int));
+  for (int I = 0; I < 8; ++I)
+    oneapiext::nd_launch(Q, Ndr, ScalarsKernel, Memory, 1, 0);
+  Q.wait();
+  for (size_t I = 0; I < N; ++I)
+    Failed += Check(Memory, 8, I, "in-order accumulation");
+
+  // An accessor cannot be bound as plain bytes, so this has to fall back to the
+  // command group path and still produce the right result.
+  {
+    std::vector<int> Data(N, 5);
+    sycl::buffer<int, 1> Buf{Data.data(), sycl::range<1>{N}};
+    sycl::kernel AccessorKernel = getKernel<addViaAccessor>(Q);
+    Q.submit([&](sycl::handler &CGH) {
+      sycl::accessor Acc{Buf, CGH, sycl::read_write};
+      oneapiext::nd_launch(CGH, Ndr, AccessorKernel, Acc, 4);
+    });
+    Q.wait();
+    sycl::host_accessor Host{Buf};
+    for (size_t I = 0; I < N; ++I)
+      Failed +=
+          Check(&Host[0], 9, I, "accessor through the command group path");
+  }
+
+  // These overloads cannot carry a work_group_scratch_size property, so a
+  // kernel that allocates work group scratch memory has to be reported rather
+  // than launched, exactly as the command group path reports it.
+  {
+    bool Reported = false;
+    try {
+      oneapiext::nd_launch(Q, Ndr, getKernel<usesScratch>(Q), Memory);
+      Q.wait_and_throw();
+    } catch (const sycl::exception &E) {
+      Reported = E.code() == sycl::errc::memory_allocation;
+    }
+    if (!Reported) {
+      std::cout << "Failed: work group scratch memory without a size property "
+                   "was not reported"
+                << std::endl;
+      ++Failed;
+    }
+  }
+
+  sycl::free(Memory, Q);
+  return Failed;
+}
diff --git a/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_direct_raw_ptr.cpp b/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_direct_raw_ptr.cpp
new file mode 100644
index 0000000000000..4d1f4c111161a
--- /dev/null
+++ b/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_direct_raw_ptr.cpp
@@ -0,0 +1,67 @@
+// REQUIRES: aspect-usm_shared_allocations
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+// Tests the nd_launch overload that takes an already built sycl::kernel with
+// every argument passed through raw_kernel_arg, the USM pointer included. That
+// is what a caller which only knows the signature as sizes has to do.
+//
+// The pointer is passed through the pointer form of raw_kernel_arg. Passing the
+// bytes of a pointer instead would bind it as a value argument, which only
+// reaches the kernel on Level Zero: the OpenCL adapter passes a value argument
+// to clSetKernelArg, which rejects a USM pointer with CL_INVALID_MEM_OBJECT,
+// and the Native CPU adapter puts the address of its own copy of the bytes into
+// the argument slot instead of the pointer itself. The RawKernelArg tests cover
+// the byte form of a pointer, on Level Zero only.
+
+#include <sycl/detail/core.hpp>
+#include <sycl/ext/oneapi/experimental/enqueue_functions.hpp>
+#include <sycl/ext/oneapi/experimental/raw_kernel_arg.hpp>
+#include <sycl/ext/oneapi/free_function_queries.hpp>
+#include <sycl/kernel_bundle.hpp>
+#include <sycl/properties/all_properties.hpp>
+#include <sycl/usm.hpp>
+
+#include "common.hpp"
+
+namespace syclext = sycl::ext::oneapi;
+
+constexpr size_t N = 1024;
+constexpr size_t WGSize = 8;
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((oneapiext::nd_range_kernel<1>))
+void addScalars(int *Ptr, int A, int B) {
+  size_t I = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
+  Ptr[I] += A + B;
+}
+
+template <auto *Func> sycl::kernel getKernel(sycl::queue &Q) {
+  auto Bundle =
+      oneapiext::get_kernel_bundle<Func, sycl::bundle_state::executable>(
+          Q.get_context());
+  return Bundle.template ext_oneapi_get_kernel<Func>();
+}
+
+int main() {
+  sycl::queue Q{sycl::property::queue::in_order{}};
+  int *Memory = sycl::malloc_shared<int>(N, Q);
+  sycl::nd_range<1> Ndr{sycl::range<1>{N}, sycl::range<1>{WGSize}};
+
+  int Failed = 0;
+
+  int A = 10, B = 20;
+  sycl::kernel ScalarsKernel = getKernel<addScalars>(Q);
+  Q.memset(Memory, 0, N * sizeof(int));
+  oneapiext::nd_launch(
+      Q, Ndr, ScalarsKernel,
+      oneapiext::raw_kernel_arg{&Memory, oneapiext::pointer_arg},
+      oneapiext::raw_kernel_arg{&A, sizeof(A)},
+      oneapiext::raw_kernel_arg{&B, sizeof(B)});
+  Q.wait();
+  for (size_t I = 0; I < N; ++I)
+    Failed +=
+        Check(Memory, 30, I, "pointer and scalars through raw_kernel_arg");
+
+  sycl::free(Memory, Q);
+  return Failed;
+}
diff --git a/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_span.cpp b/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_span.cpp
new file mode 100644
index 0000000000000..e69975b2af3b5
--- /dev/null
+++ b/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_span.cpp
@@ -0,0 +1,162 @@
+// REQUIRES: aspect-usm_shared_allocations
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+// Tests the nd_launch overloads that take the arguments of a sycl::kernel as a
+// span of raw_kernel_arg, i.e. an argument list whose length is only known at
+// run time. They have to bind the same arguments in the same order as the
+// parameter pack overloads, on the queue and on the handler alike.
+
+#include <sycl/detail/core.hpp>
+#include <sycl/ext/oneapi/experimental/enqueue_functions.hpp>
+#include <sycl/ext/oneapi/experimental/raw_kernel_arg.hpp>
+#include <sycl/ext/oneapi/free_function_queries.hpp>
+#include <sycl/kernel_bundle.hpp>
+#include <sycl/properties/all_properties.hpp>
+#include <sycl/usm.hpp>
+
+#include "common.hpp"
+
+#include <vector>
+
+namespace syclext = sycl::ext::oneapi;
+
+static_assert(SYCL_EXT_ONEAPI_ENQUEUE_FUNCTIONS >= 2,
+              "The span overloads require version 2 of the extension");
+
+constexpr size_t N = 1024;
+constexpr size_t WGSize = 8;
+
+// A mixture of argument sizes, so that a wrong size or a wrong order shows up
+// as a wrong result rather than as a silent pass.
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((oneapiext::nd_range_kernel<1>))
+void addMixed(int *Ptr, int A, long B, float C, char D) {
+  size_t I = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
+  Ptr[I] += A + static_cast<int>(B) + static_cast<int>(C) + D;
+}
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((oneapiext::nd_range_kernel<1>))
+void increment(int *Ptr) {
+  size_t I = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
+  Ptr[I] += 1;
+}
+
+template <auto *Func> sycl::kernel getKernel(sycl::queue &Q) {
+  auto Bundle =
+      oneapiext::get_kernel_bundle<Func, sycl::bundle_state::executable>(
+          Q.get_context());
+  return Bundle.template ext_oneapi_get_kernel<Func>();
+}
+
+int main() {
+  sycl::queue Q{sycl::property::queue::in_order{}};
+  sycl::kernel Kernel = getKernel<addMixed>(Q);
+
+  int *Memory = sycl::malloc_shared<int>(N, Q);
+  sycl::nd_range<1> Ndr{sycl::range<1>{N}, sycl::range<1>{WGSize}};
+
+  int A = 1;
+  long B = 20;
+  float C = 300.0f;
+  char D = 4;
+  constexpr int Sum = 1 + 20 + 300 + 4;
+
+  // The argument list is built at run time, which is the case these overloads
+  // exist for. The pointer says that it is one, so that it is bound as a
+  // pointer rather than as the bytes it is made of, which only the Level Zero
+  // backend binds as a pointer.
+  std::vector<oneapiext::raw_kernel_arg> Args;
+  Args.emplace_back(&Memory, oneapiext::pointer_arg);
+  Args.emplace_back(&A, sizeof(A));
+  Args.emplace_back(&B, sizeof(B));
+  Args.emplace_back(&C, sizeof(C));
+  Args.emplace_back(&D, sizeof(D));
+  sycl::span<const oneapiext::raw_kernel_arg> ArgSpan{Args.data(), Args.size()};
+
+  int Failed = 0;
+
+  // A run of launches through one span, so that an argument bound to storage
+  // that does not outlive a single call would show up.
+  constexpr int Launches = 8;
+  Q.memset(Memory, 0, N * sizeof(int));
+  for (int I = 0; I < Launches; ++I)
+    oneapiext::nd_launch(Q, Ndr, Kernel, ArgSpan);
+  Q.wait();
+  for (size_t I = 0; I < N; ++I)
+    Failed += Check(Memory, Sum * Launches, I, "span overload");
+
+  // The parameter pack overload has to agree element for element.
+  Q.memset(Memory, 0, N * sizeof(int));
+  oneapiext::nd_launch(
+      Q, Ndr, Kernel,
+      oneapiext::raw_kernel_arg{&Memory, oneapiext::pointer_arg},
+      oneapiext::raw_kernel_arg{&A, sizeof(A)},
+      oneapiext::raw_kernel_arg{&B, sizeof(B)},
+      oneapiext::raw_kernel_arg{&C, sizeof(C)},
+      oneapiext::raw_kernel_arg{&D, sizeof(D)});
+  Q.wait();
+  for (size_t I = 0; I < N; ++I)
+    Failed += Check(Memory, Sum, I, "parameter pack overload");
+
+  // And so does the handler form of the span overload.
+  Q.memset(Memory, 0, N * sizeof(int));
+  Q.submit([&](sycl::handler &CGH) {
+     oneapiext::nd_launch(CGH, Ndr, Kernel, ArgSpan);
+   }).wait();
+  for (size_t I = 0; I < N; ++I)
+    Failed += Check(Memory, Sum, I, "handler form of the span overload");
+
+  // The container holding the arguments converts to that span, so passing it
+  // has to bind the arguments it holds rather than the container object.
+  Q.memset(Memory, 0, N * sizeof(int));
+  oneapiext::nd_launch(Q, Ndr, Kernel, Args);
+  Q.wait();
+  for (size_t I = 0; I < N; ++I)
+    Failed += Check(Memory, Sum, I, "argument list passed as a container");
+
+  Q.memset(Memory, 0, N * sizeof(int));
+  Q.submit([&](sycl::handler &CGH) {
+     oneapiext::nd_launch(CGH, Ndr, Kernel, Args);
+   }).wait();
+  for (size_t I = 0; I < N; ++I)
+    Failed += Check(Memory, Sum, I, "container through the handler form");
+
+  // A mutable span of the same sequence names the same argument list.
+  Q.memset(Memory, 0, N * sizeof(int));
+  oneapiext::nd_launch(
+      Q, Ndr, Kernel,
+      sycl::span<oneapiext::raw_kernel_arg>{Args.data(), Args.size()});
+  Q.wait();
+  for (size_t I = 0; I < N; ++I)
+    Failed += Check(Memory, Sum, I, "argument list as a mutable span");
+
+  // A one element span is the boundary against the parameter pack overload,
+  // which a single raw_kernel_arg selects instead.
+  std::vector<oneapiext::raw_kernel_arg> OneArg{
+      {&Memory, oneapiext::pointer_arg}};
+  Q.memset(Memory, 0, N * sizeof(int));
+  oneapiext::nd_launch(Q, Ndr, getKernel<increment>(Q),
+                       sycl::span<const oneapiext::raw_kernel_arg>{
+                           OneArg.data(), OneArg.size()});
+  Q.wait();
+  for (size_t I = 0; I < N; ++I)
+    Failed += Check(Memory, 1, I, "one element span");
+
+  // A dependency the scheduler has to track forces the command group path,
+  // which the span form has to take as well.
+  {
+    sycl::buffer<int, 1> Buf{sycl::range<1>{N}};
+    Q.memset(Memory, 0, N * sizeof(int));
+    Q.submit([&](sycl::handler &CGH) {
+      sycl::accessor Acc{Buf, CGH, sycl::write_only, sycl::no_init};
+      CGH.fill(Acc, 7);
+    });
+    oneapiext::nd_launch(Q, Ndr, Kernel, ArgSpan);
+    Q.wait();
+    for (size_t I = 0; I < N; ++I)
+      Failed += Check(Memory, Sum, I, "span form behind a buffer dependency");
+  }
+
+  sycl::free(Memory, Q);
+  return Failed;
+}
diff --git a/sycl/test-e2e/Graph/Update/update_with_raw_kernel_arg_pointer.cpp b/sycl/test-e2e/Graph/Update/update_with_raw_kernel_arg_pointer.cpp
new file mode 100644
index 0000000000000..e899c9d53c0fd
--- /dev/null
+++ b/sycl/test-e2e/Graph/Update/update_with_raw_kernel_arg_pointer.cpp
@@ -0,0 +1,83 @@
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+// Extra run to check for leaks in Level Zero using UR_L0_LEAKS_DEBUG
+// RUN: %if level_zero %{%{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
+
+// REQUIRES: ocloc && level_zero
+
+// Tests updating a raw_kernel_arg that was built as a pointer argument, which
+// has to rebind the pointer rather than the bytes one is made of.
+
+#include "../graph_common.hpp"
+
+auto constexpr CLSource = R"===(
+__kernel void RawArgKernel(int scalar, __global int *out) {
+  size_t id = get_global_id(0);
+  out[id] = id + scalar;
+}
+)===";
+
+int main() {
+  queue Queue{};
+
+  auto SourceKB =
+      sycl::ext::oneapi::experimental::create_kernel_bundle_from_source(
+          Queue.get_context(),
+          sycl::ext::oneapi::experimental::source_language::opencl, CLSource);
+  auto ExecKB = sycl::ext::oneapi::experimental::build(SourceKB);
+
+  exp_ext::command_graph Graph{Queue};
+
+  const size_t N = 1024;
+  int32_t *PtrA = malloc_device<int32_t>(N, Queue);
+  int32_t *PtrB = malloc_device<int32_t>(N, Queue);
+  Queue.memset(PtrA, 0, N * sizeof(int32_t));
+  Queue.memset(PtrB, 0, N * sizeof(int32_t));
+  Queue.wait();
+
+  int32_t Scalar = 42;
+  exp_ext::raw_kernel_arg RawScalar(&Scalar, sizeof(int32_t));
+  exp_ext::raw_kernel_arg RawPtrA(&PtrA, exp_ext::pointer_arg);
+  exp_ext::raw_kernel_arg RawPtrB(&PtrB, exp_ext::pointer_arg);
+
+  exp_ext::dynamic_parameter ScalarParam(RawScalar);
+  exp_ext::dynamic_parameter PtrParam(RawPtrA);
+
+  auto KernelNode = Graph.add([&](handler &cgh) {
+    cgh.set_arg(0, ScalarParam);
+    cgh.set_arg(1, PtrParam);
+    cgh.parallel_for(sycl::range<1>{N},
+                     ExecKB.ext_oneapi_get_kernel("RawArgKernel"));
+  });
+
+  auto ExecGraph = Graph.finalize(exp_ext::property::graph::updatable{});
+
+  // PtrA is the one the pointer argument was built from.
+  Queue.ext_oneapi_graph(ExecGraph).wait();
+
+  std::vector<int> HostDataA(N);
+  std::vector<int> HostDataB(N);
+  Queue.copy(PtrA, HostDataA.data(), N);
+  Queue.copy(PtrB, HostDataB.data(), N);
+  Queue.wait();
+  for (size_t i = 0; i < N; i++) {
+    assert(HostDataA[i] == static_cast<int>(i + Scalar));
+    assert(HostDataB[i] == 0);
+  }
+
+  // Rebind the pointer argument to the other allocation.
+  PtrParam.update(RawPtrB);
+  ExecGraph.update(KernelNode);
+  Queue.ext_oneapi_graph(ExecGraph).wait();
+
+  Queue.copy(PtrA, HostDataA.data(), N).wait();
+  Queue.copy(PtrB, HostDataB.data(), N).wait();
+  for (size_t i = 0; i < N; i++) {
+    assert(HostDataA[i] == static_cast<int>(i + Scalar));
+    assert(HostDataB[i] == static_cast<int>(i + Scalar));
+  }
+
+  sycl::free(PtrA, Queue);
+  sycl::free(PtrB, Queue);
+  return 0;
+}
diff --git a/sycl/test-e2e/RawKernelArg/pointer_arg.cpp b/sycl/test-e2e/RawKernelArg/pointer_arg.cpp
new file mode 100644
index 0000000000000..e0e894e69ac92
--- /dev/null
+++ b/sycl/test-e2e/RawKernelArg/pointer_arg.cpp
@@ -0,0 +1,71 @@
+// REQUIRES: aspect-usm_shared_allocations
+// REQUIRES: ocloc && (opencl || level_zero)
+
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+// Tests the pointer form of raw_kernel_arg, which says that an argument is a
+// pointer instead of leaving the runtime to bind the bytes it is made of. That
+// is what makes a pointer argument reach the kernel on a backend which takes a
+// pointer through a different entry point than a value: OpenCL passes a value
+// argument to clSetKernelArg, which rejects a USM pointer, and Native CPU puts
+// the address of its own copy of the bytes into the argument slot. Hence no
+// Level Zero requirement here, unlike the tests of the byte form.
+
+#include <sycl/detail/core.hpp>
+#include <sycl/ext/oneapi/experimental/raw_kernel_arg.hpp>
+#include <sycl/kernel_bundle.hpp>
+#include <sycl/usm.hpp>
+
+#include <cassert>
+
+namespace oneapiext = sycl::ext::oneapi::experimental;
+
+auto constexpr CLSource = R"===(
+__kernel void WriteScalar(int in, __global int *out) {
+  out[get_global_id(0)] = in;
+}
+)===";
+
+constexpr size_t N = 8;
+
+int main() {
+  sycl::queue Q;
+
+  auto SourceKB = oneapiext::create_kernel_bundle_from_source(
+      Q.get_context(), oneapiext::source_language::opencl, CLSource);
+  auto ExecKB = oneapiext::build(SourceKB);
+  sycl::kernel Kernel = ExecKB.ext_oneapi_get_kernel("WriteScalar");
+
+  int *Out = sycl::malloc_shared<int>(N, Q);
+  int In = 42;
+
+  // Both arguments are raw, the pointer as a pointer and the scalar as bytes.
+  // The queue is out of order, so the sentinel has to be waited for rather than
+  // left to race with the kernel that overwrites it.
+  Q.memset(Out, 0xFF, N * sizeof(int)).wait();
+  Q.submit([&](sycl::handler &CGH) {
+     CGH.set_arg(0, oneapiext::raw_kernel_arg{&In, sizeof(In)});
+     CGH.set_arg(1, oneapiext::raw_kernel_arg{&Out, oneapiext::pointer_arg});
+     CGH.parallel_for(sycl::range<1>{N}, Kernel);
+   }).wait();
+  for (size_t I = 0; I < N; ++I)
+    assert(Out[I] == In);
+
+  // The same through set_args, and with the pointer bound to a different
+  // allocation, so that a stale argument would show up.
+  int *Other = sycl::malloc_shared<int>(N, Q);
+  In = 7;
+  Q.memset(Other, 0xFF, N * sizeof(int)).wait();
+  Q.submit([&](sycl::handler &CGH) {
+     CGH.set_args(oneapiext::raw_kernel_arg{&In, sizeof(In)},
+                  oneapiext::raw_kernel_arg{&Other, oneapiext::pointer_arg});
+     CGH.parallel_for(sycl::range<1>{N}, Kernel);
+   }).wait();
+  for (size_t I = 0; I < N; ++I)
+    assert(Other[I] == In);
+
+  sycl::free(Other, Q);
+  sycl::free(Out, Q);
+  return 0;
+}
diff --git a/sycl/test/abi/layout_kernel_arg_view.cpp b/sycl/test/abi/layout_kernel_arg_view.cpp
new file mode 100644
index 0000000000000..5ebd25d87eaea
--- /dev/null
+++ b/sycl/test/abi/layout_kernel_arg_view.cpp
@@ -0,0 +1,17 @@
+// RUN: %clangxx -fsycl -c -fno-color-diagnostics -Xclang -fdump-record-layouts %s -o %t.out | FileCheck %s
+// RUN: %clangxx -fsycl -fsycl-device-only -c -fno-color-diagnostics -Xclang -fdump-record-layouts %s -o %t.out | FileCheck %s
+// REQUIRES: linux
+// UNSUPPORTED: libcxx
+
+// clang-format off
+
+#include <sycl/detail/kernel_arg_view.hpp>
+
+
+SYCL_EXTERNAL void kernel_arg_view(sycl::detail::kernel_arg_view_v1::KernelArgView) {}
+// CHECK: 0 | struct sycl::detail::KernelArgView
+// CHECK-NEXT: 0 |   const void * MPtr
+// CHECK-NEXT: 8 |   size_t MSize
+// CHECK-NEXT: 16 |  kernel_param_kind_t MKind
+// CHECK-NEXT: | [sizeof=24, dsize=24, align=8,
+// CHECK-NEXT: |  nvsize=24, nvalign=8]
diff --git a/sycl/test/abi/layout_raw_kernel_arg.cpp b/sycl/test/abi/layout_raw_kernel_arg.cpp
new file mode 100644
index 0000000000000..0b34c35316a98
--- /dev/null
+++ b/sycl/test/abi/layout_raw_kernel_arg.cpp
@@ -0,0 +1,21 @@
+// RUN: %clangxx -fsycl -c -fno-color-diagnostics -Xclang -fdump-record-layouts %s -o %t.out | FileCheck %s
+// RUN: %clangxx -fsycl -fsycl-device-only -c -fno-color-diagnostics -Xclang -fdump-record-layouts %s -o %t.out | FileCheck %s
+// REQUIRES: linux
+// UNSUPPORTED: libcxx
+
+// clang-format off
+
+#include <sycl/detail/defines_elementary.hpp> // for SYCL_EXTERNAL
+#include <sycl/ext/oneapi/experimental/raw_kernel_arg.hpp>
+
+// A raw_kernel_arg crosses the ABI boundary as an element of the span that the
+// nd_launch overloads take, and the graph extension copies one as bytes, so its
+// layout is fixed here. MIsPointer comes last, so that the two members the
+// library read before it existed keep their offsets.
+SYCL_EXTERNAL void takeRawKernelArg(sycl::ext::oneapi::experimental::raw_kernel_arg) {}
+// CHECK: 0 | class sycl::ext::oneapi::experimental::raw_kernel_arg
+// CHECK-NEXT: 0 |   const void * MArgData
+// CHECK-NEXT: 8 |   size_t MArgSize
+// CHECK-NEXT: 16 |  _Bool MIsPointer
+// CHECK-NEXT: | [sizeof=24, dsize=17, align=8,
+// CHECK-NEXT: |  nvsize=17, nvalign=8]
diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump
index 01137ab1ab81c..a460dec11dea2 100644
--- a/sycl/test/abi/sycl_symbols_linux.dump
+++ b/sycl/test/abi/sycl_symbols_linux.dump
@@ -3163,6 +3163,8 @@ _ZN4sycl3_V13ext6oneapi15filter_selectorC1ENS0_6detail11string_viewE
 _ZN4sycl3_V13ext6oneapi15filter_selectorC2ENS0_6detail11string_viewE
 _ZN4sycl3_V13ext6oneapilsERSoRKNS2_8bfloat16E
 _ZN4sycl3_V13ext6oneapirsERSiRNS2_8bfloat16E
+_ZN4sycl3_V143submit_kernel_obj_direct_without_event_implERKNS0_5queueERKNS0_6detail16nd_range_view_v113nd_range_viewERKNS0_6kernelENS0_4spanIKNS0_3ext6oneapi12experimental14raw_kernel_argELm18446744073709551615EEERKNS4_13code_locationEb
+_ZN4sycl3_V143submit_kernel_obj_direct_without_event_implERKNS0_5queueERKNS0_6detail16nd_range_view_v113nd_range_viewERKNS0_6kernelENS0_4spanIKNS4_18kernel_arg_view_v113KernelArgViewELm18446744073709551615EEERKNS4_13code_locationEb
 _ZN4sycl3_V14freeEPvRKNS0_5queueERKNS0_6detail13code_locationE
 _ZN4sycl3_V14freeEPvRKNS0_7contextERKNS0_6detail13code_locationE
 _ZN4sycl3_V15event13get_wait_listEv
diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump
index 70a39a2743b10..13741eb2d21f7 100644
--- a/sycl/test/abi/sycl_symbols_windows.dump
+++ b/sycl/test/abi/sycl_symbols_windows.dump
@@ -4388,6 +4388,8 @@
 ?submit_graph_direct_without_event_impl@_V1@sycl@@YAXAEBVqueue@12@AEAV?$command_graph@$00@experimental@oneapi@ext@12@V?$span@$$CBVevent@_V1@sycl@@$0?0@12@AEBUcode_location@detail@12@@Z
 ?submit_kernel_direct_with_event_impl@_V1@sycl@@YA?AVevent@12@AEBVqueue@12@AEBVnd_range_view@nd_range_view_v1@detail@12@AEAVHostKernelRefBase@712@PEAVDeviceKernelInfo@712@V?$span@$$CBVevent@_V1@sycl@@$0?0@12@AEBU?$PropsHolder@Uwork_group_scratch_size@experimental@oneapi@ext@_V1@sycl@@Ucache_config@2intel@456@Uuse_root_sync_key@23456@Uwork_group_progress_key@23456@Usub_group_progress_key@23456@Uwork_item_progress_key@23456@U?$cluster_size@$00@cuda@23456@U?$cluster_size@$01@cuda@23456@U?$cluster_size@$02@cuda@23456@@kernel_launch_properties_v1@712@AEBUcode_location@712@_N@Z
 ?submit_kernel_direct_without_event_impl@_V1@sycl@@YAXAEBVqueue@12@AEBVnd_range_view@nd_range_view_v1@detail@12@AEAVHostKernelRefBase@612@PEAVDeviceKernelInfo@612@V?$span@$$CBVevent@_V1@sycl@@$0?0@12@AEBU?$PropsHolder@Uwork_group_scratch_size@experimental@oneapi@ext@_V1@sycl@@Ucache_config@2intel@456@Uuse_root_sync_key@23456@Uwork_group_progress_key@23456@Usub_group_progress_key@23456@Uwork_item_progress_key@23456@U?$cluster_size@$00@cuda@23456@U?$cluster_size@$01@cuda@23456@U?$cluster_size@$02@cuda@23456@@kernel_launch_properties_v1@612@AEBUcode_location@612@_N@Z
+?submit_kernel_obj_direct_without_event_impl@_V1@sycl@@YAXAEBVqueue@12@AEBVnd_range_view@nd_range_view_v1@detail@12@AEBVkernel@12@V?$span@$$CBUKernelArgView@kernel_arg_view_v1@detail@_V1@sycl@@$0?0@12@AEBUcode_location@612@_N@Z
+?submit_kernel_obj_direct_without_event_impl@_V1@sycl@@YAXAEBVqueue@12@AEBVnd_range_view@nd_range_view_v1@detail@12@AEBVkernel@12@V?$span@$$CBVraw_kernel_arg@experimental@oneapi@ext@_V1@sycl@@$0?0@12@AEBUcode_location@612@_N@Z
 ?submit_with_event_impl@queue@_V1@sycl@@AEBA?AVevent@23@AEBVtype_erased_cgfo_ty@detail@23@AEBVSubmissionInfo@2623@AEBUcode_location@623@_N@Z
 ?submit_without_event_impl@queue@_V1@sycl@@AEBAXAEBVtype_erased_cgfo_ty@detail@23@AEBVSubmissionInfo@2523@AEBUcode_location@523@_N@Z
 ?supportsUSMFill2D@handler@_V1@sycl@@AEAA_NXZ
diff --git a/sycl/test/extensions/enqueue_functions/kernel_arg_span.cpp b/sycl/test/extensions/enqueue_functions/kernel_arg_span.cpp
new file mode 100644
index 0000000000000..208b73f310ead
--- /dev/null
+++ b/sycl/test/extensions/enqueue_functions/kernel_arg_span.cpp
@@ -0,0 +1,57 @@
+// RUN: %clangxx -fsycl -fsyntax-only %s
+// RUN: %clangxx -fsycl -fsyntax-only -std=c++20 %s
+
+// An argument list held in a container converts to the sycl::span that the
+// nd_launch span overloads take, but a parameter pack is an exact match and
+// wins overload resolution. Without the forwarding the pack overloads do, the
+// container object itself would be bound as a single kernel argument, which
+// compiles for any trivially copyable container and produces wrong results at
+// run time. Check that every spelling of an argument list is accepted, and that
+// a single raw_kernel_arg is still one argument.
+
+#include <sycl/sycl.hpp>
+
+#include <array>
+#include <vector>
+#if __cpp_lib_span
+#include <span>
+#endif
+
+namespace oneapiext = sycl::ext::oneapi::experimental;
+
+void argument_list_spellings(sycl::queue Q, sycl::handler &CGH,
+                             sycl::nd_range<1> Range,
+                             const sycl::kernel &Kernel) {
+  int Value = 1;
+  std::vector<oneapiext::raw_kernel_arg> Vector{{&Value, sizeof(Value)}};
+  std::array<oneapiext::raw_kernel_arg, 1> Array{
+      oneapiext::raw_kernel_arg{&Value, sizeof(Value)}};
+  sycl::span<const oneapiext::raw_kernel_arg> Span{Vector.data(),
+                                                   Vector.size()};
+  sycl::span<oneapiext::raw_kernel_arg> MutableSpan{Vector.data(),
+                                                    Vector.size()};
+
+  oneapiext::nd_launch(Q, Range, Kernel, Vector);
+  oneapiext::nd_launch(Q, Range, Kernel, Array);
+  oneapiext::nd_launch(Q, Range, Kernel, Span);
+  oneapiext::nd_launch(Q, Range, Kernel, MutableSpan);
+  oneapiext::nd_launch(CGH, Range, Kernel, Vector);
+  oneapiext::nd_launch(CGH, Range, Kernel, Array);
+  oneapiext::nd_launch(CGH, Range, Kernel, Span);
+  oneapiext::nd_launch(CGH, Range, Kernel, MutableSpan);
+#if __cpp_lib_span
+  std::span<const oneapiext::raw_kernel_arg> StdSpan{Vector.data(),
+                                                     Vector.size()};
+  oneapiext::nd_launch(Q, Range, Kernel, StdSpan);
+  oneapiext::nd_launch(CGH, Range, Kernel, StdSpan);
+#endif
+
+  // One raw_kernel_arg is one argument, and typed arguments are unaffected.
+  oneapiext::nd_launch(Q, Range, Kernel,
+                       oneapiext::raw_kernel_arg{&Value, sizeof(Value)});
+  oneapiext::nd_launch(CGH, Range, Kernel,
+                       oneapiext::raw_kernel_arg{&Value, sizeof(Value)});
+  int *Pointer = nullptr;
+  oneapiext::nd_launch(Q, Range, Kernel, Pointer, Value);
+  oneapiext::nd_launch(CGH, Range, Kernel, Pointer, Value);
+}
diff --git a/sycl/test/extensions/raw_kernel_arg/pointer_arg_errors.cpp b/sycl/test/extensions/raw_kernel_arg/pointer_arg_errors.cpp
new file mode 100644
index 0000000000000..eaad709fe3111
--- /dev/null
+++ b/sycl/test/extensions/raw_kernel_arg/pointer_arg_errors.cpp
@@ -0,0 +1,37 @@
+// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note %s
+
+// The pointer form of raw_kernel_arg takes the address of the pointer, in
+// keeping with the byte form taking the address of the bytes. Passing the
+// pointer itself would name an address the runtime would then read a pointer
+// from, which is why it does not compile.
+
+#include <sycl/ext/oneapi/experimental/raw_kernel_arg.hpp>
+
+#include <type_traits>
+
+namespace oneapiext = sycl::ext::oneapi::experimental;
+
+// The graph extension copies a raw_kernel_arg as bytes.
+static_assert(std::is_trivially_copyable_v<oneapiext::raw_kernel_arg>);
+
+void pointer_form(int *Ptr, const float *ConstPtr, void *VoidPtr) {
+  // Every pointer type binds through the address form without a cast.
+  oneapiext::raw_kernel_arg Typed{&Ptr, oneapiext::pointer_arg};
+  oneapiext::raw_kernel_arg Const{&ConstPtr, oneapiext::pointer_arg};
+  oneapiext::raw_kernel_arg Void{&VoidPtr, oneapiext::pointer_arg};
+
+  // The byte form is unchanged, including for the bytes of a pointer, which
+  // only bind as a pointer on the Level Zero backend.
+  oneapiext::raw_kernel_arg Bytes{&Ptr, sizeof(Ptr)};
+
+  // expected-error@+1 {{no matching constructor for initialization of 'oneapiext::raw_kernel_arg'}}
+  oneapiext::raw_kernel_arg PointerItself{Ptr, oneapiext::pointer_arg};
+
+  // expected-error@+1 {{no matching constructor for initialization of 'oneapiext::raw_kernel_arg'}}
+  oneapiext::raw_kernel_arg NotAPointer{42, oneapiext::pointer_arg};
+
+  (void)Typed;
+  (void)Const;
+  (void)Void;
+  (void)Bytes;
+}