[SYCL][Doc] Add sycl_ext_oneapi_raw_kernel_arg

sycl/doc/extensions/proposed/sycl_ext_oneapi_raw_kernel_arg.asciidoc

@@ -0,0 +1,149 @@
+= sycl_ext_oneapi_raw_kernel_arg
+
+:source-highlighter: coderay
+:coderay-linenums-mode: table
+
+// This section needs to be after the document title.
+:doctype: book
+:toc2:
+:toc: left
+:encoding: utf-8
+:lang: en
+:dpcpp: pass:[DPC++]
+:endnote: —{nbsp}end{nbsp}note
+
+// Set the default source code type in this document to C++,
+// for syntax highlighting purposes.  This is needed because
+// docbook uses c++ and html5 uses cpp.
+:language: {basebackend@docbook:c++:cpp}
+
+
+== Notice
+
+[%hardbreaks]
+Copyright (C) 2023-2023 Intel Corporation.  All rights reserved.
+
+Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are trademarks
+of The Khronos Group Inc.  OpenCL(TM) is a trademark of Apple Inc. used by
+permission by Khronos.
+
+
+== Contact
+
+To report problems with this extension, please open a new issue at:
+
+https://github.com/intel/llvm/issues
+
+
+== Dependencies
+
+This extension is written against the SYCL 2020 revision 8 specification.  All
+references below to the "core SYCL specification" or to section numbers in the
+SYCL specification refer to that revision.
+
+
+== Status
+
+This is a proposed extension specification, intended to gather community
+feedback.  Interfaces defined in this specification may not be implemented yet
+or may be in a preliminary state.  The specification itself may also change in
+incompatible ways before it is finalized.  *Shipping software products should
+not rely on APIs defined in this specification.*
+
+
+== Overview
+
+When launching kernels that are represented as `sycl::kernel` objects,
+developers must pass arguments via the `set_arg` or `set_args` functions.
+
+There are situations where developers would like to pass a raw byte
+representation of the kernel argument to the backend (e.g., when a single
+`parallel_for` in the source code is used to invoke multiple kernels with
+different arguments types, or when passing an argument to a built-in kernel
+for which there is no equivalent type defined on the host).
+
+=== Usage example
+
+[source,c++]
+----
+int a;
+char* opaque_type;
+int nbytes;
+...
+h.set_arg(0, a);
+h.set_arg(1, sycl::ext::oneapi::raw_kernel_arg(opaque_type, nbytes));
+h.parallel_for(range, kernel);
+----
+
+
+== Specification
+
+=== Feature test macro
+
+This extension provides a feature-test macro as described in the core SYCL
+specification.  An implementation supporting this extension must predefine the
+macro `SYCL_EXT_ONEAPI_RAW_KERNEL_ARG` to one of the values defined in the
+table below.  Applications can test for the existence of this macro to
+determine if the implementation supports this feature, or applications can test
+the macro's value to determine which of the extension's features the
+implementation supports.
+
+[%header,cols="1,5"]
+|===
+|Value
+|Description
+
+|1
+|Initial version of this extension.
+|===
+
+=== The `raw_kernel_arg` class
+
+This extension adds a new `raw_kernel_arg` class that can be used to declare
+kernel arguments via a raw byte representation.
+
+[source,c++]
+----
+namespace sycl::ext::oneapi {
+
+  class raw_kernel_arg {
+  public:
+    raw_kernel_arg(void* bytes, size_t count);
+  };
+
+} // namespace sycl::ext::oneapi
+----
+
+[source,c++]
+----
+raw_kernel_arg(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.
+
+_Effects_: Constructs a `raw_kernel_arg` representing a view of the `count`
+bytes starting at the address specified by `bytes`.
+
+=== Using a raw kernel argument
+
+Instances of `raw_kernel_arg` are passed to kernels via the existing `set_arg`
+and `set_args` functions defined by the SYCL specification.
+
+This extension adds a new overload of `set_arg`, as defined below.
+
+[_Note:_ Since the definition of `set_args` says that it acts "as if each
+argument in `args` was passed to `set_arg` ", adding a new overload of
+`set_arg` is sufficient to change the behavior of `set_args`. _{endnote}_]
+
+[source,c++]
+----
+void set_arg(int argIndex, sycl::ext::oneapi::raw_kernel_arg&& arg);
+----
+_Effects_: Sets the kernel argument associated with index `argIndex` using the
+bytes represented by `arg`.
+
+
+== Issues
+
+None.