Skip to content

[SYCL][Clang] Add sentinel value to kernel_args_sizes #20540

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

Merged
steffenlarsen merged 4 commits into from
Nov 21, 2025
Merged

[SYCL][Clang] Add sentinel value to kernel_args_sizes #20540

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
fb94687
[SYCL][Clang] Add sentinel value to kernel_args_sizes
elizabethandrews Nov 3, 2025
e1e1c74
Change sentinel value to 0 to avoid type change and ABI break
elizabethandrews Nov 5, 2025
8728f65
Use std::numeric_limits<uint32_t>::max()
elizabethandrews Nov 12, 2025
ac51b4d
clang format fix
elizabethandrews Nov 12, 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
6 changes: 5 additions & 1 deletion clang/lib/Sema/SemaSYCL.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -7052,11 +7052,15 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
O << " \"\",\n";
O << "};\n\n";

O << "static constexpr unsigned kernel_args_sizes[] = {";
O << "static constexpr int kernel_args_sizes[] = {";
Copy link
Contributor

@Fznamznon Fznamznon Nov 4, 2025

Choose a reason for hiding this comment

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

Is that array used at all? I can't grep it in sycl folder.

Copy link
Contributor Author

@elizabethandrews elizabethandrews Nov 4, 2025

Choose a reason for hiding this comment

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

I'm not sure actually. @aelovikov-intel would you know if this is just a relic of the past or whether we actually need this generated?

Copy link
Contributor

@Fznamznon Fznamznon Nov 4, 2025

Choose a reason for hiding this comment

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

CI says

  # | C:\Users\GH_RUN~1\AppData\Local\Temp\lit-tmp-qc0j5t24\L0_interop_test-header-a28efe.h:82:75: error: cannot initialize a parameter of type 'const unsigned int *' with an lvalue of type 'const int[2]'
  # |    82 |     sycl::detail::free_function_info_map::add(sycl::detail::kernel_names, sycl::detail::kernel_args_sizes, 1);
  # |       |                                                                           ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~

so, this seems to be used within the integration header?

Copy link
Contributor

@aelovikov-intel aelovikov-intel Nov 4, 2025

Choose a reason for hiding this comment

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

+ @dklochkov-emb

Copy link
Contributor Author

@elizabethandrews elizabethandrews Nov 4, 2025

Choose a reason for hiding this comment

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

Yea it looks like this is used for the implementation of free functions/device globals.

Copy link
Contributor

@dklochkov-emb dklochkov-emb Nov 5, 2025

Choose a reason for hiding this comment

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

@elizabethandrews @aelovikov-intel yes, it is used in free functions:
https://github.com/intel/llvm/blob/sycl/clang/lib/Sema/SemaSYCL.cpp#L7055
https://github.com/intel/llvm/blob/sycl/clang/lib/Sema/SemaSYCL.cpp#L7317
https://github.com/intel/llvm/blob/sycl/clang/lib/Sema/SemaSYCL.cpp#L7323

llvm/sycl/source/detail/kernel_global_info.cpp

Line 13 in e05e82b

namespace detail::free_function_info_map {

If it is needed to change the type of kernel_args_sizes, then it is needed to change the whole chain..

Also, it is supposed that both arrays kernel arg sizes and kernel names have the same size. If -1 is added into kernel arg sizes, please, check which size is used during add and remove calls

Copy link
Contributor Author

@elizabethandrews elizabethandrews Nov 5, 2025

Choose a reason for hiding this comment

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

Changing the type would be an ABI break. So I have change the PR to use a sentinel value of 0 so we can keep it unsigned. I don't see any harm with that but please let me know if you disagree.

Copy link
Contributor

@Fznamznon Fznamznon Nov 6, 2025 •
edited
Loading

Choose a reason for hiding this comment

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

I'm slightly concerned that we can actually put 0 for some arguments there. A capturless lambda is an example.
I think alternatively we could put there something like std::numeric_limits<uint32_t>::max(); Hopefully there is no argument of that size :)

Copy link
Contributor Author

@elizabethandrews elizabethandrews Nov 6, 2025

