Skip to content

Commit

Permalink
HIP: Directly call isnan builtin
Browse files Browse the repository at this point in the history
  • Loading branch information
@arsenm
arsenm committed Jun 8, 2023
1 parent 7a52f79 commit 31dfd2b
Show file tree
Hide file tree
Showing 4 changed files with 24 additions and 28 deletions.
4 changes: 2 additions & 2 deletions clang/lib/Headers/__clang_hip_math.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -274,7 +274,7 @@ __DEVICE__
__RETURN_TYPE __isinff(float __x) { return __builtin_isinf(__x); }

__DEVICE__
__RETURN_TYPE __isnanf(float __x) { return __ocml_isnan_f32(__x); }
__RETURN_TYPE __isnanf(float __x) { return __builtin_isnan(__x); }

__DEVICE__
float j0f(float __x) { return __ocml_j0_f32(__x); }
Expand Down Expand Up @@ -823,7 +823,7 @@ __DEVICE__
__RETURN_TYPE __isinf(double __x) { return __builtin_isinf(__x); }

__DEVICE__
__RETURN_TYPE __isnan(double __x) { return __ocml_isnan_f64(__x); }
__RETURN_TYPE __isnan(double __x) { return __builtin_isnan(__x); }

__DEVICE__
double j0(double __x) { return __ocml_j0_f64(__x); }
Expand Down
20 changes: 6 additions & 14 deletions clang/test/Headers/__clang_hip_math.hip
View file
Original file line number Diff line number Diff line change
Expand Up @@ -1221,35 +1221,27 @@ extern "C" __device__ BOOL_TYPE test___isinf(double x) {

// DEFAULT-LABEL: @test___isnanf(
// DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isnan_f32(float noundef [[X:%.*]]) #[[ATTR14]]
// DEFAULT-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0
// DEFAULT-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32
// DEFAULT-NEXT: [[CMP_I:%.*]] = fcmp contract uno float [[X:%.*]], 0.000000e+00
// DEFAULT-NEXT: [[CONV:%.*]] = zext i1 [[CMP_I]] to i32
// DEFAULT-NEXT: ret i32 [[CONV]]
//
// FINITEONLY-LABEL: @test___isnanf(
// FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isnan_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]]
// FINITEONLY-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0
// FINITEONLY-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32
// FINITEONLY-NEXT: ret i32 [[CONV]]
// FINITEONLY-NEXT: ret i32 0
//
extern "C" __device__ BOOL_TYPE test___isnanf(float x) {
return __isnanf(x);
}

// DEFAULT-LABEL: @test___isnan(
// DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isnan_f64(double noundef [[X:%.*]]) #[[ATTR14]]
// DEFAULT-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0
// DEFAULT-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32
// DEFAULT-NEXT: [[CMP_I:%.*]] = fcmp contract uno double [[X:%.*]], 0.000000e+00
// DEFAULT-NEXT: [[CONV:%.*]] = zext i1 [[CMP_I]] to i32
// DEFAULT-NEXT: ret i32 [[CONV]]
//
// FINITEONLY-LABEL: @test___isnan(
// FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_isnan_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]]
// FINITEONLY-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0
// FINITEONLY-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32
// FINITEONLY-NEXT: ret i32 [[CONV]]
// FINITEONLY-NEXT: ret i32 0
//
extern "C" __device__ BOOL_TYPE test___isnan(double x) {
return __isnan(x);
Expand Down
8 changes: 4 additions & 4 deletions clang/test/Headers/hip-header.hip
View file
Original file line number Diff line number Diff line change
Expand Up @@ -134,12 +134,12 @@ __device__ double test_isnan() {
double d = 5.0;
float f = 5.0;

// AMD_INT_RETURN: call i32 @__ocml_isnan_f32(float
// AMD_BOOL_RETURN: call i32 @__ocml_isnan_f32(float
// AMD_INT_RETURN: fcmp contract uno float
// AMD_BOOL_RETURN: fcmp contract uno float
r += isnan(f);

// AMD_INT_RETURN: call i32 @__ocml_isnan_f64(double
// AMD_BOOL_RETURN: call i32 @__ocml_isnan_f64(double
// AMD_INT_RETURN: fcmp contract uno double
// AMD_BOOL_RETURN: fcmp contract uno double
r += isnan(d);

return r ;
Expand Down
20 changes: 12 additions & 8 deletions clang/test/Headers/openmp_device_math_isnan.cpp
View file
Original file line number Diff line number Diff line change
@@ -1,34 +1,38 @@
// RUN: %clang_cc1 -x c++ -internal-isystem %S/Inputs/include -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
// RUN: %clang_cc1 -x c++ -internal-isystem %S/Inputs/include -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-ppc-host.bc
// RUN: %clang_cc1 -x c++ -include __clang_openmp_device_functions.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -fopenmp -triple nvptx64-nvidia-cuda -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=BOOL_RETURN
// RUN: %clang_cc1 -x c++ -include __clang_openmp_device_functions.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -fopenmp -triple amdgcn-amd-amdhsa -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=AMD_BOOL_RETURN
// RUN: %clang_cc1 -x c++ -include __clang_openmp_device_functions.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -fopenmp -triple amdgcn-amd-amdhsa -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=AMD_BOOL_RETURN_SAFE
// RUN: %clang_cc1 -x c++ -internal-isystem %S/Inputs/include -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc -ffast-math -ffp-contract=fast
// RUN: %clang_cc1 -x c++ -internal-isystem %S/Inputs/include -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-ppc-host.bc -ffast-math -ffp-contract=fast
// RUN: %clang_cc1 -x c++ -include __clang_openmp_device_functions.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -fopenmp -triple nvptx64-nvidia-cuda -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -ffast-math -ffp-contract=fast | FileCheck %s --check-prefix=BOOL_RETURN
// RUN: %clang_cc1 -x c++ -include __clang_openmp_device_functions.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -fopenmp -triple amdgcn-amd-amdhsa -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -ffast-math -ffp-contract=fast | FileCheck %s --check-prefix=AMD_BOOL_RETURN
// RUN: %clang_cc1 -x c++ -include __clang_openmp_device_functions.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -fopenmp -triple amdgcn-amd-amdhsa -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -ffast-math -ffp-contract=fast | FileCheck %s --check-prefix=AMD_BOOL_RETURN_FAST
// RUN: %clang_cc1 -x c++ -internal-isystem %S/Inputs/include -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc -DUSE_ISNAN_WITH_INT_RETURN
// RUN: %clang_cc1 -x c++ -internal-isystem %S/Inputs/include -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-ppc-host.bc -DUSE_ISNAN_WITH_INT_RETURN
// RUN: %clang_cc1 -x c++ -include __clang_openmp_device_functions.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -fopenmp -triple nvptx64-nvidia-cuda -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -DUSE_ISNAN_WITH_INT_RETURN | FileCheck %s --check-prefix=INT_RETURN
// RUN: %clang_cc1 -x c++ -include __clang_openmp_device_functions.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -fopenmp -triple amdgcn-amd-amdhsa -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -DUSE_ISNAN_WITH_INT_RETURN | FileCheck %s --check-prefix=AMD_INT_RETURN
// RUN: %clang_cc1 -x c++ -include __clang_openmp_device_functions.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -fopenmp -triple amdgcn-amd-amdhsa -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -DUSE_ISNAN_WITH_INT_RETURN | FileCheck %s --check-prefix=AMD_INT_RETURN_SAFE
// RUN: %clang_cc1 -x c++ -internal-isystem %S/Inputs/include -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc -ffast-math -ffp-contract=fast -DUSE_ISNAN_WITH_INT_RETURN
// RUN: %clang_cc1 -x c++ -internal-isystem %S/Inputs/include -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-ppc-host.bc -ffast-math -ffp-contract=fast -DUSE_ISNAN_WITH_INT_RETURN
// RUN: %clang_cc1 -x c++ -include __clang_openmp_device_functions.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -fopenmp -triple nvptx64-nvidia-cuda -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -ffast-math -ffp-contract=fast -DUSE_ISNAN_WITH_INT_RETURN | FileCheck %s --check-prefix=INT_RETURN
// RUN: %clang_cc1 -x c++ -include __clang_openmp_device_functions.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -fopenmp -triple amdgcn-amd-amdhsa -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -ffast-math -ffp-contract=fast -DUSE_ISNAN_WITH_INT_RETURN | FileCheck %s --check-prefix=AMD_INT_RETURN
// RUN: %clang_cc1 -x c++ -include __clang_openmp_device_functions.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -fopenmp -triple amdgcn-amd-amdhsa -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -ffast-math -ffp-contract=fast -DUSE_ISNAN_WITH_INT_RETURN | FileCheck %s --check-prefix=AMD_INT_RETURN_FAST
// expected-no-diagnostics

#include <cmath>

double math(float f, double d) {
double r = 0;
// INT_RETURN: call i32 @__nv_isnanf(float
// AMD_INT_RETURN: call i32 @__ocml_isnan_f32(float
// AMD_INT_RETURN_SAFE: fcmp uno float
// AMD_INT_RETURN_FAST: sitofp i32 0 to double
// BOOL_RETURN: call i32 @__nv_isnanf(float
// AMD_BOOL_RETURN: call i32 @__ocml_isnan_f32(float
// AMD_BOOL_RETURN_SAFE: fcmp uno float
// AMD_BOOL_RETURN_FAST: icmp ne i32 0, 0
r += std::isnan(f);
// INT_RETURN: call i32 @__nv_isnand(double
// AMD_INT_RETURN: call i32 @__ocml_isnan_f64(double
// AMD_INT_RETURN_SAFE: fcmp uno double
// AMD_INT_RETURN_FAST: sitofp i32 0 to double
// BOOL_RETURN: call i32 @__nv_isnand(double
// AMD_BOOL_RETURN: call i32 @__ocml_isnan_f64(double
// AMD_BOOL_RETURN_SAFE: fcmp uno double
// AMD_BOOL_RETURN_FAST: icmp ne i32 0, 0
r += std::isnan(d);
return r;
}
Expand Down

0 comments on commit 31dfd2b

Please sign in to comment.