Skip to content

Commit

Permalink
Don't use host-only functions in host-device contexts.
Browse files Browse the repository at this point in the history
This fixes the issue reported in NVIDIA#299. There's no
clear reason why this should use `RandomBits` unconditionally.
  • Loading branch information
alliepiper committed Mar 24, 2022
1 parent 2ba1dbb commit 2ec04b9
Showing 1 changed file with 9 additions and 16 deletions.
25 changes: 9 additions & 16 deletions test/test_util.h
Original file line number Diff line number Diff line change
Expand Up @@ -713,22 +713,15 @@ __host__ __device__ __forceinline__ void InitValue(
// This specialization only appears to be used by test_warp_scan.
// It initializes with uniform values and random keys, so we need to
// protect the call to the host-only RandomBits.
if (CUB_IS_DEVICE_CODE)
{
#if CUB_INCLUDE_DEVICE_CODE
_CubLog("%s\n",
"cub::InitValue cannot generate random numbers on device.");
CUB_NS_QUALIFIER::ThreadTrap();
#endif // CUB_INCLUDE_DEVICE_CODE
}
else
{
#if CUB_INCLUDE_HOST_CODE
// Assign corresponding flag with a likelihood of the last bit being set with entropy-reduction level 3
RandomBits(value.key, 3);
value.key = (value.key & 0x1);
#endif // CUB_INCLUDE_HOST_CODE
}
NV_IF_TARGET(NV_IS_HOST,
// Assign corresponding flag with a likelihood of the last bit
// being set with entropy-reduction level 3
(RandomBits(value.key, 3); value.key = (value.key & 0x1);),
( // NV_IS_DEVICE
_CubLog("%s\n",
"cub::InitValue cannot generate random numbers on "
"device.");
CUB_NS_QUALIFIER::ThreadTrap();));
}


Expand Down

0 comments on commit 2ec04b9

Please sign in to comment.