[clang] 31dfd2b - HIP: Directly call isnan builtin
Matt Arsenault via cfe-commits
cfe-commits at lists.llvm.org
Thu Jun 8 13:08:30 PDT 2023
Author: Matt Arsenault
Date: 2023-06-08T16:08:24-04:00
New Revision: 31dfd2b99ce31ae076aa35867dbede052693a6c9
URL: https://github.com/llvm/llvm-project/commit/31dfd2b99ce31ae076aa35867dbede052693a6c9
DIFF: https://github.com/llvm/llvm-project/commit/31dfd2b99ce31ae076aa35867dbede052693a6c9.diff
LOG: HIP: Directly call isnan builtin
Added:
Modified:
clang/lib/Headers/__clang_hip_math.h
clang/test/Headers/__clang_hip_math.hip
clang/test/Headers/hip-header.hip
clang/test/Headers/openmp_device_math_isnan.cpp
Removed:
################################################################################
diff --git a/clang/lib/Headers/__clang_hip_math.h b/clang/lib/Headers/__clang_hip_math.h
index b92d7b74f25cf..40ee1598200b5 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 d099691fcfd8f..453acda646a95 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 c11730f8717f5..8264b4e2c8e5d 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 ddb3d75ff1157..21419a184a325 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 <cmath>
@@ -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;
}
More information about the cfe-commits
mailing list