Skip to content
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.

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

sycl::kernel_handler::get_specialization_constant calls __sycl_getComposite2020SpecConstantValue on enumerated value_type #12703

Closed
victor-eds opened this issue Feb 13, 2024 · 5 comments
Assignees
Labels
bug Something isn't working confirmed

Comments

@victor-eds
Copy link
Contributor

victor-eds commented Feb 13, 2024

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 victor-eds added the bug Something isn't working label Feb 13, 2024
@maksimsab
Copy link
Contributor

@victor-eds please, attach the code.

@victor-eds
Copy link
Contributor Author

@victor-eds please, attach the code.

Added inline

@AlexeySachkov
Copy link
Contributor

Some massaging of our enable_ifs is required here:

#ifdef __SYCL_DEVICE_ONLY__
template <
auto &S,
typename T = typename std::remove_reference_t<decltype(S)>::value_type,
std::enable_if_t<std::is_fundamental_v<T>> * = nullptr>
__SYCL_ALWAYS_INLINE T getSpecializationConstantOnDevice() {
const char *SymbolicID = __builtin_sycl_unique_stable_id(S);
return __sycl_getScalar2020SpecConstantValue<T>(
SymbolicID, &S, MSpecializationConstantsBuffer);
}
template <
auto &S,
typename T = typename std::remove_reference_t<decltype(S)>::value_type,
std::enable_if_t<std::is_compound_v<T>> * = nullptr>
__SYCL_ALWAYS_INLINE T getSpecializationConstantOnDevice() {
const char *SymbolicID = __builtin_sycl_unique_stable_id(S);
return __sycl_getComposite2020SpecConstantValue<T>(
SymbolicID, &S, MSpecializationConstantsBuffer);
}
#endif // __SYCL_DEVICE_ONLY__

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++

@victor-eds
Copy link
Contributor Author

Some massaging of our enable_ifs is required here:

#ifdef __SYCL_DEVICE_ONLY__
template <
auto &S,
typename T = typename std::remove_reference_t<decltype(S)>::value_type,
std::enable_if_t<std::is_fundamental_v<T>> * = nullptr>
__SYCL_ALWAYS_INLINE T getSpecializationConstantOnDevice() {
const char *SymbolicID = __builtin_sycl_unique_stable_id(S);
return __sycl_getScalar2020SpecConstantValue<T>(
SymbolicID, &S, MSpecializationConstantsBuffer);
}
template <
auto &S,
typename T = typename std::remove_reference_t<decltype(S)>::value_type,
std::enable_if_t<std::is_compound_v<T>> * = nullptr>
__SYCL_ALWAYS_INLINE T getSpecializationConstantOnDevice() {
const char *SymbolicID = __builtin_sycl_unique_stable_id(S);
return __sycl_getComposite2020SpecConstantValue<T>(
SymbolicID, &S, MSpecializationConstantsBuffer);
}
#endif // __SYCL_DEVICE_ONLY__

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
Copy link
Contributor Author

Fixed by #12925

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
bug Something isn't working confirmed
Projects
None yet
Development

No branches or pull requests

3 participants