Compute bit parity in CUDA

Question

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;
}

can't you just take the popcount and determine if that number is even or odd, i.e. `__popcll(n) & 1` ? — Robert Crovella, Oct 06 '17 at 22:31
The [CUDA documentation](http://docs.nvidia.com/cuda/cuda-math-api/group__CUDA__MATH__INTRINSIC__INT.html) doesn't list *bit parity* as an intrinsic, so I may have to go with your suggestion — étale-cohomology, Oct 06 '17 at 23:14

score 1 · Accepted Answer · answered Oct 07 '17 at 17:28

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);}

Compute bit parity in CUDA

1 Answers1