Backport bit manipulation facility from C++20 standard library header `<bit>` #4577

dalg24 · 2021-12-02T16:55:58Z

~~Fix #3828~~

See documentation in kokkos/kokkos-core-wiki#300

~~Deprecate Kokkos::log2(unsigned) -> int~~ done in Deprecate log2(unsigned) -> int #4595
~~Provide alternative Impl::bit_log2(UnsignedInteger) -> UnsignedInteger for internal use only~~ withdrawn (might come back with it in a follow up PR)
Backport bit manipulation facility from C++20 https://eel.is/c++draft/bit
~~Remove Impl::bit_count(unsigned int) in favor of Experimental::popcount(UnsignedInteger) -> UnsignedInteger~~ withdrawn (will propose in a follow up PR)

NOTES

I have already implemented [bit.rot] and I intend to remove Impl::rotate_right(unsigned, int) -> unsigned in the Bitset implementation but I was planing to do that in a follow-up PR
I intend to incrementally replace everything in <impl/Kokkos_BitOps.hpp> until we can get rid of the header
I will also attempt to replace the integral_power_of_two* facility in impl/Kokkos_Traits.hpp but I would rather deal with it later
I only used GCC built-in functions on the host-side for Experimental::*_builtin function templates. We can specialize further and add Intel intrinsics and such later.

ajpowelsnl

