From 31dfd2b99ce31ae076aa35867dbede052693a6c9 Mon Sep 17 00:00:00 2001 From: Matt Arsenault Date: Sun, 20 Nov 2022 08:48:29 -0800 Subject: [PATCH] HIP: Directly call isnan builtin --- clang/lib/Headers/__clang_hip_math.h | 4 ++-- clang/test/Headers/__clang_hip_math.hip | 20 ++++++-------------- clang/test/Headers/hip-header.hip | 8 ++++---- clang/test/Headers/openmp_device_math_isnan.cpp | 20 ++++++++++++-------- 4 files changed, 24 insertions(+), 28 deletions(-) diff --git a/clang/lib/Headers/__clang_hip_math.h b/clang/lib/Headers/__clang_hip_math.h index b92d7b7..40ee159 100644 --- a/clang/lib/Headers/__clang_hip_math.h +++ b/clang/lib/Headers/__clang_hip_math.h @@ -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); } @@ -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); } diff --git a/clang/test/Headers/__clang_hip_math.hip b/clang/test/Headers/__clang_hip_math.hip index d099691..453acda 100644 --- a/clang/test/Headers/__clang_hip_math.hip +++ b/clang/test/Headers/__clang_hip_math.hip @@ -1221,17 +1221,13 @@ 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); @@ -1239,17 +1235,13 @@ extern "C" __device__ BOOL_TYPE test___isnanf(float 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); diff --git a/clang/test/Headers/hip-header.hip b/clang/test/Headers/hip-header.hip index c11730f..8264b4e 100644 --- a/clang/test/Headers/hip-header.hip +++ b/clang/test/Headers/hip-header.hip @@ -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 ; diff --git a/clang/test/Headers/openmp_device_math_isnan.cpp b/clang/test/Headers/openmp_device_math_isnan.cpp index ddb3d75..21419a1 100644 --- a/clang/test/Headers/openmp_device_math_isnan.cpp +++ b/clang/test/Headers/openmp_device_math_isnan.cpp @@ -1,19 +1,19 @@ // 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 @@ -21,14 +21,18 @@ 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; } -- 2.7.4