Skip to content

Deprecate and Replace cub::BFE #4031

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.

Sign up for GitHub

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

Merged
fbusato merged 11 commits into from
Mar 13, 2025
Merged

Deprecate and Replace cub::BFE #4031

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
11 commits
Select commit Hold shift + click to select a range
4695f1a
deprecate and replace BFE
fbusato Mar 6, 2025
815cb9e
prevent begin_bit == end_bit
fbusato Mar 6, 2025
148ab8a
use Catch2 SKIP()
fbusato Mar 6, 2025
1b5c9d5
fix msvc warning
fbusato Mar 6, 2025
2f7f86d
fix block radix sort test
fbusato Mar 7, 2025
3321906
Merge branch 'main' into deprecate-cub-bfe-bfi
fbusato Mar 11, 2025
cfa5d93
revert zero-size bit test cases
fbusato Mar 11, 2025
de7d898
add maybe_unused
fbusato Mar 11, 2025
fba5c0d
adapt extended bit to bit_start == width
fbusato Mar 12, 2025
a103ce8
allow start bit == width
fbusato Mar 12, 2025
80609d1
one more [[maybe_unused]]
fbusato Mar 13, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
3 changes: 2 additions & 1 deletion cub/cub/block/radix_rank_sort_operations.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -49,6 +49,7 @@

#include <thrust/type_traits/integer_sequence.h>

#include <cuda/bit>
#include <cuda/std/__algorithm/max.h>
#include <cuda/std/__algorithm/min.h>
#include <cuda/std/cstdint>
Expand Down Expand Up @@ -119,7 +120,7 @@ struct BFEDigitExtractor : BaseDigitExtractor<KeyT>

_CCCL_DEVICE _CCCL_FORCEINLINE ::cuda::std::uint32_t Digit(UnsignedBits key) const
{
return BFE(this->ProcessFloatMinusZero(key), bit_start, num_bits);
return ::cuda::bitfield_extract(this->ProcessFloatMinusZero(key), bit_start, num_bits);
}
};

Expand Down
19 changes: 12 additions & 7 deletions cub/cub/util_ptx.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -58,8 +58,9 @@ CUB_NAMESPACE_BEGIN
* Bitfield-extract.
*/
template <typename UnsignedBits, int BYTE_LEN>
_CCCL_DEVICE _CCCL_FORCEINLINE unsigned int
BFE(UnsignedBits source, unsigned int bit_start, unsigned int num_bits, detail::constant_t<BYTE_LEN> /*byte_len*/)
//! deprecated [Since 3.0]
CCCL_DEPRECATED_BECAUSE("Use cuda::bitfield_extract()") _CCCL_DEVICE _CCCL_FORCEINLINE unsigned int BFE(
UnsignedBits source, unsigned int bit_start, unsigned int num_bits, detail::constant_t<BYTE_LEN> /*byte_len*/)
{
unsigned int bits;
asm("bfe.u32 %0, %1, %2, %3;" : "=r"(bits) : "r"((unsigned int) source), "r"(bit_start), "r"(num_bits));
Expand All @@ -70,8 +71,9 @@ BFE(UnsignedBits source, unsigned int bit_start, unsigned int num_bits, detail::
* Bitfield-extract for 64-bit types.
*/
template <typename UnsignedBits>
_CCCL_DEVICE _CCCL_FORCEINLINE unsigned int
BFE(UnsignedBits source, unsigned int bit_start, unsigned int num_bits, detail::constant_t<8> /*byte_len*/)
//! deprecated [Since 3.0]
CCCL_DEPRECATED_BECAUSE("Use cuda::bitfield_extract()") _CCCL_DEVICE _CCCL_FORCEINLINE unsigned int BFE(
UnsignedBits source, unsigned int bit_start, unsigned int num_bits, detail::constant_t<8> /*byte_len*/)
{
const unsigned long long MASK = (1ull << num_bits) - 1;
return (source >> bit_start) & MASK;
Expand All @@ -82,8 +84,9 @@ BFE(UnsignedBits source, unsigned int bit_start, unsigned int num_bits, detail::
* Bitfield-extract for 128-bit types.
*/
template <typename UnsignedBits>
_CCCL_DEVICE _CCCL_FORCEINLINE unsigned int
BFE(UnsignedBits source, unsigned int bit_start, unsigned int num_bits, detail::constant_t<16> /*byte_len*/)
//! deprecated [Since 3.0]
CCCL_DEPRECATED_BECAUSE("Use cuda::bitfield_extract()") _CCCL_DEVICE _CCCL_FORCEINLINE unsigned int BFE(
UnsignedBits source, unsigned int bit_start, unsigned int num_bits, detail::constant_t<16> /*byte_len*/)
{
const __uint128_t MASK = (__uint128_t{1} << num_bits) - 1;
return (source >> bit_start) & MASK;
Expand All @@ -97,7 +100,9 @@ BFE(UnsignedBits source, unsigned int bit_start, unsigned int num_bits, detail::
* source may be an 8b, 16b, 32b, or 64b unsigned integer type.
*/
template <typename UnsignedBits>
_CCCL_DEVICE _CCCL_FORCEINLINE unsigned int BFE(UnsignedBits source, unsigned int bit_start, unsigned int num_bits)
//! deprecated [Since 3.0]
CCCL_DEPRECATED_BECAUSE("Use cuda::bitfield_extract()") _CCCL_DEVICE
_CCCL_FORCEINLINE unsigned int BFE(UnsignedBits source, unsigned int bit_start, unsigned int num_bits)
{
return BFE(source, bit_start, num_bits, detail::constant_v<int{sizeof(UnsignedBits)}>);
}
Expand Down
Loading