Skip to content

Commit

Permalink
Do not bother with sycl::rotate
Browse files Browse the repository at this point in the history
  • Loading branch information
dalg24 committed Feb 23, 2023
1 parent 22ee14e commit fcdedf7
Showing 1 changed file with 2 additions and 32 deletions.
34 changes: 2 additions & 32 deletions core/src/Kokkos_BitManipulation.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -304,34 +304,6 @@ KOKKOS_IMPL_HOST_FUNCTION

#undef KOKKOS_IMPL_USE_GCC_BUILT_IN_FUNCTIONS

template <class T>
KOKKOS_FUNCTION T rotl_builtin_host(T x, int s) noexcept {
return rotl(x, s);
}

template <class T>
KOKKOS_FUNCTION T rotl_builtin_device(T x, int s) noexcept {
#ifdef KOKKOS_ENABLE_SYCL
return sycl::rotate(x, s);
#else
return rotl(x, s);
#endif
}

template <class T>
KOKKOS_FUNCTION T rotr_builtin_host(T x, int s) noexcept {
return rotr(x, s);
}

template <class T>
KOKKOS_FUNCTION T rotr_builtin_device(T x, int s) noexcept {
#ifdef KOKKOS_ENABLE_SYCL
return sycl::rotate(x, -s);
#else
return rotr(x, s);
#endif
}

} // namespace Kokkos::Impl

namespace Kokkos::Experimental {
Expand Down Expand Up @@ -411,16 +383,14 @@ template <class T>
[[nodiscard]] KOKKOS_FUNCTION
std::enable_if_t<::Kokkos::Impl::is_standard_unsigned_integer_type_v<T>, T>
rotl_builtin(T x, int s) noexcept {
KOKKOS_IF_ON_DEVICE((return ::Kokkos::Impl::rotl_builtin_device(x, s);))
KOKKOS_IF_ON_HOST((return ::Kokkos::Impl::rotl_builtin_host(x, s);))
return rotl(x, s); // no benefit to call the _builtin variant
}

template <class T>
[[nodiscard]] KOKKOS_FUNCTION
std::enable_if_t<::Kokkos::Impl::is_standard_unsigned_integer_type_v<T>, T>
rotr_builtin(T x, int s) noexcept {
KOKKOS_IF_ON_DEVICE((return ::Kokkos::Impl::rotr_builtin_device(x, s);))
KOKKOS_IF_ON_HOST((return ::Kokkos::Impl::rotr_builtin_host(x, s);))
return rotr(x, s); // no benefit to call the _builtin variant
}

} // namespace Kokkos::Experimental
Expand Down

0 comments on commit fcdedf7

Please sign in to comment.