From 53cdd5e09cb2eac1433bd2dd850431a744043c28 Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Mon, 16 Dec 2024 11:15:41 +0000 Subject: [PATCH] Add sycl_khr_work_item_queries extension This extension allows developers to access instances of the sycl::nd_item, sycl::group and sycl::sub_group classes without having to explicitly pass them as arguments to each function used on the device. --- adoc/extensions/index.adoc | 1 + .../sycl_khr_work_item_queries.adoc | 169 ++++++++++++++++++ 2 files changed, 170 insertions(+) create mode 100644 adoc/extensions/sycl_khr_work_item_queries.adoc diff --git a/adoc/extensions/index.adoc b/adoc/extensions/index.adoc index 1a13364a..e2e98096 100644 --- a/adoc/extensions/index.adoc +++ b/adoc/extensions/index.adoc @@ -11,3 +11,4 @@ specification, but their design is subject to change. // include::sycl_khr_extension_name.adoc[leveloffset=2] include::sycl_khr_default_context.adoc[leveloffset=2] +include::sycl_khr_work_item_queries.adoc[leveloffset=2] diff --git a/adoc/extensions/sycl_khr_work_item_queries.adoc b/adoc/extensions/sycl_khr_work_item_queries.adoc new file mode 100644 index 00000000..fc9c4133 --- /dev/null +++ b/adoc/extensions/sycl_khr_work_item_queries.adoc @@ -0,0 +1,169 @@ +[[sec:khr-work-item-queries]] += SYCL_KHR_WORK_ITEM_QUERIES + +This extension allows developers to access instances of the +[code]#sycl::nd_item#, [code]#sycl::group# and [code]#sycl::sub_group# classes +without having to explicitly pass them as arguments to each function used on the +device. + +{note} Passing such instances as arguments can result in a clearer interface +that is less error-prone to use. +For example, if a function accepts a [code]#sycl::group#, the caller must assume +that function may call a [code]#sycl::group_barrier# and ensure that associated +control flow requirements are satisfied. +It is recommended that this extension is used only when modifying existing +interfaces is not feasible. +{endnote} + +[[sec:khr-work-item-queries-dependencies]] +== Dependencies + +This extension has no dependencies on other extensions. + +[[sec:khr-work-item-queries-feature-test]] +== Feature test macro + +An implementation supporting this extension must predefine the macro +[code]#SYCL_KHR_WORK_ITEM_QUERIES# to one of the values defined in the table +below. + +[%header,cols="1,5"] +|=== +|Value +|Description + +|1 +|Initial version of this extension. +|=== + +[[sec:khr-work-item-queries-free-funcs]] +== New free functions to access instances of special SYCL classes + +This extension adds the following free functions to the +[code]#sycl::khr::this_work_item# namespace, which provide information about the +currently executing work-item. + +It is the user's responsibility to ensure that these functions are called in a +manner that is compatible with the kernel's launch parameters, as detailed in +the definition of each function. +Calling these functions from an incompatible kernel results in undefined +behavior. + +''' + +.[apidef]#khr::this_work_item::get_nd_item# +[source,role=synopsis,id=api:khr-this-work-item-get-nd-item] +---- +namespace sycl::khr::this_work_item { + +template +nd_item get_nd_item(); + +} +---- + +_Preconditions_: [code]#Dimensions# must match the dimensionality of the +currently executing kernel. +The currently executing kernel must have been launched with a +[code]#sycl::nd_range# argument. + +_Returns_: A [code]#sycl::nd_item# instance representing the calling work-item +in the [code]#sycl::nd_range#. + +''' + +.[apidef]#khr::this_work_item::get_work_group# +[source,role=synopsis,id=api:khr-this-work-item-get-work-group] +---- +namespace sycl::khr::this_work_item { + +template +group get_work_group(); + +} +---- + +_Preconditions_: [code]#Dimensions# must match the dimensionality of the +currently executing kernel. +The currently executing kernel must have been launched with a +[code]#sycl::nd_range# argument. + +_Returns_: A [code]#sycl::group# instance representing the <> to +which the calling work-item belongs. + +''' + +.[apidef]#khr::this_work_item::get_sub_group# +[source,role=synopsis,id=api:khr-this-work-item-get-sub-group] +---- +namespace sycl::khr::this_work_item { + +sub_group get_sub_group(); + +} +---- + +_Preconditions_: The currently executing kernel must have been launched with a +[code]#sycl::nd_range# argument. + +_Returns_: A [code]#sycl::sub_group# instance representing the <> to +which the calling work-item belongs. + +''' + +[[sec:khr-work-item-queries-example]] +== Example + +The example below demonstrates the usage of this extension with a simple kernel +calling a device function. + +[source,,linenums] +---- +#include +#include +#include +#include +using namespace sycl; // (optional) avoids need for "sycl::" before SYCL names + +void vector_add(float* a, float* b, float* c) +{ + // Access this work-item's nd_item class directly. + size_t i = khr::this_work_item::get_nd_item<1>().get_global_linear_id(); + + c[i] = a[i] + b[i]; +} + +constexpr size_t N = 1024; + +int main() { + + queue q; + + float* a = malloc_shared(N, q); + float* b = malloc_shared(N, q); + float* c = malloc_shared(N, q); + + std::iota(a, a + N, 0); + std::iota(b, b + N, 0); + std::fill(c, c + N, 0); + + range<1> global{N}; + range<1> local{32}; + q.parallel_for(nd_range<1>{global, local}, [=](nd_item<1> it) { + // Function does not take nd_item as an argument. + vector_add(a, b, c); + }); + + std::cout << std::endl << "Result:" << std::endl; + for (size_t i = 0; i < N; i++) { + if (c[i] != a[i] + b[i]) { + std::cout << "Wrong value " << c[i] << " on element " << i << std::endl; + exit(-1); + } + } + + std::cout << "Good computation!" << std::endl; + return 0; + +} +----