Comparison of negative zero and positive zero in CUDA's BFloat16 datatype

Accelerated Computing CUDA CUDA Programming and Performance
1

Greetings!

I am implementing vector bit quantization and reached the point where I have to handle negative zero. I am seeking guidance on what should be the result of an equality comparision of negative zero and positive zero if implementing own BFloat16. From CUDA documentation, it is not clear what does __heq() return if the operands are CUDART_NEG_ZERO_BF16 and CUDARD_ZERO_BF16. Please let me know!

2 
# cat t355.cu
#include <cuda_fp16.h>
#include <cuda_bf16.h>
#include <cstdio>
__global__ void k(){ printf("%d\n", (int)__heq(CUDART_NEG_ZERO_BF16, CUDART_ZERO_BF16));}

int main(){
  k<<<1,1>>>();
  cudaDeviceSynchronize();
  }
# nvcc -arch=sm_89 -o t355 t355.cu
# compute-sanitizer ./t355
========= COMPUTE-SANITIZER
1
========= ERROR SUMMARY: 0 errors
#

(CUDA 12.2)

Iā€™m not suggesting this is IEEE-754 arithemetic, but it seems at least consistent with a similar operation in IEEE-754.

1 Like