Skip to content

[AMDGPU] The __builtin_round intrinsic does not preserve -0.0 #65629

@jhuber6

Description

@jhuber6

The standard states the following about the behavior of the round functoin:

If the implementation supports IEEE floating-point arithmetic (IEC 60559), for the std::round function:

  • If num is ±0, it is returned, unmodified

This is currently not respected by the AMDGPU implementation of __builtin_round. See the following reproducer for an example.

#include <assert.h> #include <math.h> #include <stdio.h> int main() { double r; double in = -0.0; #pragma omp target map(from : r) { r = __builtin_round(in); } assert(signbit(r) && "Result is no longer negative"); }
$ clang round.c -fopenmp --offload-arch=gfx1030 -fopenmp-offload-mandatory $ ./a.out a.out: round.c:10: int main(): Assertion `signbit(r) && "Result is no longer negative"' failed. [1] 316479 IOT instruction (core dumped) ./a.out  $ clang round.c -fopenmp --offload-arch=sm_89 -fopenmp-offload-mandatory $ ./a.out

Metadata

Metadata

Assignees

Type

No type

Projects

No projects

Milestone

No milestone

Relationships

None yet

Development

No branches or pull requests

Issue actions