Choose a reason for hiding this comment

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

Yes I had talked to @premanandrao about this. You're right in that there is an overlap about what 0 means here. Kernels without captures and no kernels will both have 0 here now. We thought it shouldn't matter since hopefully the runtime will have some other way of detecting no kernels. Maybe @dklochkov-emb can weigh in here?

Copy link
Contributor

@dklochkov-emb dklochkov-emb Nov 7, 2025

Choose a reason for hiding this comment

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

I think, using std::numeric_limits<uint32_t>::max(); is better, it looks reasonable and large enough to avoid collisions

for (unsigned I = 0; I < KernelDescs.size(); I++) {
O << KernelDescs[I].Params.size() << ", ";
}
// Add a sentinel to avoid warning if the collection is empty
// (similar to what we do for kernel_signatures below).
O << "-1, \n";
O << "};\n\n";

O << "// array representing signatures of all kernels defined in the\n";
O << "// corresponding source\n";
O << "static constexpr\n";
Expand Down
2 changes: 1 addition & 1 deletion clang/test/CodeGenSYCL/int-header-empty-signatures.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10,7 +10,7 @@
// CHECK-NEXT: ""
// CHECK-NEXT: };

// CHECK: static constexpr unsigned kernel_args_sizes[] = {
// CHECK: static constexpr int kernel_args_sizes[] = {0, -1

// CHECK: static constexpr
// CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = {
Expand Down
2 changes: 1 addition & 1 deletion clang/test/CodeGenSYCL/integration_header.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -25,7 +25,7 @@
// CHECK-NEXT: ""
// CHECK-NEXT: };
//
// CHECK: static constexpr unsigned kernel_args_sizes[] = {
// CHECK: static constexpr int kernel_args_sizes[] = {5, 3, 2, 8, 3, 0, -1,
//
// CHECK: static constexpr
// CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = {
Expand Down
2 changes: 1 addition & 1 deletion clang/test/CodeGenSYCL/kernel-param-acc-array-ih.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18,7 +18,7 @@
// CHECK-NEXT: ""
// CHECK-NEXT: };

// CHECK: static constexpr unsigned kernel_args_sizes[] = {
// CHECK: static constexpr int kernel_args_sizes[] = {2, -1,

// CHECK: static constexpr
// CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = {
Expand Down
2 changes: 1 addition & 1 deletion clang/test/CodeGenSYCL/kernel-param-member-acc-array-ih.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18,7 +18,7 @@
// CHECK-NEXT: ""
// CHECK-NEXT: };

// CHECK: static constexpr unsigned kernel_args_sizes[] = {
// CHECK: static constexpr int kernel_args_sizes[] = {2, -1,

// CHECK: static constexpr
// CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = {
Expand Down
2 changes: 1 addition & 1 deletion clang/test/CodeGenSYCL/kernel-param-pod-array-ih.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -19,7 +19,7 @@
// CHECK-NEXT: ""
// CHECK-NEXT: };

// CHECK: static constexpr unsigned kernel_args_sizes[] = {
// CHECK: static constexpr int kernel_args_sizes[] = {1, 1, 1, -1,

// CHECK: static constexpr
// CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = {
Expand Down
2 changes: 1 addition & 1 deletion clang/test/CodeGenSYCL/union-kernel-param-ih.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18,7 +18,7 @@
// CHECK-NEXT: ""
// CHECK-NEXT: };

// CHECK: static constexpr unsigned kernel_args_sizes[] = {
// CHECK: static constexpr int kernel_args_sizes[] = {1, -1,

// CHECK: static constexpr
// CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = {
Expand Down
2 changes: 1 addition & 1 deletion clang/test/CodeGenSYCL/wrapped-accessor.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -15,7 +15,7 @@
// CHECK-NEXT: ""
// CHECK-NEXT: };

// CHECK: static constexpr unsigned kernel_args_sizes[] = {
// CHECK: static constexpr int kernel_args_sizes[] = {1, -1,

// CHECK: static constexpr
// CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = {
Expand Down
Loading