The PR generally looks good, and there's no reason to think it shouldn't be merged. But I do have a basic question (perhaps I should know the answer, but I don't), why does Kokkos need to manipulate at the bit level? To deal with various proclivities and folkways of different compilers and backends?

PhilMiller · 2023-02-21T17:58:34Z

This is part of a broader theme of Kokkos broadly mimicking C++ standard library functionality with __host__ __device__ annotations (or their moral equivalent) so that applications can use them on GPUs. See also Kokkos::pair, Kokkos::complex, Kokkos::array, math routines (pow, sqrt, etc), and much more.

core/src/Kokkos_BitManipulation.hpp

masterleinad · 2023-02-21T18:59:48Z

core/src/Kokkos_BitManipulation.hpp

+
+namespace Kokkos::Impl {
+
+#if defined(KOKKOS_COMPILER_CLANG) || defined(KOKKOS_COMPILER_GCC)


IntelLLVM would probably also define it.

I'd prefer fiddling with more compiler and other intrinsics later.

masterleinad · 2023-02-21T19:01:10Z

core/src/Kokkos_BitManipulation.hpp

+countl_zero(T x) noexcept {
+  using ::Kokkos::Experimental::digits_v;
+  if (x == 0) return digits_v<T>;
+  // TODO use compiler intrinsics when available


Suggested change

// TODO use compiler intrinsics when available

// TODO use constexpr compiler intrinsics when available

core/unit_test/TestBitManipulationBuiltins.hpp

masterleinad · 2023-02-22T03:45:20Z

The PR generally looks good, and there's no reason to think it shouldn't be merged. But I do have a basic question (perhaps I should know the answer, but I don't), why does Kokkos need to manipulate at the bit level? To deal with various proclivities and folkways of different compilers and backends?

We already have a much of places where we would use these functions in the library.

masterleinad

Looks OK to me.

dalg24 · 2023-02-22T04:06:21Z

(I figured out the failure with NVCC. It was a few stray constexpr qualifiers on builtin device side function templates.)

dalg24 · 2023-02-22T05:05:58Z

Retest this please

dalg24 · 2023-02-22T12:20:35Z

Retest this please

dalg24 · 2023-02-22T14:33:15Z

core/unit_test/TestBitManipulationBuiltins.hpp

+}
+
+TEST(TEST_CATEGORY, bit_manip_countr_zero) {
+  test_bit_manip_countr_zero<unsigned char>();


There is an issue with SYCL

4: value at 1 which is 248 was expected to be 0 4: value at 2 which is 249 was expected to be 1 4: value at 3 which is 248 was expected to be 0 4: value at 4 which is 250 was expected to be 2 4: value at 5 which is 248 was expected to be 0 4: value at 6 which is 249 was expected to be 1 4: value at 7 which is 248 was expected to be 0 4: value at 8 which is 251 was expected to be 3 4: value at 9 which is 248 was expected to be 0 4: value at 7e which is 249 was expected to be 1 4: value at 7f which is 248 was expected to be 0 4: value at 80 which is 255 was expected to be 7 4: value at 81 which is 248 was expected to be 0 4: value at 82 which is 249 was expected to be 1 4: value at 7f which is 248 was expected to be 0 4: /var/jenkins/workspace/Kokkos/core/unit_test/TestBitManipulationBuiltins.hpp:67: Failure 4: Expected equality of these values: 4: errors 4: Which is: 15 4: 0 4: Failed check no error for countr_zero(unsigned char) 4: [ FAILED ] sycl.bit_manip_countr_zero (9 ms)

Daniel reported it works on Intel GPUs. I disabled on NVIDIA.

dalg24 · 2023-02-22T14:33:28Z

core/unit_test/TestBitManipulationBuiltins.hpp

+}
+
+TEST(TEST_CATEGORY, bit_manip_countr_one) {
+  test_bit_manip_countr_one<unsigned char>();


4: [ RUN ] sycl.bit_manip_countr_one 4: value at 0 which is 248 was expected to be 0 4: value at 1 which is 249 was expected to be 1 4: value at 2 which is 248 was expected to be 0 4: value at 3 which is 250 was expected to be 2 4: value at 4 which is 248 was expected to be 0 4: value at 5 which is 249 was expected to be 1 4: value at 6 which is 248 was expected to be 0 4: value at 7 which is 251 was expected to be 3 4: value at 8 which is 248 was expected to be 0 4: value at 9 which is 249 was expected to be 1 4: value at 7e which is 248 was expected to be 0 4: value at 7f which is 255 was expected to be 7 4: value at 80 which is 248 was expected to be 0 4: value at 7e which is 248 was expected to be 0 4: value at 7f which is 255 was expected to be 7 4: /var/jenkins/workspace/Kokkos/core/unit_test/TestBitManipulationBuiltins.hpp:67: Failure 4: Expected equality of these values: 4: errors 4: Which is: 15 4: 0 4: Failed check no error for countr_one(unsigned char)

dalg24 · 2023-02-22T16:41:09Z

SYCL build passed with the last commit. HIP failure is unrelated.

…plate functions

dalg24 · 2023-02-22T16:46:06Z

core/unit_test/TestBitManipulationBuiltins.hpp

+  static_assert(std::is_same_v<decltype(countl_zero_builtin(UInt())), int>);
+  constexpr auto max =
+      Kokkos::Experimental::finite_max_v<std::make_signed_t<UInt>>;
+  //                                     ^^^^^^^^^^^^^^^^^^


Didn't mean to keep that change...

nmm0

Looks good to me, we just need to file an issue for the eventual merging of these paths

nmm0 · 2023-02-21T18:13:41Z

core/src/Kokkos_BitManipulation.hpp

+countl_zero(T x) noexcept {
+  using ::Kokkos::Experimental::digits_v;
+  if (x == 0) return digits_v<T>;
+  // TODO use compiler intrinsics when available


We should file a new issue to avoid TODOs otherwise these will be forgotten, especially since we can't do this properly until C++20 support

nmm0 · 2023-02-21T18:15:17Z

core/src/Kokkos_BitManipulation.hpp

+countr_zero(T x) noexcept {
+  using ::Kokkos::Experimental::digits_v;
+  if (x == 0) return digits_v<T>;
+  // TODO use compiler intrinsics when available


Same as above

nmm0 · 2023-02-21T18:15:33Z

core/src/Kokkos_BitManipulation.hpp

+    Impl::is_standard_unsigned_integer_type_v<T>, int>
+popcount(T x) noexcept {
+  if (x == 0) return 0;
+  // TODO use compiler intrinsics when available


Same as above

nmm0 · 2023-02-22T17:59:04Z

I promised to myself that when I would review I would ask for docs/changelog entry. You already have the docs, so I guess we should add to the changelog of whatever this is targeting.

dalg24 · 2023-02-22T18:23:41Z

Ignoring HIP failure

dalg24 added the Refactor Tidying up: make code code, cleaner, simpler, understandable and robust label Dec 2, 2021

dalg24 mentioned this pull request Dec 3, 2021

Remove (unused) Impl::[is_]power_of_two traits #4578

Merged

dalg24 mentioned this pull request May 31, 2022

Allow non-power-of-two team sizes for team reductions and scans #4809

Merged

dalg24 mentioned this pull request Jun 9, 2022

Moving SIMD types into Kokkos #5016

Merged

dalg24 mentioned this pull request Aug 11, 2022

Team- and thread-level sort, sort_by_key #5317

Merged

dalg24 force-pushed the bit_manip branch from 487d60f to 6a4d560 Compare January 31, 2023 03:52

dalg24 mentioned this pull request Feb 21, 2023

Add documentation for <Kokkos_BitManipulation.hpp> header kokkos/kokkos-core-wiki#300

Merged

dalg24 force-pushed the bit_manip branch from 6c22658 to 8640d07 Compare February 21, 2023 03:25

dalg24 marked this pull request as ready for review February 21, 2023 03:27

dalg24 force-pushed the bit_manip branch 3 times, most recently from 7cf153a to 747071f Compare February 21, 2023 12:53

Backport function templates from <bit> standard library header

b3a8182

dalg24 force-pushed the bit_manip branch 3 times, most recently from 4d4be48 to e914ebf Compare February 21, 2023 16:11

ajpowelsnl approved these changes Feb 21, 2023

View reviewed changes

masterleinad reviewed Feb 21, 2023

View reviewed changes

nliber approved these changes Feb 22, 2023

View reviewed changes

masterleinad approved these changes Feb 22, 2023

View reviewed changes

dalg24 force-pushed the bit_manip branch from e914ebf to 4ed2113 Compare February 22, 2023 04:03

dalg24 force-pushed the bit_manip branch 2 times, most recently from 320378a to 9823ec8 Compare February 22, 2023 13:01

dalg24 commented Feb 22, 2023

View reviewed changes

dalg24 added 2 commits February 22, 2023 11:45

Add Experimental::*_builtin counterpart to the bit manipulation tem…

b166d77

…plate functions

Disable uchar test to work around broken sycl::ctz on NVIDIA GPUs

75a3e80

dalg24 commented Feb 22, 2023

View reviewed changes

dalg24 force-pushed the bit_manip branch from a08421c to 75a3e80 Compare February 22, 2023 16:46

masterleinad mentioned this pull request Feb 22, 2023

Wrong results with sycl::ctz for unsigned char with SYCL+Cuda intel/llvm#8426

Closed

nmm0 approved these changes Feb 22, 2023

View reviewed changes

dalg24 mentioned this pull request Feb 22, 2023

CHANGELOG: 4.1.0 #5902

Closed

dalg24 merged commit fb3d754 into kokkos:develop Feb 22, 2023

dalg24 deleted the bit_manip branch February 22, 2023 18:36

dalg24 mentioned this pull request Feb 23, 2023

Implement rot{l,r} function templates #5907

Merged

dalg24 mentioned this pull request Mar 8, 2023

Bit manipulation: implement byteswap #5967

Merged

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Backport bit manipulation facility from C++20 standard library header `<bit>` #4577

Backport bit manipulation facility from C++20 standard library header `<bit>` #4577

dalg24 commented Dec 2, 2021 •

edited

ajpowelsnl left a comment

PhilMiller commented Feb 21, 2023

masterleinad Feb 21, 2023

dalg24 Feb 22, 2023

masterleinad Feb 21, 2023

masterleinad commented Feb 22, 2023

masterleinad left a comment

dalg24 commented Feb 22, 2023

dalg24 commented Feb 22, 2023

dalg24 commented Feb 22, 2023

dalg24 Feb 22, 2023

dalg24 Feb 22, 2023

dalg24 Feb 22, 2023

dalg24 commented Feb 22, 2023

dalg24 Feb 22, 2023

nmm0 left a comment

nmm0 Feb 21, 2023

nmm0 Feb 21, 2023

nmm0 Feb 21, 2023

nmm0 commented Feb 22, 2023

dalg24 commented Feb 22, 2023


		namespace Kokkos::Impl {

		#if defined(KOKKOS_COMPILER_CLANG) \|\| defined(KOKKOS_COMPILER_GCC)

	// TODO use compiler intrinsics when available
	// TODO use constexpr compiler intrinsics when available

Backport bit manipulation facility from C++20 standard library header <bit> #4577

Backport bit manipulation facility from C++20 standard library header <bit> #4577

Conversation

dalg24 commented Dec 2, 2021 • edited

NOTES

ajpowelsnl left a comment

Choose a reason for hiding this comment

PhilMiller commented Feb 21, 2023

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

masterleinad commented Feb 22, 2023

masterleinad left a comment

Choose a reason for hiding this comment

dalg24 commented Feb 22, 2023

dalg24 commented Feb 22, 2023

dalg24 commented Feb 22, 2023

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

dalg24 commented Feb 22, 2023

Choose a reason for hiding this comment

nmm0 left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

nmm0 commented Feb 22, 2023

dalg24 commented Feb 22, 2023

Backport bit manipulation facility from C++20 standard library header `<bit>` #4577

Backport bit manipulation facility from C++20 standard library header `<bit>` #4577

dalg24 commented Dec 2, 2021 •

edited