Search code examples
cudagpgpugpuparity

Compute bit parity in CUDA

CUDA has popcount intrinsics for 32-bit and 64-bit types: __popc() and __popcll().

Does CUDA also have intrinsics to get the parity of 32-bit and 64-bit types? (The parity refers to whether an integer has an even or odd amount of 1-bits.)

For example, GCC has __builtin_parityl() for 64-bit integers.

And here's a C function that does the same thing:

inline uint parity64(uint64 n){
  n ^= n >> 1;
  n ^= n >> 2;
  n  = (n & 0x1111111111111111lu) * 0x1111111111111111lu;
  return (n >> 60) & 1;
}
Solution

  • I'm not aware of a parity intrinsic for CUDA.

    However you should be able to create a fairly simple function to do it using either the __popc() (32-bit unsigned case) or __popcll() (64-bit unsigned case) intrinsics.

    For example, the following function should indicate whether the number of 1 bits in a 64-bit unsigned quantity is odd (true) or even (false):

    __device__ bool my_parity(unsigned long long d){
  return (__popcll(d) & 1);}