[SYCL][Clang] Add sentinel value to kernel_args_sizes

[elizabethandrews]

Add sentinel value to compiler generated array kernel_args_sizes in integration header to avoid warnings about empty arrays. This array is used to hold the number of elements in each kernel. This PR changes the type to int to add a sentinel value of std::numeric_limits<uint32_t>::max().

[Fznamznon]

CI fails might be related.

[Fznamznon]

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

[elizabethandrews]

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?

[Fznamznon]

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?

[aelovikov-intel]

+ @dklochkov-emb

[elizabethandrews]

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

[dklochkov-emb]

@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

[elizabethandrews]

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.

[Fznamznon]

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 :)

[elizabethandrews]

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?

[dklochkov-emb]

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

[elizabethandrews]

@intel/llvm-gatekeepers How do I re-run the failed test. I tried to do re-run jobs and get an error about "failing to download artifact"

[sarnex]

That error means too much time has passed between when build was run and now so we don't have the build artifacts anymore. You can restart the build job to fix it, that will also rerun the tests. I'll restart it.

[elizabethandrews]

@intel/llvm-gatekeepers PR is ready for merge