[SYCL][CUDA][HIP] Throw a runtime error with invalid sub-group size to kernel

[abagusetty]

Describe the bug
Using an invalid sub-group size to the kernel doesn't result in a runtime error. Possibly throw a RT error as seen on Intel devices with invalid sub-group sizes. This behavior was observed for both the CUDA and HIP Plugins.

For CUDA possibly 32, for HIP possibly 64.

From the SYCL spec:
Each device supports only certain sub-group sizes as defined by info::device::sub_group_sizes. In addition, some device features may be incompatible with certain sub-group sizes. If a kernel is decorated with this attribute and then submitted to a device that does not support the sub-group size or if the kernel uses a feature that the device does not support with this sub-group size, the implementation must throw a synchronous exception with the errc::kernel_not_supported error code.

To Reproduce

#include <sycl/sycl.hpp>

int main() {
  auto const& gpu_devices = sycl::device::get_devices(sycl::info::device_type::gpu);
  std::cout << "Number of Root GPUs: " << gpu_devices.size() << std::endl;

  for(const auto& d : gpu_devices) {
    std::cout << "Found Root GPU-ID: " << d.get_info<sycl::info::device::name>() << std::endl;
    std::vector<size_t> sg_sizes = d.get_info<sycl::info::device::sub_group_sizes>();
    std::cout << "Supported sub-group sizes: ";
    for (int i=0; i<sg_sizes.size(); i++) {
      std::cout << sg_sizes[i];
    }
    std::cout << std::endl;
  }

  const int N{1024};
  sycl::queue Q{sycl::gpu_selector{}};
  int* ptr = sycl::malloc_device<int>(N, Q);
  Q.parallel_for(N, [=](sycl::item<1> id) [[sycl::reqd_sub_group_size(8)]] { ptr[id] = id; }).wait();
  return 0;
}

Environment (please complete the following information):
CUDA PI built with CUDA-11.6.2, HIP PI built with room-5.1.0
llvm compiler: ac6a4f5

[JackAKirk]

CUDA only currently supports one subgroup (warp) size : 32 for all devices.
As detailed in 5.7 of the SYCL 2020 Revision 5 spec, reqd_sub_group_size() is an optional kernel feature which means it does not have to be supported for all devices but the spec states:

In order to guarantee source code portability of SYCL applications that use optional kernel features, all SYCL implementations must be able to compile device code that uses these optional features regardless of whether the implementation supports the features on any of its devices.

This statement is unclear for me. Should a SYCL implementation be able to compile device code using reqd_sub_group_size for all backends, or only those in which reqd_sub_group_size can be sensibly implemented/has a use.

