Skip to content

Implement sycl_khr_work_item_queries extension #18519

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
wants to merge 2 commits into
base: sycl
Choose a base branch
from
Open

Implement sycl_khr_work_item_queries extension #18519

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
a96112b
Implement `sycl_khr_work_item_queries` extension
0x12CC May 16, 2025
bfb2d89
Add license to extension header
0x12CC May 16, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
36 changes: 36 additions & 0 deletions sycl/include/sycl/khr/work_item_queries.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,36 @@
//===-- work_item_queries.hpp --- KHR work item queries extension ---------===//
//
// 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

#ifdef __DPCPP_ENABLE_UNFINISHED_KHR_EXTENSIONS

#include <sycl/ext/oneapi/free_function_queries.hpp>

#define SYCL_KHR_WORK_ITEM_QUERIES 1
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I see this is how we do it in khr/group_interface.hpp too, but would it be more appropriate to guard the definition in sycl/source/feature_test.hpp.in together with all the other feature test macros?

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I agree on this because there are already two khr extension macros set within sycl/source/feature_test.hpp.in


namespace sycl {
inline namespace _V1 {
namespace khr {

template <int Dimensions> nd_item<Dimensions> this_nd_item() {
return ext::oneapi::experimental::this_nd_item<Dimensions>();
}

template <int Dimensions> group<Dimensions> this_group() {
return ext::oneapi::this_work_item::get_work_group<Dimensions>();
}

inline sub_group this_sub_group() {
return ext::oneapi::this_work_item::get_sub_group();
}

} // namespace khr
} // namespace _V1
} // namespace sycl

#endif
1 change: 1 addition & 0 deletions sycl/include/sycl/sycl.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -124,3 +124,4 @@
#include <sycl/ext/oneapi/virtual_mem/virtual_mem.hpp>
#include <sycl/ext/oneapi/weak_object.hpp>
#include <sycl/khr/group_interface.hpp>
#include <sycl/khr/work_item_queries.hpp>
99 changes: 99 additions & 0 deletions sycl/test-e2e/Basic/work_item_queries/work_item_queries.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,99 @@
// RUN: %{build} -o %t.out
// RUN: %{run} %t.out

//===- work_item_queries.cpp - KHR work item queries test -----------------===//
//
// 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
//
//===----------------------------------------------------------------------===//

#define __DPCPP_ENABLE_UNFINISHED_KHR_EXTENSIONS

#include <cassert>
#include <sycl/detail/core.hpp>
#include <sycl/khr/work_item_queries.hpp>

template <size_t... Dims> static void check_this_nd_item_api() {
// Define the kernel ranges.
constexpr int Dimensions = sizeof...(Dims);
const sycl::range<Dimensions> local_range{Dims...};
const sycl::range<Dimensions> global_range = local_range;
const sycl::nd_range<Dimensions> nd_range{global_range, local_range};
// Launch an ND-range kernel.
sycl::queue q;
sycl::buffer<bool, Dimensions> results{global_range};
q.submit([&](sycl::handler &cgh) {
sycl::accessor acc{results, cgh, sycl::write_only};
cgh.parallel_for(nd_range, [=](sycl::nd_item<Dimensions> it) {
// Compare it to this_nd_item<Dimensions>().
acc[it.get_global_id()] = (it == sycl::khr::this_nd_item<Dimensions>());
});
});
// Check the test results.
sycl::host_accessor acc{results};
for (const auto &result : acc)
assert(result);
Comment on lines +36 to +37
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Nit; Asserts stop execution immediately, so personally I prefer to do something like:

Suggested change
for (const auto &result : acc)
assert(result);
for (const auto &result : acc) {
if (!result) {
std::cout << "Failed check_this_nd_item_api check for dimensionality " << Dimensions << ".\n";
return 1;
}
}
return 0;

and then in the main function you do:

  int Failed = 0;
  // nd_item
  Failed += check_this_nd_item_api<2>();
  Failed += check_this_nd_item_api<2, 3>();
  Failed += check_this_nd_item_api<2, 3, 4>();
  // group
  Failed += check_this_group_api<2>();
  Failed += check_this_group_api<2, 3>();
  Failed += check_this_group_api<2, 3, 4>();
  // sub_group
  Failed += check_this_sub_group_api<2>();
  Failed += check_this_sub_group_api<2, 3>();
  Failed += check_this_sub_group_api<2, 3, 4>();

  return Failed;

The more information, the better.

}

template <size_t... Dims> static void check_this_group_api() {
// Define the kernel ranges.
constexpr int Dimensions = sizeof...(Dims);
const sycl::range<Dimensions> local_range{Dims...};
const sycl::range<Dimensions> global_range = local_range;
const sycl::nd_range<Dimensions> nd_range{global_range, local_range};
// Launch an ND-range kernel.
sycl::queue q;
sycl::buffer<bool, Dimensions> results{global_range};
q.submit([&](sycl::handler &cgh) {
sycl::accessor acc{results, cgh, sycl::write_only};
cgh.parallel_for(nd_range, [=](sycl::nd_item<Dimensions> it) {
// Compare it.get_group() to this_group<Dimensions>().
acc[it.get_global_id()] =
(it.get_group() == sycl::khr::this_group<Dimensions>());
});
});
// Check the test results.
sycl::host_accessor acc{results};
for (const auto &result : acc)
assert(result);
}

template <size_t... Dims> static void check_this_sub_group_api() {
// Define the kernel ranges.
constexpr int Dimensions = sizeof...(Dims);
const sycl::range<Dimensions> local_range{Dims...};
const sycl::range<Dimensions> global_range = local_range;
const sycl::nd_range<Dimensions> nd_range{global_range, local_range};
// Launch an ND-range kernel.
sycl::queue q;
sycl::buffer<bool, Dimensions> results{global_range};
q.submit([&](sycl::handler &cgh) {
sycl::accessor acc{results, cgh, sycl::write_only};
cgh.parallel_for(nd_range, [=](sycl::nd_item<Dimensions> it) {
// Compare it.get_sub_group() to this_sub_group().
acc[it.get_global_id()] =
(it.get_sub_group() == sycl::khr::this_sub_group());
});
});
// Check the test results.
sycl::host_accessor acc{results};
for (const auto &result : acc)
assert(result);
}

int main() {
// nd_item
check_this_nd_item_api<2>();
check_this_nd_item_api<2, 3>();
check_this_nd_item_api<2, 3, 4>();
// group
check_this_group_api<2>();
check_this_group_api<2, 3>();
check_this_group_api<2, 3, 4>();
// sub_group
check_this_sub_group_api<2>();
check_this_sub_group_api<2, 3>();
check_this_sub_group_api<2, 3, 4>();
}