llvm icon indicating copy to clipboard operation
llvm copied to clipboard

`sycl::kernel_handler::get_specialization_constant` calls `__sycl_getComposite2020SpecConstantValue` on enumerated `value_type`

Open victor-eds opened this issue 1 year ago • 4 comments

Describe the bug

sycl::kernel_handler::get_specialization_constant calls __sycl_getComposite2020SpecConstantValue instead of __sycl_getScalar2020SpecConstantValue. Enumerated types have the same size, value representation and alignment requirements as their underlying integral types, so we could be simply using the scalar version of the builtin.

To Reproduce

Compile the following code:

#include <sycl/sycl.hpp>

enum class my_enum { a, b, c };

constexpr sycl::specialization_id<my_enum> val{my_enum::a};

SYCL_EXTERNAL my_enum f(sycl::kernel_handler kh) {
  return kh.get_specialization_constant<val>();
}
clang++ -fsycl -fsycl-device-only -o - -S -emit-llvm -O3

You'll see the @_Z40__sycl_getComposite2020SpecConstantValueI7my_enumET_PKcPKvS5_ function being called to get the specialization constant value instead of the scalar version of the function.

Environment (please complete the following information):

  • DPC++ version: effbbabcd6e3

Additional context

This should have the exact same semantics, but using the scalar builtin leads to a simpler compilation path.

Changing the check in kernel_handler.hpp so that the scalar version accepts enumerated types might be enough.

victor-eds avatar Feb 13 '24 11:02 victor-eds

@victor-eds please, attach the code.

maksimsab avatar Feb 13 '24 16:02 maksimsab

@victor-eds please, attach the code.

Added inline

victor-eds avatar Feb 14 '24 08:02 victor-eds

Some massaging of our enable_ifs is required here:

https://github.com/intel/llvm/blob/5fae0aaa4f932c6ecddb052a0cb3d873ff2eee64/sycl/include/sycl/kernel_handler.hpp#L66-L85

We have a simple switch between compound and fundamental types, but we should probably use less C++-y terms like composite and scalar to treat more data types as scalar ones, even though they are compound per C++

AlexeySachkov avatar Feb 20 '24 14:02 AlexeySachkov

Some massaging of our enable_ifs is required here:

https://github.com/intel/llvm/blob/5fae0aaa4f932c6ecddb052a0cb3d873ff2eee64/sycl/include/sycl/kernel_handler.hpp#L66-L85

We have a simple switch between compound and fundamental types, but we should probably use less C++-y terms like composite and scalar to treat more data types as scalar ones, even though they are compound per C++

According to cppreference, enums are scalar types. Maybe using std::is_scalar would fix this?

victor-eds avatar Feb 20 '24 15:02 victor-eds

Fixed by #12925

victor-eds avatar Mar 08 '24 10:03 victor-eds