A problem with such code being able to compile even when a reqd subgroup size that is specified with reqd_sub_group_size() is invalid for the given device (that doesn't support the reqd_sub_group_size optional kernel feature), is that code might then run but give unexpected results if the user is assuming a different subgroup size to the one actually used on the device: an example is: #5971.

I think that #6183 introduces a sensible solution for backends which only support a single subgroup size: if reqd_sub_group_size() is used with the supported subgroup size then it will compile and behave as the programmer intends. If reqd_sub_group_size() is used with another incompatible subgroup size the following error is returned when compiling:

invalid_sg_orig.cpp:20:21: error: CUDA backend requires subgroup size 32
  Q.parallel_for(N, [=](sycl::item<1> id) [[sycl::reqd_sub_group_size(8)]] { ptr[id] = id; }).wait();
                    ^
1 error generated.

The other alternative would be to allow reqd_sub_group_size with an invalid subgroup size to compile but throw errc::ker­nel_not_supported at runtime. A potential technical difficulty with this is that, since CUDA only supports a single subgroup and doesn't throw any corresponding device errors for an incorrect subgroup size, it isn't possible to progagate an error from the device to the runtime via the PI_Plugin (which may be how it would be dealt with in other backends). The only other solution I imagine would be to attached a subgroup attribute to a kernel, that can be checked prior to a kernel submission. This brings me to another point: I observe that, using OpenCL CPU with the sycl::reqd_sub_group_size kernel label using some invalid subgroup sizes, I get the following:

Failed to build: : -11 (CL_BUILD_PROGRAM_FAILURE)

llvm-foreach: 
clang-15: error: x86_64 compiler command failed with exit code 245 (use -v to see invocation)

With some other invalid subgroup sizes the code compiles and executes with no errors. I am not sure if the implementation is complete (i.e. reqd_sub_group_size(val) leads to errc::ker­nel_not_supported if val is an unsupported size) on another backend that I don't have access to.

For the HIP AMD case older devices only support subgroup (wavefront) of 64, but newer "RDNA" cards also support 32. @npmiller do you know if DPC++ HIP AMD supports both subgroup sizes for the "RDNA" cards?
If DPC++ only currently supports the older 64 subgroup size on all devices then it may be sensible to apply a similar solution to #6183 for the HIP AMD, until a time when both subgroup sizes are supported and we should also implement reqd_sub_group_size for the HIP AMD backend? At the moment the metadata that is attached via reqd_sub_group_size is apparently ignored in the HIP AMD backend and only used in the OpenCL backend. Presumably the level_zero backend also does not support reqd_sub_group_size at this time, and would result in similar issues to those raised here?

[npmiller]

We do support both sizes, in fact RDNA cards will default to the 32 thread wavefront size, and I believe that's what the CI hardware does.

[AerialMantis]

So the motivation for the spec to require that all SYCL applications compile, even if a device may not support a feature at runtime, is for portability, so an application will always compile as long as it's valid SYCL code. I agree it does seem a little counter intuitive for sycl::reqd_sub_group_size to compile for sub-groups sizes which are known to be invalid since this will fail at runtime and the compiler may know that it cannot support that size, however, you may want to compile for multiple backends so in that case the compiler can't know which binary will be used.

It might be worth emitting a compiler warning for this, to help catch the case where an invalid sub-group size is being used by mistake rather than intentionally?

[JackAKirk]

So the motivation for the spec to require that all SYCL applications compile, even if a device may not support a feature at runtime, is for portability, so an application will always compile as long as it's valid SYCL code. I agree it does seem a little counter intuitive for sycl::reqd_sub_group_size to compile for sub-groups sizes which are known to be invalid since this will fail at runtime and the compiler may know that it cannot support that size, however, you may want to compile for multiple backends so in that case the compiler can't know which binary will be used.

It might be worth emitting a compiler warning for this, to help catch the case where an invalid sub-group size is being used by mistake rather than intentionally?

Yep Compiler warning added here: #6183

[JackAKirk]

Update:

reqd_sub_group_size() kernel attribute does not have a complete implementation in any backend: it does not throw the errc::kernel_not_­supported for invalid cases as described in the SYCL 2020 spec in any backend. It sounds like there is a plan to implement a general solution for this in the works.

In the meantime reqd_sub_group_size() appears to have a partially working implementation for OpenCL (I did not test it fully: I described my experiments in an earlier comment), but for other backends the kernel metadata is unused.

A warning has been added for the CUDA case in this PR: #6183
Hopefully this will warn users of potential unexpected behaviour when attempting to port OpenCL code to CUDA that makes use of the reqd_sub_group_size kernel attribute.

[abagusetty]

I think, I came across invalid sub-group size set to the kernel attributes error at RT when running on a Intel Gen9, ATS, PVC devices with L0 backend specifically. Might be worth taking a peek at L0 ?

[JackAKirk]

I think, I came across invalid sub-group size set to the kernel attributes error at RT when running on a Intel Gen9, ATS, PVC devices with L0 backend specifically. Might be worth taking a peek at L0 ?

@gmlueck Do you know if this is the case? Could you clarify the current support for reqd_sub_group_size() in L0?
Thanks

[abagusetty]

This is a sample output from SYCL on Intel Gen9 (support SG size: 8,16,32) using L0 backend.
Error generated with: Q.parallel_for(N, [=](sycl::item<1> id) [[sycl::reqd_sub_group_size(7)]] { ptr[id] = id; }).wait();

abagusetty@iris09:~/soft/sycl_req_sub_group_size$ dpcpp sycl_free_mem.cpp
abagusetty@iris09:~/soft/sycl_req_sub_group_size$ ./a.out 
Number of Root GPUs: 1
Found Root GPU-ID: Intel(R) Iris(R) Pro Graphics P580 [0x193a]
Supported sub-group sizes: 81632
terminate called after throwing an instance of 'cl::sycl::compile_program_error'
  what():  The program was built for 1 devices
Build program log for 'Intel(R) Iris(R) Pro Graphics P580 [0x193a]':

error: Unsupported required sub group size
in kernel: 'typeinfo name for main::'lambda'(cl::sycl::item<1, true>)'
error: backend compiler failed build.
 -11 (CL_BUILD_PROGRAM_FAILURE)
Aborted

[JackAKirk]

This is a sample output from SYCL on Intel Gen9 (support SG size: 8,16,32) using L0 backend. Error generated with: Q.parallel_for(N, [=](sycl::item<1> id) [[sycl::reqd_sub_group_size(7)]] { ptr[id] = id; }).wait();

abagusetty@iris09:~/soft/sycl_req_sub_group_size$ ./a.out 
Number of Root GPUs: 1
Found Root GPU-ID: Intel(R) Iris(R) Pro Graphics P580 [0x193a]
Supported sub-group sizes: 81632
terminate called after throwing an instance of 'cl::sycl::compile_program_error'
  what():  The program was built for 1 devices
Build program log for 'Intel(R) Iris(R) Pro Graphics P580 [0x193a]':

error: Unsupported required sub group size
in kernel: 'typeinfo name for main::'lambda'(cl::sycl::item<1, true>)'
error: backend compiler failed build.
 -11 (CL_BUILD_PROGRAM_FAILURE)
Aborted

That's at compile time right? Same error that I observed for the OpenCL backend (incidentally using the same sub_group size, 7). I also observed that if you try very large subgroup sizes (> 64 I think) or very small subgroup sizes (e.g. size 1) then OpenCL backend will compile and run without any errors, presumably with a default subgroup size.

[abagusetty]

That's at compile time right? Same error that I observed for the OpenCL backend. I also observed that if you try very large subgroup sizes (> 64 I think) or very small subgroup sizes (e.g. size 1) then OpenCL backend will compile and run without any errors, presumably with a default subgroup size.

The error was at Runtime, the compilation didn't really result in any warning or error. Also I check with using -Wsycl-strict option if that has any guards but no-avail