Skip to content

Commit

Permalink
Fix NVCC warnings (kokkos#6483)
Browse files Browse the repository at this point in the history
* Fix NVCC warnings

* Use ternary instead of Kokkos::min
  • Loading branch information
masterleinad committed Oct 5, 2023
1 parent 4ce289b commit 890148e
Show file tree
Hide file tree
Showing 5 changed files with 20 additions and 26 deletions.
7 changes: 3 additions & 4 deletions algorithms/src/sorting/Kokkos_BinSortPublicAPI.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -387,10 +387,9 @@ class BinSort {
// Switching to std::sort for more than 10 elements has been found
// reasonable experimentally.
if (use_std_sort && bin_size > 10) {
if constexpr (use_std_sort) {
std::sort(&sort_order(lower_bound), &sort_order(upper_bound),
[this](int p, int q) { return bin_op(keys_rnd, p, q); });
}
KOKKOS_IF_ON_HOST(
(std::sort(&sort_order(lower_bound), &sort_order(upper_bound),
[this](int p, int q) { return bin_op(keys_rnd, p, q); });))
} else {
for (int k = lower_bound + 1; k < upper_bound; ++k) {
int old_idx = sort_order(k);
Expand Down
3 changes: 2 additions & 1 deletion algorithms/unit_tests/TestStdAlgorithmsCommon.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -198,7 +198,8 @@ auto create_deep_copyable_compatible_view_with_same_extent(ViewType view) {

// this is needed for intel to avoid
// error #1011: missing return statement at end of non-void function
#if defined KOKKOS_COMPILER_INTEL
#if defined KOKKOS_COMPILER_INTEL || \
(defined(KOKKOS_COMPILER_NVCC) && KOKKOS_COMPILER_NVCC >= 1130)
__builtin_unreachable();
#endif
}
Expand Down
8 changes: 4 additions & 4 deletions core/src/Kokkos_Array.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -234,14 +234,14 @@ struct Array<T, KOKKOS_INVALID_INDEX, Array<>::contiguous> {

KOKKOS_INLINE_FUNCTION
Array& operator=(const Array& rhs) {
const size_t n = std::min(m_size, rhs.size());
const size_t n = size() < rhs.size() ? size() : rhs.size();
for (size_t i = 0; i < n; ++i) m_elem[i] = rhs[i];
return *this;
}

template <size_t N, class P>
KOKKOS_INLINE_FUNCTION Array& operator=(const Array<T, N, P>& rhs) {
const size_t n = std::min(m_size, rhs.size());
const size_t n = size() < rhs.size() ? size() : rhs.size();
for (size_t i = 0; i < n; ++i) m_elem[i] = rhs[i];
return *this;
}
Expand Down Expand Up @@ -303,14 +303,14 @@ struct Array<T, KOKKOS_INVALID_INDEX, Array<>::strided> {

KOKKOS_INLINE_FUNCTION
Array& operator=(const Array& rhs) {
const size_t n = std::min(m_size, rhs.size());
const size_t n = size() < rhs.size() ? size() : rhs.size();
for (size_t i = 0; i < n; ++i) m_elem[i * m_stride] = rhs[i];
return *this;
}

template <size_t N, class P>
KOKKOS_INLINE_FUNCTION Array& operator=(const Array<T, N, P>& rhs) {
const size_t n = std::min(m_size, rhs.size());
const size_t n = size() < rhs.size() ? size() : rhs.size();
for (size_t i = 0; i < n; ++i) m_elem[i * m_stride] = rhs[i];
return *this;
}
Expand Down
25 changes: 9 additions & 16 deletions core/src/Kokkos_BitManipulation.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -287,15 +287,13 @@ KOKKOS_IMPL_DEVICE_FUNCTION
std::enable_if_t<is_standard_unsigned_integer_type_v<T>, int>
countl_zero_builtin_device(T x) noexcept {
#if defined(KOKKOS_ENABLE_CUDA) || defined(KOKKOS_ENABLE_HIP)
if constexpr (sizeof(T) == sizeof(long long int)) {
if constexpr (sizeof(T) == sizeof(long long int))
return __clzll(reinterpret_cast<long long int&>(x));
} else if constexpr (sizeof(T) == sizeof(int)) {
if constexpr (sizeof(T) == sizeof(int))
return __clz(reinterpret_cast<int&>(x));
} else {
using ::Kokkos::Experimental::digits_v;
constexpr int shift = digits_v<unsigned int> - digits_v<T>;
return __clz(x) - shift;
}
using ::Kokkos::Experimental::digits_v;
constexpr int shift = digits_v<unsigned int> - digits_v<T>;
return __clz(x) - shift;
#elif defined(KOKKOS_ENABLE_SYCL)
return sycl::clz(x);
#else
Expand Down Expand Up @@ -332,11 +330,9 @@ KOKKOS_IMPL_DEVICE_FUNCTION
using ::Kokkos::Experimental::digits_v;
if (x == 0) return digits_v<T>;
#if defined(KOKKOS_ENABLE_CUDA) || defined(KOKKOS_ENABLE_HIP)
if constexpr (sizeof(T) == sizeof(long long int)) {
if constexpr (sizeof(T) == sizeof(long long int))
return __ffsll(reinterpret_cast<long long int&>(x)) - 1;
} else {
return __ffs(reinterpret_cast<int&>(x)) - 1;
}
return __ffs(reinterpret_cast<int&>(x)) - 1;
#elif defined(KOKKOS_ENABLE_SYCL)
return sycl::ctz(x);
#else
Expand Down Expand Up @@ -368,11 +364,8 @@ KOKKOS_IMPL_DEVICE_FUNCTION
std::enable_if_t<is_standard_unsigned_integer_type_v<T>, int>
popcount_builtin_device(T x) noexcept {
#if defined(KOKKOS_ENABLE_CUDA) || defined(KOKKOS_ENABLE_HIP)
if constexpr (sizeof(T) == sizeof(long long int)) {
return __popcll(x);
} else {
return __popc(x);
}
if constexpr (sizeof(T) == sizeof(long long int)) return __popcll(x);
return __popc(x);
#elif defined(KOKKOS_ENABLE_SYCL)
return sycl::popcount(x);
#else
Expand Down
3 changes: 2 additions & 1 deletion core/unit_test/TestMathematicalFunctions.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -30,7 +30,8 @@
#define MATHEMATICAL_FUNCTIONS_HAVE_LONG_DOUBLE_OVERLOADS
#endif

#if defined KOKKOS_COMPILER_INTEL
#if defined KOKKOS_COMPILER_INTEL || \
(defined(KOKKOS_COMPILER_NVCC) && KOKKOS_COMPILER_NVCC >= 1130)
#define MATHEMATICAL_FUNCTIONS_TEST_UNREACHABLE __builtin_unreachable();
#else
#define MATHEMATICAL_FUNCTIONS_TEST_UNREACHABLE
Expand Down

0 comments on commit 890148e

Please sign in to comment.