[SYCL][Doc][RTC] Add device global queries

sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler.asciidoc

@@ -48,6 +48,8 @@ This extension also depends on the following other SYCL extensions:
   sycl_ext_oneapi_properties]
 * link:../proposed/sycl_ext_oneapi_free_function_kernels.asciidoc[
   sycl_ext_oneapi_free_function_kernels]
+* link:../experimental/sycl_ext_oneapi_device_global.asciidoc[
+  sycl_ext_oneapi_device_global]
 
 
 == Status
@@ -582,6 +584,8 @@ class kernel_bundle {
   bool ext_oneapi_has_kernel(const std::string &name);
   kernel ext_oneapi_get_kernel(const std::string &name);
   std::string ext_oneapi_get_raw_kernel_name(const std::string &name);
+
+  // Continued below in "New kernel bundle member functions for device globals"
 };
 
 } // namespace sycl
@@ -790,6 +794,113 @@ sycl::kernel k_float = kb.ext_oneapi_get_kernel("bartmpl<float>");
 sycl::kernel k_int = kb.ext_oneapi_get_kernel("bartmpl<int>");
 ----
 
+=== New kernel bundle member functions for device globals
+
+This extensions adds the following new `kernel_bundle` member functions to let
+the host application interact with device globals defined in runtime-compiled
+code. Device globals are only supported for the `source_language::sycl`
+language.
+
+This extension currently supports only a subset of the
+link:../experimental/sycl_ext_oneapi_device_global.asciidoc[
+sycl_ext_oneapi_device_global] extension:
+
+* Device globals must be declared at global scope.
+* Device globals declared with the `device_image_scope` property can be used in
+the runtime-compiled device code, but cannot be accessed from the host.
+
+We plan to lift both limitations in a future version of this extension.
+
+[source,c++]
+----
+namespace sycl {
+
+template <bundle_state State>
+class kernel_bundle {
+  // Continued from "New kernel bundle member functions"
+
+  bool ext_oneapi_has_device_global(const std::string &name);
+  void *ext_oneapi_get_device_global_address(const std::string &name,
+                                             const device &dev);
+  size_t ext_oneapi_get_device_global_size(const std::string &name);
+};
+
+} // namespace sycl
+----
+
+|====
+a|
+[frame=all,grid=none]
+!====
+a!
+[source,c++]
+----
+bool ext_oneapi_has_device_global(const std::string &name)
+----
+!====
+
+_Constraints:_ This function is not available when `State` is
+`bundle_state::ext_oneapi_source`.
+
+_Returns:_ The value `true` only if
+
+* the kernel bundle was created from a bundle of state
+  `bundle_state::ext_oneapi_source` in the language `source_language::sycl`, and
+* it defines a device global whose name is `name` and which was declared without
+  the `device_image_scope` property.
+
+a|
+[frame=all,grid=none]
+!====
+a!
+[source,c++]
+----
+void *ext_oneapi_get_device_global_address(const std::string &name,
+                                           const device &dev)
+----
+!====
+
+_Constraints:_ This function is not available when `State` is
+`bundle_state::ext_oneapi_source`.
+
+_Returns:_ A device USM pointer to the storage for the device global `name` on
+device `dev`.
+
+_Remarks:_ The contents of the device global may be read or written from the
+host by reading from or writing to this address. If the address is read before
+any kernel writes to the device global, the read operation returns the device
+global's initial value.
+
+_Throws:_
+
+* An `exception` with the `errc::invalid` error code if
+  `ext_oneapi_has_device_global(name)` returns `false`.
+* An `exception` with the `errc::invalid` error code if the context associated
+  with this bundle does not contain device `dev`.
+* An `exception` with the `errc::memory_allocation` error code if the allocation
+  or initialization of the device global's storage fails.
+
+a|
+[frame=all,grid=none]
+!====
+a!
+[source,c++]
+----
+size_t ext_oneapi_get_device_global_size(const std::string &name)
+----
+!====
+
+_Constraints:_ This function is not available when `State` is
+`bundle_state::ext_oneapi_source`.
+
+_Returns:_ The size in bytes of the USM storage for device global `name`.
+
+_Throws:_
+
+* An `exception` with the `errc::invalid` error code if
+  `ext_oneapi_has_device_global(name)` returns `false`.
+|====
+
 
 == Examples
 
@@ -915,6 +1026,71 @@ int main() {
 }
 ----
 
+=== Using device globals
+
+This examples demonstrates how a device global defined in runtime-compiled code
+can be accessed from the host and the device.
+
+[source,c++]
+----
+#include <sycl/sycl.hpp>
+namespace syclexp = sycl::ext::oneapi::experimental;
+
+static constexpr size_t NUM = 1024;
+static constexpr size_t WGSIZE = 16;
+
+int main() {
+  sycl::queue q;
+
+  // The source code for a kernel, defined as a SYCL "free function kernel".
+  std::string source = R"""(
+    #include <sycl/sycl.hpp>
+    namespace syclext = sycl::ext::oneapi;
+    namespace syclexp = sycl::ext::oneapi::experimental;
+
+    syclexp::device_global<float> scale;
+
+    extern "C"
+    SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
+    void scaled_iota(float start, float *ptr) {
+      size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
+      ptr[id] = start + scale * static_cast<float>(id);
+    }
+  )""";
+
+  // Create a kernel bundle in "source" state.
+  sycl::kernel_bundle<sycl::bundle_state::ext_oneapi_source> kb_src =
+    syclexp::create_kernel_bundle_from_source(
+      q.get_context(),
+      syclexp::source_language::sycl,
+      source);
+
+  // Compile the kernel.
+  sycl::kernel_bundle<sycl::bundle_state::executable> kb_exe = syclexp::build(kb_src);
+
+  // Initialize the device global.
+  float scale = 0.1f;
+  void *scale_addr =
+    kb_exe.ext_oneapi_get_device_global_address("scale", q.get_device());
+  size_t scale_size = kb_exe.ext_oneapi_get_device_global_size("scale");
+  q.memcpy(scale_addr, &scale, scale_size).wait();
+
+  // Get the kernel via its compiler-generated name, and launch it as before.
+  sycl::kernel scaled_iota = kb_exe.ext_oneapi_get_kernel("scaled_iota");
+
+  float *ptr = sycl::malloc_shared<float>(NUM, q);
+  q.submit([&](sycl::handler &cgh) {
+    // Set the values of the kernel arguments.
+    cgh.set_args(3.14f, ptr);
+
+    // Launch the kernel according to its type, in this case an nd-range kernel.
+    sycl::nd_range ndr{{NUM}, {WGSIZE}};
+    cgh.parallel_for(ndr, scaled_iota);
+  }).wait();
+
+  sycl::free(ptr, q);
+}
+----
 
 == Issues