[clang] 7d6c2e1 - [clang] Use llvm.is_fpclass to implement FP classification functions
Serge Pavlov via cfe-commits
cfe-commits at lists.llvm.org
Tue Jul 11 07:35:41 PDT 2023
Author: Serge Pavlov
Date: 2023-07-11T21:34:53+07:00
New Revision: 7d6c2e18114de9900d1b012cf9c219803b183f63
URL: https://github.com/llvm/llvm-project/commit/7d6c2e18114de9900d1b012cf9c219803b183f63
DIFF: https://github.com/llvm/llvm-project/commit/7d6c2e18114de9900d1b012cf9c219803b183f63.diff
LOG: [clang] Use llvm.is_fpclass to implement FP classification functions
Builtin floating-point number classification functions:
- __builtin_isnan,
- __builtin_isinf,
- __builtin_finite, and
- __builtin_isnormal
now are implemented using `llvm.is_fpclass`.
This change makes the target callback `TargetCodeGenInfo::testFPKind`
unneeded. It is preserved in this change and should be removed later.
Differential Revision: https://reviews.llvm.org/D112932
Added:
Modified:
clang/lib/CodeGen/CGBuiltin.cpp
clang/test/CodeGen/X86/strictfp_builtins.c
clang/test/CodeGen/aarch64-strictfp-builtins.c
clang/test/CodeGen/builtin_float.c
clang/test/CodeGen/builtin_float_strictfp.c
clang/test/CodeGen/builtins.c
clang/test/CodeGen/isfpclass.c
clang/test/CodeGen/strictfp_builtins.c
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/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp
index 088fb46faee5d8..1fbf29abcaae54 100644
--- a/clang/lib/CodeGen/CGBuiltin.cpp
+++ b/clang/lib/CodeGen/CGBuiltin.cpp
@@ -31,6 +31,7 @@
#include "clang/Frontend/FrontendDiagnostic.h"
#include "llvm/ADT/APFloat.h"
#include "llvm/ADT/APInt.h"
+#include "llvm/ADT/FloatingPointMode.h"
#include "llvm/ADT/SmallPtrSet.h"
#include "llvm/ADT/StringExtras.h"
#include "llvm/Analysis/ValueTracking.h"
@@ -2239,6 +2240,17 @@ static unsigned mutateLongDoubleBuiltin(unsigned BuiltinID) {
}
}
+static Value *tryUseTestFPKind(CodeGenFunction &CGF, unsigned BuiltinID,
+ Value *V) {
+ if (CGF.Builder.getIsFPConstrained() &&
+ CGF.Builder.getDefaultConstrainedExcept() != fp::ebIgnore) {
+ if (Value *Result =
+ CGF.getTargetHooks().testFPKind(V, BuiltinID, CGF.Builder, CGF.CGM))
+ return Result;
+ }
+ return nullptr;
+}
+
RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID,
const CallExpr *E,
ReturnValueSlot ReturnValue) {
@@ -3122,37 +3134,49 @@ RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID,
// ZExt bool to int type.
return RValue::get(Builder.CreateZExt(LHS, ConvertType(E->getType())));
}
+
case Builtin::BI__builtin_isnan: {
CodeGenFunction::CGFPOptionsRAII FPOptsRAII(*this, E);
Value *V = EmitScalarExpr(E->getArg(0));
- llvm::Type *Ty = V->getType();
- const llvm::fltSemantics &Semantics = Ty->getFltSemantics();
- if (!Builder.getIsFPConstrained() ||
- Builder.getDefaultConstrainedExcept() == fp::ebIgnore ||
- !Ty->isIEEE()) {
- V = Builder.CreateFCmpUNO(V, V, "cmp");
- return RValue::get(Builder.CreateZExt(V, ConvertType(E->getType())));
- }
+ if (Value *Result = tryUseTestFPKind(*this, BuiltinID, V))
+ return RValue::get(Result);
+ return RValue::get(
+ Builder.CreateZExt(Builder.createIsFPClass(V, FPClassTest::fcNan),
+ ConvertType(E->getType())));
+ }
- if (Value *Result = getTargetHooks().testFPKind(V, BuiltinID, Builder, CGM))
+ case Builtin::BI__builtin_isinf: {
+ CodeGenFunction::CGFPOptionsRAII FPOptsRAII(*this, E);
+ Value *V = EmitScalarExpr(E->getArg(0));
+ if (Value *Result = tryUseTestFPKind(*this, BuiltinID, V))
return RValue::get(Result);
+ return RValue::get(
+ Builder.CreateZExt(Builder.createIsFPClass(V, FPClassTest::fcInf),
+ ConvertType(E->getType())));
+ }
- // NaN has all exp bits set and a non zero significand. Therefore:
- // isnan(V) == ((exp mask - (abs(V) & exp mask)) < 0)
- unsigned bitsize = Ty->getScalarSizeInBits();
- llvm::IntegerType *IntTy = Builder.getIntNTy(bitsize);
- Value *IntV = Builder.CreateBitCast(V, IntTy);
- APInt AndMask = APInt::getSignedMaxValue(bitsize);
- Value *AbsV =
- Builder.CreateAnd(IntV, llvm::ConstantInt::get(IntTy, AndMask));
- APInt ExpMask = APFloat::getInf(Semantics).bitcastToAPInt();
- Value *Sub =
- Builder.CreateSub(llvm::ConstantInt::get(IntTy, ExpMask), AbsV);
- // V = sign bit (Sub) <=> V = (Sub < 0)
- V = Builder.CreateLShr(Sub, llvm::ConstantInt::get(IntTy, bitsize - 1));
- if (bitsize > 32)
- V = Builder.CreateTrunc(V, ConvertType(E->getType()));
- return RValue::get(V);
+ case Builtin::BIfinite:
+ case Builtin::BI__finite:
+ case Builtin::BIfinitef:
+ case Builtin::BI__finitef:
+ case Builtin::BIfinitel:
+ case Builtin::BI__finitel:
+ case Builtin::BI__builtin_isfinite: {
+ CodeGenFunction::CGFPOptionsRAII FPOptsRAII(*this, E);
+ Value *V = EmitScalarExpr(E->getArg(0));
+ if (Value *Result = tryUseTestFPKind(*this, BuiltinID, V))
+ return RValue::get(Result);
+ return RValue::get(
+ Builder.CreateZExt(Builder.createIsFPClass(V, FPClassTest::fcFinite),
+ ConvertType(E->getType())));
+ }
+
+ case Builtin::BI__builtin_isnormal: {
+ CodeGenFunction::CGFPOptionsRAII FPOptsRAII(*this, E);
+ Value *V = EmitScalarExpr(E->getArg(0));
+ return RValue::get(
+ Builder.CreateZExt(Builder.createIsFPClass(V, FPClassTest::fcNormal),
+ ConvertType(E->getType())));
}
case Builtin::BI__builtin_isfpclass: {
@@ -3388,52 +3412,6 @@ RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID,
return RValue::get(Result);
}
- case Builtin::BIfinite:
- case Builtin::BI__finite:
- case Builtin::BIfinitef:
- case Builtin::BI__finitef:
- case Builtin::BIfinitel:
- case Builtin::BI__finitel:
- case Builtin::BI__builtin_isinf:
- case Builtin::BI__builtin_isfinite: {
- // isinf(x) --> fabs(x) == infinity
- // isfinite(x) --> fabs(x) != infinity
- // x != NaN via the ordered compare in either case.
- CodeGenFunction::CGFPOptionsRAII FPOptsRAII(*this, E);
- Value *V = EmitScalarExpr(E->getArg(0));
- llvm::Type *Ty = V->getType();
- if (!Builder.getIsFPConstrained() ||
- Builder.getDefaultConstrainedExcept() == fp::ebIgnore ||
- !Ty->isIEEE()) {
- Value *Fabs = EmitFAbs(*this, V);
- Constant *Infinity = ConstantFP::getInfinity(V->getType());
- CmpInst::Predicate Pred = (BuiltinID == Builtin::BI__builtin_isinf)
- ? CmpInst::FCMP_OEQ
- : CmpInst::FCMP_ONE;
- Value *FCmp = Builder.CreateFCmp(Pred, Fabs, Infinity, "cmpinf");
- return RValue::get(Builder.CreateZExt(FCmp, ConvertType(E->getType())));
- }
-
- if (Value *Result = getTargetHooks().testFPKind(V, BuiltinID, Builder, CGM))
- return RValue::get(Result);
-
- // Inf values have all exp bits set and a zero significand. Therefore:
- // isinf(V) == ((V << 1) == ((exp mask) << 1))
- // isfinite(V) == ((V << 1) < ((exp mask) << 1)) using unsigned comparison
- unsigned bitsize = Ty->getScalarSizeInBits();
- llvm::IntegerType *IntTy = Builder.getIntNTy(bitsize);
- Value *IntV = Builder.CreateBitCast(V, IntTy);
- Value *Shl1 = Builder.CreateShl(IntV, 1);
- const llvm::fltSemantics &Semantics = Ty->getFltSemantics();
- APInt ExpMask = APFloat::getInf(Semantics).bitcastToAPInt();
- Value *ExpMaskShl1 = llvm::ConstantInt::get(IntTy, ExpMask.shl(1));
- if (BuiltinID == Builtin::BI__builtin_isinf)
- V = Builder.CreateICmpEQ(Shl1, ExpMaskShl1);
- else
- V = Builder.CreateICmpULT(Shl1, ExpMaskShl1);
- return RValue::get(Builder.CreateZExt(V, ConvertType(E->getType())));
- }
-
case Builtin::BI__builtin_isinf_sign: {
// isinf_sign(x) -> fabs(x) == infinity ? (signbit(x) ? -1 : 1) : 0
CodeGenFunction::CGFPOptionsRAII FPOptsRAII(*this, E);
@@ -3453,26 +3431,6 @@ RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID,
return RValue::get(Result);
}
- case Builtin::BI__builtin_isnormal: {
- // isnormal(x) --> x == x && fabsf(x) < infinity && fabsf(x) >= float_min
- CodeGenFunction::CGFPOptionsRAII FPOptsRAII(*this, E);
- // FIXME: for strictfp/IEEE-754 we need to not trap on SNaN here.
- Value *V = EmitScalarExpr(E->getArg(0));
- Value *Eq = Builder.CreateFCmpOEQ(V, V, "iseq");
-
- Value *Abs = EmitFAbs(*this, V);
- Value *IsLessThanInf =
- Builder.CreateFCmpULT(Abs, ConstantFP::getInfinity(V->getType()),"isinf");
- APFloat Smallest = APFloat::getSmallestNormalized(
- getContext().getFloatTypeSemantics(E->getArg(0)->getType()));
- Value *IsNormal =
- Builder.CreateFCmpUGE(Abs, ConstantFP::get(V->getContext(), Smallest),
- "isnormal");
- V = Builder.CreateAnd(Eq, IsLessThanInf, "and");
- V = Builder.CreateAnd(V, IsNormal, "and");
- return RValue::get(Builder.CreateZExt(V, ConvertType(E->getType())));
- }
-
case Builtin::BI__builtin_flt_rounds: {
Function *F = CGM.getIntrinsic(Intrinsic::get_rounding);
diff --git a/clang/test/CodeGen/X86/strictfp_builtins.c b/clang/test/CodeGen/X86/strictfp_builtins.c
index 1fccb2665ee901..43e4060bef259b 100644
--- a/clang/test/CodeGen/X86/strictfp_builtins.c
+++ b/clang/test/CodeGen/X86/strictfp_builtins.c
@@ -1,10 +1,6 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
// RUN: %clang_cc1 %s -emit-llvm -ffp-exception-behavior=maytrap -o - -triple x86_64-unknown-unknown | FileCheck %s
-// Test that the constrained intrinsics are picking up the exception
-// metadata from the AST instead of the global default from the command line.
-// FIXME: these functions shouldn't trap on SNaN.
-
#pragma float_control(except, on)
int printf(const char *, ...);
@@ -17,7 +13,7 @@ int printf(const char *, ...);
// CHECK-NEXT: store i32 [[X:%.*]], ptr [[X_ADDR]], align 4
// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[STR_ADDR]], align 8
// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[X_ADDR]], align 4
-// CHECK-NEXT: [[CALL:%.*]] = call i32 (ptr, ...) @printf(ptr noundef @.str, ptr noundef [[TMP0]], i32 noundef [[TMP1]]) [[ATTR4:#.*]]
+// CHECK-NEXT: [[CALL:%.*]] = call i32 (ptr, ...) @printf(ptr noundef @.str, ptr noundef [[TMP0]], i32 noundef [[TMP1]]) #[[ATTR3:[0-9]+]]
// CHECK-NEXT: ret void
//
void p(char *str, int x) {
@@ -29,13 +25,11 @@ void p(char *str, int x) {
// CHECK-LABEL: @test_long_double_isinf(
// CHECK-NEXT: entry:
// CHECK-NEXT: [[LD_ADDR:%.*]] = alloca x86_fp80, align 16
-// CHECK-NEXT: store x86_fp80 [[D:%.*]], ptr [[LD_ADDR]], align 16
+// CHECK-NEXT: store x86_fp80 [[LD:%.*]], ptr [[LD_ADDR]], align 16
// CHECK-NEXT: [[TMP0:%.*]] = load x86_fp80, ptr [[LD_ADDR]], align 16
-// CHECK-NEXT: [[BITCAST:%.*]] = bitcast x86_fp80 [[TMP0]] to i80
-// CHECK-NEXT: [[SHL1:%.*]] = shl i80 [[BITCAST]], 1
-// CHECK-NEXT: [[CMP:%.*]] = icmp eq i80 [[SHL1]], -18446744073709551616
-// CHECK-NEXT: [[RES:%.*]] = zext i1 [[CMP]] to i32
-// CHECK-NEXT: call void @p(ptr noundef @.str.[[#STRID:1]], i32 noundef [[RES]]) [[ATTR4]]
+// CHECK-NEXT: [[TMP1:%.*]] = call i1 @llvm.is.fpclass.f80(x86_fp80 [[TMP0]], i32 516) #[[ATTR3]]
+// CHECK-NEXT: [[TMP2:%.*]] = zext i1 [[TMP1]] to i32
+// CHECK-NEXT: call void @p(ptr noundef @.str.1, i32 noundef [[TMP2]]) #[[ATTR3]]
// CHECK-NEXT: ret void
//
void test_long_double_isinf(long double ld) {
@@ -47,13 +41,11 @@ void test_long_double_isinf(long double ld) {
// CHECK-LABEL: @test_long_double_isfinite(
// CHECK-NEXT: entry:
// CHECK-NEXT: [[LD_ADDR:%.*]] = alloca x86_fp80, align 16
-// CHECK-NEXT: store x86_fp80 [[D:%.*]], ptr [[LD_ADDR]], align 16
+// CHECK-NEXT: store x86_fp80 [[LD:%.*]], ptr [[LD_ADDR]], align 16
// CHECK-NEXT: [[TMP0:%.*]] = load x86_fp80, ptr [[LD_ADDR]], align 16
-// CHECK-NEXT: [[BITCAST:%.*]] = bitcast x86_fp80 [[TMP0]] to i80
-// CHECK-NEXT: [[SHL1:%.*]] = shl i80 [[BITCAST]], 1
-// CHECK-NEXT: [[CMP:%.*]] = icmp ult i80 [[SHL1]], -18446744073709551616
-// CHECK-NEXT: [[RES:%.*]] = zext i1 [[CMP]] to i32
-// CHECK-NEXT: call void @p(ptr noundef @.str.[[#STRID:STRID+1]], i32 noundef [[RES]]) [[ATTR4]]
+// CHECK-NEXT: [[TMP1:%.*]] = call i1 @llvm.is.fpclass.f80(x86_fp80 [[TMP0]], i32 504) #[[ATTR3]]
+// CHECK-NEXT: [[TMP2:%.*]] = zext i1 [[TMP1]] to i32
+// CHECK-NEXT: call void @p(ptr noundef @.str.2, i32 noundef [[TMP2]]) #[[ATTR3]]
// CHECK-NEXT: ret void
//
void test_long_double_isfinite(long double ld) {
@@ -65,14 +57,11 @@ void test_long_double_isfinite(long double ld) {
// CHECK-LABEL: @test_long_double_isnan(
// CHECK-NEXT: entry:
// CHECK-NEXT: [[LD_ADDR:%.*]] = alloca x86_fp80, align 16
-// CHECK-NEXT: store x86_fp80 [[D:%.*]], ptr [[LD_ADDR]], align 16
+// CHECK-NEXT: store x86_fp80 [[LD:%.*]], ptr [[LD_ADDR]], align 16
// CHECK-NEXT: [[TMP0:%.*]] = load x86_fp80, ptr [[LD_ADDR]], align 16
-// CHECK-NEXT: [[BITCAST:%.*]] = bitcast x86_fp80 [[TMP0]] to i80
-// CHECK-NEXT: [[ABS:%.*]] = and i80 [[BITCAST]], 604462909807314587353087
-// CHECK-NEXT: [[TMP1:%.*]] = sub i80 604453686435277732577280, [[ABS]]
-// CHECK-NEXT: [[ISNAN:%.*]] = lshr i80 [[TMP1]], 79
-// CHECK-NEXT: [[RES:%.*]] = trunc i80 [[ISNAN]] to i32
-// CHECK-NEXT: call void @p(ptr noundef @.str.[[#STRID:STRID+1]], i32 noundef [[RES]]) [[ATTR4]]
+// CHECK-NEXT: [[TMP1:%.*]] = call i1 @llvm.is.fpclass.f80(x86_fp80 [[TMP0]], i32 3) #[[ATTR3]]
+// CHECK-NEXT: [[TMP2:%.*]] = zext i1 [[TMP1]] to i32
+// CHECK-NEXT: call void @p(ptr noundef @.str.3, i32 noundef [[TMP2]]) #[[ATTR3]]
// CHECK-NEXT: ret void
//
void test_long_double_isnan(long double ld) {
diff --git a/clang/test/CodeGen/aarch64-strictfp-builtins.c b/clang/test/CodeGen/aarch64-strictfp-builtins.c
index 285156311ac145..d762ae96965cb3 100644
--- a/clang/test/CodeGen/aarch64-strictfp-builtins.c
+++ b/clang/test/CodeGen/aarch64-strictfp-builtins.c
@@ -29,10 +29,8 @@ void p(char *str, int x) {
// CHECK-NEXT: [[LD_ADDR:%.*]] = alloca fp128, align 16
// CHECK-NEXT: store fp128 [[D:%.*]], ptr [[LD_ADDR]], align 16
// CHECK-NEXT: [[TMP0:%.*]] = load fp128, ptr [[LD_ADDR]], align 16
-// CHECK-NEXT: [[BITCAST:%.*]] = bitcast fp128 [[TMP0]] to i128
-// CHECK-NEXT: [[SHL1:%.*]] = shl i128 [[BITCAST]], 1
-// CHECK-NEXT: [[CMP:%.*]] = icmp eq i128 [[SHL1]], -10384593717069655257060992658440192
-// CHECK-NEXT: [[RES:%.*]] = zext i1 [[CMP]] to i32
+// CHECK-NEXT: [[TMP1:%.*]] = call i1 @llvm.is.fpclass.f128(fp128 [[TMP0]], i32 516)
+// CHECK-NEXT: [[RES:%.*]] = zext i1 [[TMP1]] to i32
// CHECK-NEXT: call void @p(ptr noundef @.str.[[#STRID:1]], i32 noundef [[RES]]) [[ATTR4]]
// CHECK-NEXT: ret void
//
@@ -47,10 +45,8 @@ void test_long_double_isinf(long double ld) {
// CHECK-NEXT: [[LD_ADDR:%.*]] = alloca fp128, align 16
// CHECK-NEXT: store fp128 [[D:%.*]], ptr [[LD_ADDR]], align 16
// CHECK-NEXT: [[TMP0:%.*]] = load fp128, ptr [[LD_ADDR]], align 16
-// CHECK-NEXT: [[BITCAST:%.*]] = bitcast fp128 [[TMP0]] to i128
-// CHECK-NEXT: [[SHL1:%.*]] = shl i128 [[BITCAST]], 1
-// CHECK-NEXT: [[CMP:%.*]] = icmp ult i128 [[SHL1]], -10384593717069655257060992658440192
-// CHECK-NEXT: [[RES:%.*]] = zext i1 [[CMP]] to i32
+// CHECK-NEXT: [[TMP1:%.*]] = call i1 @llvm.is.fpclass.f128(fp128 [[TMP0]], i32 504)
+// CHECK-NEXT: [[RES:%.*]] = zext i1 [[TMP1]] to i32
// CHECK-NEXT: call void @p(ptr noundef @.str.[[#STRID:STRID+1]], i32 noundef [[RES]]) [[ATTR4]]
// CHECK-NEXT: ret void
//
@@ -65,11 +61,8 @@ void test_long_double_isfinite(long double ld) {
// CHECK-NEXT: [[LD_ADDR:%.*]] = alloca fp128, align 16
// CHECK-NEXT: store fp128 [[D:%.*]], ptr [[LD_ADDR]], align 16
// CHECK-NEXT: [[TMP0:%.*]] = load fp128, ptr [[LD_ADDR]], align 16
-// CHECK-NEXT: [[BITCAST:%.*]] = bitcast fp128 [[TMP0]] to i128
-// CHECK-NEXT: [[ABS:%.*]] = and i128 [[BITCAST]], 170141183460469231731687303715884105727
-// CHECK-NEXT: [[TMP1:%.*]] = sub i128 170135991163610696904058773219554885632, [[ABS]]
-// CHECK-NEXT: [[ISNAN:%.*]] = lshr i128 [[TMP1]], 127
-// CHECK-NEXT: [[RES:%.*]] = trunc i128 [[ISNAN]] to i32
+// CHECK-NEXT: [[TMP1:%.*]] = call i1 @llvm.is.fpclass.f128(fp128 [[TMP0]], i32 3)
+// CHECK-NEXT: [[RES:%.*]] = zext i1 [[TMP1]] to i32
// CHECK-NEXT: call void @p(ptr noundef @.str.[[#STRID:STRID+1]], i32 noundef [[RES]])
// CHECK-NEXT: ret void
//
diff --git a/clang/test/CodeGen/builtin_float.c b/clang/test/CodeGen/builtin_float.c
index b257854278cb60..f79b6bd4ff5e08 100644
--- a/clang/test/CodeGen/builtin_float.c
+++ b/clang/test/CodeGen/builtin_float.c
@@ -49,8 +49,8 @@ void test_half(__fp16 *H, __fp16 *H2) {
// CHECK: fcmp ogt float
// CHECK-NEXT: zext i1
(void)__builtin_isinf(*H);
- // NOFP16: fcmp oeq float %{{.+}}, 0x7FF
- // FP16: fcmp oeq half %{{.+}}, 0xH7C
+ // FP16: call i1 @llvm.is.fpclass.f16(half %{{.*}}, i32 516)
+ // NOFP16: call i1 @llvm.is.fpclass.f32(float %{{.*}}, i32 516)
}
void test_mixed(double d1, float f2) {
diff --git a/clang/test/CodeGen/builtin_float_strictfp.c b/clang/test/CodeGen/builtin_float_strictfp.c
index c7940adc9f368a..b7cf567ccd66f7 100644
--- a/clang/test/CodeGen/builtin_float_strictfp.c
+++ b/clang/test/CodeGen/builtin_float_strictfp.c
@@ -19,15 +19,11 @@ void test_half(__fp16 *H, __fp16 *H2) {
// NOFP16: [[LDADDR:%.*]] = load ptr, ptr %{{.*}}, align 8
// NOFP16-NEXT: [[IHALF:%.*]] = load i16, ptr [[LDADDR]], align 2
// NOFP16-NEXT: [[CONV:%.*]] = call float @llvm.convert.from.fp16.f32(i16 [[IHALF]])
- // NOFP16-NEXT: [[IFLOAT:%.*]] = bitcast float [[CONV]] to i32
- // NOFP16-NEXT: [[SHL:%.*]] = shl i32 [[IFLOAT]], 1
- // NOFP16-NEXT: [[RES1:%.*]] = icmp eq i32 [[SHL]], -16777216
+ // NOFP16-NEXT: [[RES1:%.*]] = call i1 @llvm.is.fpclass.f32(float [[CONV]], i32 516)
// NOFP16-NEXT: zext i1 [[RES1]] to i32
// FP16: [[LDADDR:%.*]] = load ptr, ptr %{{.*}}, align 8
// FP16-NEXT: [[HALF:%.*]] = load half, ptr [[LDADDR]], align 2
- // FP16-NEXT: [[IHALF:%.*]] = bitcast half [[HALF]] to i16
- // FP16-NEXT: [[SHL:%.*]] = shl i16 [[IHALF]], 1
- // FP16-NEXT: [[RES1:%.*]] = icmp eq i16 [[SHL]], -2048
+ // FP16-NEXT: [[RES1:%.*]] = call i1 @llvm.is.fpclass.f16(half [[HALF]], i32 516)
// FP16-NEXT: zext i1 [[RES1]] to i32
}
diff --git a/clang/test/CodeGen/builtins.c b/clang/test/CodeGen/builtins.c
index c42860923568fd..1b1b2cd6413a34 100644
--- a/clang/test/CodeGen/builtins.c
+++ b/clang/test/CodeGen/builtins.c
@@ -63,6 +63,7 @@ int main(void) {
P(isinf, (1.));
P(isinf_sign, (1.));
P(isnan, (1.));
+ P(isfinite, (1.));
P(isfpclass, (1., 1));
// Bitwise & Numeric Functions
@@ -202,20 +203,20 @@ void test_conditional_bzero(void) {
void test_float_builtins(__fp16 *H, float F, double D, long double LD) {
volatile int res;
res = __builtin_isinf(*H);
- // CHECK: call half @llvm.fabs.f16(half
- // CHECK: fcmp oeq half {{.*}}, 0xH7C00
+ // CHECK: [[TMP:%.*]] = call i1 @llvm.is.fpclass.f16(half {{.*}}, i32 516)
+ // CHECK: zext i1 [[TMP]] to i32
res = __builtin_isinf(F);
- // CHECK: call float @llvm.fabs.f32(float
- // CHECK: fcmp oeq float {{.*}}, 0x7FF0000000000000
+ // CHECK: [[TMP:%.*]] = call i1 @llvm.is.fpclass.f32(float {{.*}}, i32 516)
+ // CHECK: zext i1 [[TMP]] to i32
res = __builtin_isinf(D);
- // CHECK: call double @llvm.fabs.f64(double
- // CHECK: fcmp oeq double {{.*}}, 0x7FF0000000000000
+ // CHECK: [[TMP:%.*]] = call i1 @llvm.is.fpclass.f64(double {{.*}}, i32 516)
+ // CHECK: zext i1 [[TMP]] to i32
res = __builtin_isinf(LD);
- // CHECK: call x86_fp80 @llvm.fabs.f80(x86_fp80
- // CHECK: fcmp oeq x86_fp80 {{.*}}, 0xK7FFF8000000000000000
+ // CHECK: [[TMP:%.*]] = call i1 @llvm.is.fpclass.f80(x86_fp80 {{.*}}, i32 516)
+ // CHECK: zext i1 [[TMP]] to i32
res = __builtin_isinf_sign(*H);
// CHECK: %[[ABS:.*]] = call half @llvm.fabs.f16(half %[[ARG:.*]])
@@ -250,32 +251,24 @@ void test_float_builtins(__fp16 *H, float F, double D, long double LD) {
// CHECK: select i1 %[[ISINF]], i32 %[[SIGN]], i32 0
res = __builtin_isfinite(*H);
- // CHECK: call half @llvm.fabs.f16(half
- // CHECK: fcmp one half {{.*}}, 0xH7C00
+ // CHECK: [[TMP:%.*]] = call i1 @llvm.is.fpclass.f16(half {{.*}}, i32 504)
+ // CHECK: zext i1 [[TMP]] to i32
res = __builtin_isfinite(F);
- // CHECK: call float @llvm.fabs.f32(float
- // CHECK: fcmp one float {{.*}}, 0x7FF0000000000000
+ // CHECK: [[TMP:%.*]] = call i1 @llvm.is.fpclass.f32(float {{.*}}, i32 504)
+ // CHECK: zext i1 [[TMP]] to i32
res = finite(D);
- // CHECK: call double @llvm.fabs.f64(double
- // CHECK: fcmp one double {{.*}}, 0x7FF0000000000000
+ // CHECK: [[TMP:%.*]] = call i1 @llvm.is.fpclass.f64(double {{.*}}, i32 504)
+ // CHECK: zext i1 [[TMP]] to i32
res = __builtin_isnormal(*H);
- // CHECK: fcmp oeq half
- // CHECK: call half @llvm.fabs.f16(half
- // CHECK: fcmp ult half {{.*}}, 0xH7C00
- // CHECK: fcmp uge half {{.*}}, 0xH0400
- // CHECK: and i1
- // CHECK: and i1
+ // CHECK: [[TMP:%.*]] = call i1 @llvm.is.fpclass.f16(half {{.*}}, i32 264)
+ // CHECK: zext i1 [[TMP]] to i32
res = __builtin_isnormal(F);
- // CHECK: fcmp oeq float
- // CHECK: call float @llvm.fabs.f32(float
- // CHECK: fcmp ult float {{.*}}, 0x7FF0000000000000
- // CHECK: fcmp uge float {{.*}}, 0x3810000000000000
- // CHECK: and i1
- // CHECK: and i1
+ // CHECK: [[TMP:%.*]] = call i1 @llvm.is.fpclass.f32(float {{.*}}, i32 264)
+ // CHECK: zext i1 [[TMP]] to i32
res = __builtin_flt_rounds();
// CHECK: call i32 @llvm.get.rounding(
diff --git a/clang/test/CodeGen/isfpclass.c b/clang/test/CodeGen/isfpclass.c
index 38d5ed85531b66..4b436f8f9e1feb 100644
--- a/clang/test/CodeGen/isfpclass.c
+++ b/clang/test/CodeGen/isfpclass.c
@@ -1,5 +1,5 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 2
-// RUN: %clang_cc1 -triple aarch64-linux-gnu -target-feature +avx512fp16 -S -O1 -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -triple aarch64-linux-gnu -S -O1 -emit-llvm %s -o - | FileCheck %s
// CHECK-LABEL: define dso_local i1 @check_isfpclass_finite
// CHECK-SAME: (float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {
@@ -84,3 +84,24 @@ _Bool check_isfpclass_zero_f16_strict(_Float16 x) {
#pragma STDC FENV_ACCESS ON
return __builtin_isfpclass(x, 96 /*Zero*/);
}
+
+_Bool check_isnan(float x) {
+#pragma STDC FENV_ACCESS ON
+ return __builtin_isnan(x);
+}
+
+_Bool check_isinf(float x) {
+#pragma STDC FENV_ACCESS ON
+ return __builtin_isinf(x);
+}
+
+_Bool check_isfinite(float x) {
+#pragma STDC FENV_ACCESS ON
+ return __builtin_isfinite(x);
+}
+
+_Bool check_isnormal(float x) {
+#pragma STDC FENV_ACCESS ON
+ return __builtin_isnormal(x);
+}
+
diff --git a/clang/test/CodeGen/strictfp_builtins.c b/clang/test/CodeGen/strictfp_builtins.c
index 4d34dfd2be25f9..58815c7de4fa94 100644
--- a/clang/test/CodeGen/strictfp_builtins.c
+++ b/clang/test/CodeGen/strictfp_builtins.c
@@ -17,7 +17,7 @@ int printf(const char *, ...);
// CHECK-NEXT: store i32 [[X:%.*]], ptr [[X_ADDR]], align 4
// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[STR_ADDR]], align 8
// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[X_ADDR]], align 4
-// CHECK-NEXT: [[CALL:%.*]] = call i32 (ptr, ...) @printf(ptr noundef @.str, ptr noundef [[TMP0]], i32 noundef [[TMP1]]) [[ATTR4:#.*]]
+// CHECK-NEXT: [[CALL:%.*]] = call i32 (ptr, ...) @printf(ptr noundef @.str, ptr noundef [[TMP0]], i32 noundef [[TMP1]]) #[[ATTR4:[0-9]+]]
// CHECK-NEXT: ret void
//
void p(char *str, int x) {
@@ -31,21 +31,21 @@ void p(char *str, int x) {
// CHECK-NEXT: [[D_ADDR:%.*]] = alloca double, align 8
// CHECK-NEXT: store double [[D:%.*]], ptr [[D_ADDR]], align 8
// CHECK-NEXT: [[TMP0:%.*]] = load double, ptr [[D_ADDR]], align 8
-// CHECK-NEXT: [[ISZERO:%.*]] = call i1 @llvm.experimental.constrained.fcmp.f64(double [[TMP0]], double 0.000000e+00, metadata !"oeq", metadata !"fpexcept.strict") [[ATTR4]]
+// CHECK-NEXT: [[ISZERO:%.*]] = call i1 @llvm.experimental.constrained.fcmp.f64(double [[TMP0]], double 0.000000e+00, metadata !"oeq", metadata !"fpexcept.strict") #[[ATTR4]]
// CHECK-NEXT: br i1 [[ISZERO]], label [[FPCLASSIFY_END:%.*]], label [[FPCLASSIFY_NOT_ZERO:%.*]]
// CHECK: fpclassify_end:
// CHECK-NEXT: [[FPCLASSIFY_RESULT:%.*]] = phi i32 [ 4, [[ENTRY:%.*]] ], [ 0, [[FPCLASSIFY_NOT_ZERO]] ], [ 1, [[FPCLASSIFY_NOT_NAN:%.*]] ], [ [[TMP2:%.*]], [[FPCLASSIFY_NOT_INF:%.*]] ]
-// CHECK-NEXT: call void @p(ptr noundef @.str.1, i32 noundef [[FPCLASSIFY_RESULT]]) [[ATTR4]]
+// CHECK-NEXT: call void @p(ptr noundef @.str.1, i32 noundef [[FPCLASSIFY_RESULT]]) #[[ATTR4]]
// CHECK-NEXT: ret void
// CHECK: fpclassify_not_zero:
-// CHECK-NEXT: [[CMP:%.*]] = call i1 @llvm.experimental.constrained.fcmp.f64(double [[TMP0]], double [[TMP0]], metadata !"uno", metadata !"fpexcept.strict") [[ATTR4]]
+// CHECK-NEXT: [[CMP:%.*]] = call i1 @llvm.experimental.constrained.fcmp.f64(double [[TMP0]], double [[TMP0]], metadata !"uno", metadata !"fpexcept.strict") #[[ATTR4]]
// CHECK-NEXT: br i1 [[CMP]], label [[FPCLASSIFY_END]], label [[FPCLASSIFY_NOT_NAN]]
// CHECK: fpclassify_not_nan:
-// CHECK-NEXT: [[TMP1:%.*]] = call double @llvm.fabs.f64(double [[TMP0]]) [[ATTR5:#.*]]
-// CHECK-NEXT: [[ISINF:%.*]] = call i1 @llvm.experimental.constrained.fcmp.f64(double [[TMP1]], double 0x7FF0000000000000, metadata !"oeq", metadata !"fpexcept.strict") [[ATTR4]]
+// CHECK-NEXT: [[TMP1:%.*]] = call double @llvm.fabs.f64(double [[TMP0]]) #[[ATTR5:[0-9]+]]
+// CHECK-NEXT: [[ISINF:%.*]] = call i1 @llvm.experimental.constrained.fcmp.f64(double [[TMP1]], double 0x7FF0000000000000, metadata !"oeq", metadata !"fpexcept.strict") #[[ATTR4]]
// CHECK-NEXT: br i1 [[ISINF]], label [[FPCLASSIFY_END]], label [[FPCLASSIFY_NOT_INF]]
// CHECK: fpclassify_not_inf:
-// CHECK-NEXT: [[ISNORMAL:%.*]] = call i1 @llvm.experimental.constrained.fcmp.f64(double [[TMP1]], double 0x10000000000000, metadata !"uge", metadata !"fpexcept.strict") [[ATTR4]]
+// CHECK-NEXT: [[ISNORMAL:%.*]] = call i1 @llvm.experimental.constrained.fcmp.f64(double [[TMP1]], double 0x10000000000000, metadata !"uge", metadata !"fpexcept.strict") #[[ATTR4]]
// CHECK-NEXT: [[TMP2]] = select i1 [[ISNORMAL]], i32 2, i32 3
// CHECK-NEXT: br label [[FPCLASSIFY_END]]
//
@@ -57,14 +57,12 @@ void test_fpclassify(double d) {
// CHECK-LABEL: @test_fp16_isinf(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[LD_ADDR:%.*]] = alloca half, align 2
-// CHECK-NEXT: store half [[H:%.*]], ptr [[LD_ADDR]], align 2
-// CHECK-NEXT: [[TMP0:%.*]] = load half, ptr [[LD_ADDR]], align 2
-// CHECK-NEXT: [[BITCAST:%.*]] = bitcast half [[TMP0]] to i16
-// CHECK-NEXT: [[SHL1:%.*]] = shl i16 [[BITCAST]], 1
-// CHECK-NEXT: [[CMP:%.*]] = icmp eq i16 [[SHL1]], -2048
-// CHECK-NEXT: [[RES:%.*]] = zext i1 [[CMP]] to i32
-// CHECK-NEXT: call void @p(ptr noundef @.str.[[#STRID:2]], i32 noundef [[RES]]) [[ATTR4]]
+// CHECK-NEXT: [[H_ADDR:%.*]] = alloca half, align 2
+// CHECK-NEXT: store half [[H:%.*]], ptr [[H_ADDR]], align 2
+// CHECK-NEXT: [[TMP0:%.*]] = load half, ptr [[H_ADDR]], align 2
+// CHECK-NEXT: [[TMP1:%.*]] = call i1 @llvm.is.fpclass.f16(half [[TMP0]], i32 516) #[[ATTR4]]
+// CHECK-NEXT: [[TMP2:%.*]] = zext i1 [[TMP1]] to i32
+// CHECK-NEXT: call void @p(ptr noundef @.str.2, i32 noundef [[TMP2]]) #[[ATTR4]]
// CHECK-NEXT: ret void
//
void test_fp16_isinf(_Float16 h) {
@@ -75,14 +73,12 @@ void test_fp16_isinf(_Float16 h) {
// CHECK-LABEL: @test_float_isinf(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[LD_ADDR:%.*]] = alloca float, align 4
-// CHECK-NEXT: store float [[F:%.*]], ptr [[LD_ADDR]], align 4
-// CHECK-NEXT: [[TMP0:%.*]] = load float, ptr [[LD_ADDR]], align 4
-// CHECK-NEXT: [[BITCAST:%.*]] = bitcast float [[TMP0]] to i32
-// CHECK-NEXT: [[SHL1:%.*]] = shl i32 [[BITCAST]], 1
-// CHECK-NEXT: [[CMP:%.*]] = icmp eq i32 [[SHL1]], -16777216
-// CHECK-NEXT: [[RES:%.*]] = zext i1 [[CMP]] to i32
-// CHECK-NEXT: call void @p(ptr noundef @.str.[[#STRID:STRID+1]], i32 noundef [[RES]]) [[ATTR4]]
+// CHECK-NEXT: [[F_ADDR:%.*]] = alloca float, align 4
+// CHECK-NEXT: store float [[F:%.*]], ptr [[F_ADDR]], align 4
+// CHECK-NEXT: [[TMP0:%.*]] = load float, ptr [[F_ADDR]], align 4
+// CHECK-NEXT: [[TMP1:%.*]] = call i1 @llvm.is.fpclass.f32(float [[TMP0]], i32 516) #[[ATTR4]]
+// CHECK-NEXT: [[TMP2:%.*]] = zext i1 [[TMP1]] to i32
+// CHECK-NEXT: call void @p(ptr noundef @.str.3, i32 noundef [[TMP2]]) #[[ATTR4]]
// CHECK-NEXT: ret void
//
void test_float_isinf(float f) {
@@ -93,14 +89,12 @@ void test_float_isinf(float f) {
// CHECK-LABEL: @test_double_isinf(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[LD_ADDR:%.*]] = alloca double, align 8
-// CHECK-NEXT: store double [[D:%.*]], ptr [[LD_ADDR]], align 8
-// CHECK-NEXT: [[TMP0:%.*]] = load double, ptr [[LD_ADDR]], align 8
-// CHECK-NEXT: [[BITCAST:%.*]] = bitcast double [[TMP0]] to i64
-// CHECK-NEXT: [[SHL1:%.*]] = shl i64 [[BITCAST]], 1
-// CHECK-NEXT: [[CMP:%.*]] = icmp eq i64 [[SHL1]], -9007199254740992
-// CHECK-NEXT: [[RES:%.*]] = zext i1 [[CMP]] to i32
-// CHECK-NEXT: call void @p(ptr noundef @.str.[[#STRID:STRID+1]], i32 noundef [[RES]]) [[ATTR4]]
+// CHECK-NEXT: [[D_ADDR:%.*]] = alloca double, align 8
+// CHECK-NEXT: store double [[D:%.*]], ptr [[D_ADDR]], align 8
+// CHECK-NEXT: [[TMP0:%.*]] = load double, ptr [[D_ADDR]], align 8
+// CHECK-NEXT: [[TMP1:%.*]] = call i1 @llvm.is.fpclass.f64(double [[TMP0]], i32 516) #[[ATTR4]]
+// CHECK-NEXT: [[TMP2:%.*]] = zext i1 [[TMP1]] to i32
+// CHECK-NEXT: call void @p(ptr noundef @.str.4, i32 noundef [[TMP2]]) #[[ATTR4]]
// CHECK-NEXT: ret void
//
void test_double_isinf(double d) {
@@ -111,14 +105,12 @@ void test_double_isinf(double d) {
// CHECK-LABEL: @test_fp16_isfinite(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[LD_ADDR:%.*]] = alloca half, align 2
-// CHECK-NEXT: store half [[H:%.*]], ptr [[LD_ADDR]], align 2
-// CHECK-NEXT: [[TMP0:%.*]] = load half, ptr [[LD_ADDR]], align 2
-// CHECK-NEXT: [[BITCAST:%.*]] = bitcast half [[TMP0]] to i16
-// CHECK-NEXT: [[SHL1:%.*]] = shl i16 [[BITCAST]], 1
-// CHECK-NEXT: [[CMP:%.*]] = icmp ult i16 [[SHL1]], -2048
-// CHECK-NEXT: [[RES:%.*]] = zext i1 [[CMP]] to i32
-// CHECK-NEXT: call void @p(ptr noundef @.str.[[#STRID:STRID+1]], i32 noundef [[RES]]) [[ATTR4]]
+// CHECK-NEXT: [[H_ADDR:%.*]] = alloca half, align 2
+// CHECK-NEXT: store half [[H:%.*]], ptr [[H_ADDR]], align 2
+// CHECK-NEXT: [[TMP0:%.*]] = load half, ptr [[H_ADDR]], align 2
+// CHECK-NEXT: [[TMP1:%.*]] = call i1 @llvm.is.fpclass.f16(half [[TMP0]], i32 504) #[[ATTR4]]
+// CHECK-NEXT: [[TMP2:%.*]] = zext i1 [[TMP1]] to i32
+// CHECK-NEXT: call void @p(ptr noundef @.str.5, i32 noundef [[TMP2]]) #[[ATTR4]]
// CHECK-NEXT: ret void
//
void test_fp16_isfinite(_Float16 h) {
@@ -129,14 +121,12 @@ void test_fp16_isfinite(_Float16 h) {
// CHECK-LABEL: @test_float_isfinite(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[LD_ADDR:%.*]] = alloca float, align 4
-// CHECK-NEXT: store float [[F:%.*]], ptr [[LD_ADDR]], align 4
-// CHECK-NEXT: [[TMP0:%.*]] = load float, ptr [[LD_ADDR]], align 4
-// CHECK-NEXT: [[BITCAST:%.*]] = bitcast float [[TMP0]] to i32
-// CHECK-NEXT: [[SHL1:%.*]] = shl i32 [[BITCAST]], 1
-// CHECK-NEXT: [[CMP:%.*]] = icmp ult i32 [[SHL1]], -16777216
-// CHECK-NEXT: [[RES:%.*]] = zext i1 [[CMP]] to i32
-// CHECK-NEXT: call void @p(ptr noundef @.str.[[#STRID:STRID+1]], i32 noundef [[RES]]) [[ATTR4]]
+// CHECK-NEXT: [[F_ADDR:%.*]] = alloca float, align 4
+// CHECK-NEXT: store float [[F:%.*]], ptr [[F_ADDR]], align 4
+// CHECK-NEXT: [[TMP0:%.*]] = load float, ptr [[F_ADDR]], align 4
+// CHECK-NEXT: [[TMP1:%.*]] = call i1 @llvm.is.fpclass.f32(float [[TMP0]], i32 504) #[[ATTR4]]
+// CHECK-NEXT: [[TMP2:%.*]] = zext i1 [[TMP1]] to i32
+// CHECK-NEXT: call void @p(ptr noundef @.str.6, i32 noundef [[TMP2]]) #[[ATTR4]]
// CHECK-NEXT: ret void
//
void test_float_isfinite(float f) {
@@ -147,14 +137,12 @@ void test_float_isfinite(float f) {
// CHECK-LABEL: @test_double_isfinite(
// CHECK-NEXT: entry:
-// CHECK-NEXT: [[LD_ADDR:%.*]] = alloca double, align 8
-// CHECK-NEXT: store double [[D:%.*]], ptr [[LD_ADDR]], align 8
-// CHECK-NEXT: [[TMP0:%.*]] = load double, ptr [[LD_ADDR]], align 8
-// CHECK-NEXT: [[BITCAST:%.*]] = bitcast double [[TMP0]] to i64
-// CHECK-NEXT: [[SHL1:%.*]] = shl i64 [[BITCAST]], 1
-// CHECK-NEXT: [[CMP:%.*]] = icmp ult i64 [[SHL1]], -9007199254740992
-// CHECK-NEXT: [[RES:%.*]] = zext i1 [[CMP]] to i32
-// CHECK-NEXT: call void @p(ptr noundef @.str.[[#STRID:STRID+1]], i32 noundef [[RES]]) [[ATTR4]]
+// CHECK-NEXT: [[D_ADDR:%.*]] = alloca double, align 8
+// CHECK-NEXT: store double [[D:%.*]], ptr [[D_ADDR]], align 8
+// CHECK-NEXT: [[TMP0:%.*]] = load double, ptr [[D_ADDR]], align 8
+// CHECK-NEXT: [[TMP1:%.*]] = call i1 @llvm.is.fpclass.f64(double [[TMP0]], i32 504) #[[ATTR4]]
+// CHECK-NEXT: [[TMP2:%.*]] = zext i1 [[TMP1]] to i32
+// CHECK-NEXT: call void @p(ptr noundef @.str.7, i32 noundef [[TMP2]]) #[[ATTR4]]
// CHECK-NEXT: ret void
//
void test_double_isfinite(double d) {
@@ -168,13 +156,13 @@ void test_double_isfinite(double d) {
// CHECK-NEXT: [[D_ADDR:%.*]] = alloca double, align 8
// CHECK-NEXT: store double [[D:%.*]], ptr [[D_ADDR]], align 8
// CHECK-NEXT: [[TMP0:%.*]] = load double, ptr [[D_ADDR]], align 8
-// CHECK-NEXT: [[TMP1:%.*]] = call double @llvm.fabs.f64(double [[TMP0]]) [[ATTR5]]
-// CHECK-NEXT: [[ISINF:%.*]] = call i1 @llvm.experimental.constrained.fcmp.f64(double [[TMP1]], double 0x7FF0000000000000, metadata !"oeq", metadata !"fpexcept.strict") [[ATTR4]]
+// CHECK-NEXT: [[TMP1:%.*]] = call double @llvm.fabs.f64(double [[TMP0]]) #[[ATTR5]]
+// CHECK-NEXT: [[ISINF:%.*]] = call i1 @llvm.experimental.constrained.fcmp.f64(double [[TMP1]], double 0x7FF0000000000000, metadata !"oeq", metadata !"fpexcept.strict") #[[ATTR4]]
// CHECK-NEXT: [[TMP2:%.*]] = bitcast double [[TMP0]] to i64
// CHECK-NEXT: [[TMP3:%.*]] = icmp slt i64 [[TMP2]], 0
// CHECK-NEXT: [[TMP4:%.*]] = select i1 [[TMP3]], i32 -1, i32 1
// CHECK-NEXT: [[TMP5:%.*]] = select i1 [[ISINF]], i32 [[TMP4]], i32 0
-// CHECK-NEXT: call void @p(ptr noundef @.str.[[#STRID:STRID+1]], i32 noundef [[TMP5]]) [[ATTR4]]
+// CHECK-NEXT: call void @p(ptr noundef @.str.8, i32 noundef [[TMP5]]) #[[ATTR4]]
// CHECK-NEXT: ret void
//
void test_isinf_sign(double d) {
@@ -188,12 +176,9 @@ void test_isinf_sign(double d) {
// CHECK-NEXT: [[H_ADDR:%.*]] = alloca half, align 2
// CHECK-NEXT: store half [[H:%.*]], ptr [[H_ADDR]], align 2
// CHECK-NEXT: [[TMP0:%.*]] = load half, ptr [[H_ADDR]], align 2
-// CHECK-NEXT: [[BITCAST:%.*]] = bitcast half [[TMP0]] to i16
-// CHECK-NEXT: [[ABS:%.*]] = and i16 [[BITCAST]], [[#%u,0x7FFF]]
-// CHECK-NEXT: [[TMP1:%.*]] = sub i16 [[#%u,0x7C00]], [[ABS]]
-// CHECK-NEXT: [[ISNAN:%.*]] = lshr i16 [[TMP1]], 15
-// CHECK-NEXT: [[RES:%.*]] = zext i16 [[ISNAN]] to i32
-// CHECK-NEXT: call void @p(ptr noundef @.str.[[#STRID:STRID+1]], i32 noundef [[RES]]) [[ATTR4]]
+// CHECK-NEXT: [[TMP1:%.*]] = call i1 @llvm.is.fpclass.f16(half [[TMP0]], i32 3) #[[ATTR4]]
+// CHECK-NEXT: [[TMP2:%.*]] = zext i1 [[TMP1]] to i32
+// CHECK-NEXT: call void @p(ptr noundef @.str.9, i32 noundef [[TMP2]]) #[[ATTR4]]
// CHECK-NEXT: ret void
//
void test_fp16_isnan(_Float16 h) {
@@ -207,11 +192,9 @@ void test_fp16_isnan(_Float16 h) {
// CHECK-NEXT: [[F_ADDR:%.*]] = alloca float, align 4
// CHECK-NEXT: store float [[F:%.*]], ptr [[F_ADDR]], align 4
// CHECK-NEXT: [[TMP0:%.*]] = load float, ptr [[F_ADDR]], align 4
-// CHECK-NEXT: [[BITCAST:%.*]] = bitcast float [[TMP0]] to i32
-// CHECK-NEXT: [[ABS:%.*]] = and i32 [[BITCAST]], [[#%u,0x7FFFFFFF]]
-// CHECK-NEXT: [[TMP1:%.*]] = sub i32 [[#%u,0x7F800000]], [[ABS]]
-// CHECK-NEXT: [[ISNAN:%.*]] = lshr i32 [[TMP1]], 31
-// CHECK-NEXT: call void @p(ptr noundef @.str.[[#STRID:STRID+1]], i32 noundef [[ISNAN]]) [[ATTR4]]
+// CHECK-NEXT: [[TMP1:%.*]] = call i1 @llvm.is.fpclass.f32(float [[TMP0]], i32 3) #[[ATTR4]]
+// CHECK-NEXT: [[TMP2:%.*]] = zext i1 [[TMP1]] to i32
+// CHECK-NEXT: call void @p(ptr noundef @.str.10, i32 noundef [[TMP2]]) #[[ATTR4]]
// CHECK-NEXT: ret void
//
void test_float_isnan(float f) {
@@ -225,12 +208,9 @@ void test_float_isnan(float f) {
// CHECK-NEXT: [[D_ADDR:%.*]] = alloca double, align 8
// CHECK-NEXT: store double [[D:%.*]], ptr [[D_ADDR]], align 8
// CHECK-NEXT: [[TMP0:%.*]] = load double, ptr [[D_ADDR]], align 8
-// CHECK-NEXT: [[BITCAST:%.*]] = bitcast double [[TMP0]] to i64
-// CHECK-NEXT: [[ABS:%.*]] = and i64 [[BITCAST]], [[#%u,0x7FFFFFFFFFFFFFFF]]
-// CHECK-NEXT: [[TMP1:%.*]] = sub i64 [[#%u,0x7FF0000000000000]], [[ABS]]
-// CHECK-NEXT: [[ISNAN:%.*]] = lshr i64 [[TMP1]], 63
-// CHECK-NEXT: [[RES:%.*]] = trunc i64 [[ISNAN]] to i32
-// CHECK-NEXT: call void @p(ptr noundef @.str.[[#STRID:STRID+1]], i32 noundef [[RES]]) [[ATTR4]]
+// CHECK-NEXT: [[TMP1:%.*]] = call i1 @llvm.is.fpclass.f64(double [[TMP0]], i32 3) #[[ATTR4]]
+// CHECK-NEXT: [[TMP2:%.*]] = zext i1 [[TMP1]] to i32
+// CHECK-NEXT: call void @p(ptr noundef @.str.11, i32 noundef [[TMP2]]) #[[ATTR4]]
// CHECK-NEXT: ret void
//
void test_double_isnan(double d) {
@@ -244,14 +224,9 @@ void test_double_isnan(double d) {
// CHECK-NEXT: [[D_ADDR:%.*]] = alloca double, align 8
// CHECK-NEXT: store double [[D:%.*]], ptr [[D_ADDR]], align 8
// CHECK-NEXT: [[TMP0:%.*]] = load double, ptr [[D_ADDR]], align 8
-// CHECK-NEXT: [[ISEQ:%.*]] = call i1 @llvm.experimental.constrained.fcmp.f64(double [[TMP0]], double [[TMP0]], metadata !"oeq", metadata !"fpexcept.strict") [[ATTR4]]
-// CHECK-NEXT: [[TMP1:%.*]] = call double @llvm.fabs.f64(double [[TMP0]]) [[ATTR5]]
-// CHECK-NEXT: [[ISINF:%.*]] = call i1 @llvm.experimental.constrained.fcmp.f64(double [[TMP1]], double 0x7FF0000000000000, metadata !"ult", metadata !"fpexcept.strict") [[ATTR4]]
-// CHECK-NEXT: [[ISNORMAL:%.*]] = call i1 @llvm.experimental.constrained.fcmp.f64(double [[TMP1]], double 0x10000000000000, metadata !"uge", metadata !"fpexcept.strict") [[ATTR4]]
-// CHECK-NEXT: [[AND:%.*]] = and i1 [[ISEQ]], [[ISINF]]
-// CHECK-NEXT: [[AND1:%.*]] = and i1 [[AND]], [[ISNORMAL]]
-// CHECK-NEXT: [[TMP2:%.*]] = zext i1 [[AND1]] to i32
-// CHECK-NEXT: call void @p(ptr noundef @.str.[[#STRID:STRID+1]], i32 noundef [[TMP2]]) [[ATTR4]]
+// CHECK-NEXT: [[TMP1:%.*]] = call i1 @llvm.is.fpclass.f64(double [[TMP0]], i32 264) #[[ATTR4]]
+// CHECK-NEXT: [[TMP2:%.*]] = zext i1 [[TMP1]] to i32
+// CHECK-NEXT: call void @p(ptr noundef @.str.12, i32 noundef [[TMP2]]) #[[ATTR4]]
// CHECK-NEXT: ret void
//
void test_isnormal(double d) {
diff --git a/clang/test/Headers/__clang_hip_math.hip b/clang/test/Headers/__clang_hip_math.hip
index 35f637b912f439..21864d2ecb4dfd 100644
--- a/clang/test/Headers/__clang_hip_math.hip
+++ b/clang/test/Headers/__clang_hip_math.hip
@@ -1143,9 +1143,8 @@ extern "C" __device__ int test_ilogb(double x) {
// DEFAULT-LABEL: @test___finitef(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.fabs.f32(float [[X:%.*]]) #[[ATTR17:[0-9]+]]
-// DEFAULT-NEXT: [[CMPINF_I:%.*]] = fcmp contract one float [[TMP0]], 0x7FF0000000000000
-// DEFAULT-NEXT: [[CONV:%.*]] = zext i1 [[CMPINF_I]] to i32
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call i1 @llvm.is.fpclass.f32(float [[X:%.*]], i32 504)
+// DEFAULT-NEXT: [[CONV:%.*]] = zext i1 [[TMP0]] to i32
// DEFAULT-NEXT: ret i32 [[CONV]]
//
// FINITEONLY-LABEL: @test___finitef(
@@ -1158,9 +1157,8 @@ extern "C" __device__ BOOL_TYPE test___finitef(float x) {
// DEFAULT-LABEL: @test___finite(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.fabs.f64(double [[X:%.*]]) #[[ATTR17]]
-// DEFAULT-NEXT: [[CMPINF_I:%.*]] = fcmp contract one double [[TMP0]], 0x7FF0000000000000
-// DEFAULT-NEXT: [[CONV:%.*]] = zext i1 [[CMPINF_I]] to i32
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call i1 @llvm.is.fpclass.f64(double [[X:%.*]], i32 504)
+// DEFAULT-NEXT: [[CONV:%.*]] = zext i1 [[TMP0]] to i32
// DEFAULT-NEXT: ret i32 [[CONV]]
//
// FINITEONLY-LABEL: @test___finite(
@@ -1173,9 +1171,8 @@ extern "C" __device__ BOOL_TYPE test___finite(double x) {
// DEFAULT-LABEL: @test___isinff(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.fabs.f32(float [[X:%.*]]) #[[ATTR17]]
-// DEFAULT-NEXT: [[CMPINF_I:%.*]] = fcmp contract oeq float [[TMP0]], 0x7FF0000000000000
-// DEFAULT-NEXT: [[CONV:%.*]] = zext i1 [[CMPINF_I]] to i32
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call i1 @llvm.is.fpclass.f32(float [[X:%.*]], i32 516)
+// DEFAULT-NEXT: [[CONV:%.*]] = zext i1 [[TMP0]] to i32
// DEFAULT-NEXT: ret i32 [[CONV]]
//
// FINITEONLY-LABEL: @test___isinff(
@@ -1188,9 +1185,8 @@ extern "C" __device__ BOOL_TYPE test___isinff(float x) {
// DEFAULT-LABEL: @test___isinf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.fabs.f64(double [[X:%.*]]) #[[ATTR17]]
-// DEFAULT-NEXT: [[CMPINF_I:%.*]] = fcmp contract oeq double [[TMP0]], 0x7FF0000000000000
-// DEFAULT-NEXT: [[CONV:%.*]] = zext i1 [[CMPINF_I]] to i32
+// DEFAULT-NEXT: [[TMP0:%.*]] = tail call i1 @llvm.is.fpclass.f64(double [[X:%.*]], i32 516)
+// DEFAULT-NEXT: [[CONV:%.*]] = zext i1 [[TMP0]] to i32
// DEFAULT-NEXT: ret i32 [[CONV]]
//
// FINITEONLY-LABEL: @test___isinf(
@@ -1203,8 +1199,8 @@ extern "C" __device__ BOOL_TYPE test___isinf(double x) {
// DEFAULT-LABEL: @test___isnanf(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CMP_I:%.*]] = fcmp contract uno float [[X:%.*]], 0.000000e+00
-// DEFAULT-NEXT: [[CONV:%.*]] = zext i1 [[CMP_I]] to i32
+// DEFAULT-NEXT: [[TMP0:%.*]] = fcmp uno float [[X:%.*]], 0.000000e+00
+// DEFAULT-NEXT: [[CONV:%.*]] = zext i1 [[TMP0]] to i32
// DEFAULT-NEXT: ret i32 [[CONV]]
//
// FINITEONLY-LABEL: @test___isnanf(
@@ -1217,8 +1213,8 @@ extern "C" __device__ BOOL_TYPE test___isnanf(float x) {
// DEFAULT-LABEL: @test___isnan(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CMP_I:%.*]] = fcmp contract uno double [[X:%.*]], 0.000000e+00
-// DEFAULT-NEXT: [[CONV:%.*]] = zext i1 [[CMP_I]] to i32
+// DEFAULT-NEXT: [[TMP0:%.*]] = fcmp uno double [[X:%.*]], 0.000000e+00
+// DEFAULT-NEXT: [[CONV:%.*]] = zext i1 [[TMP0]] to i32
// DEFAULT-NEXT: ret i32 [[CONV]]
//
// FINITEONLY-LABEL: @test___isnan(
@@ -1724,11 +1720,11 @@ extern "C" __device__ long int test_lround(double x) {
// DEFAULT-LABEL: @test_modff(
// DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5)
-// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18:[0-9]+]]
+// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17:[0-9]+]]
// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract float @__ocml_modf_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
// DEFAULT-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA16:![0-9]+]]
// DEFAULT-NEXT: store float [[TMP0]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA16]]
-// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
+// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_modff(
@@ -1748,11 +1744,11 @@ extern "C" __device__ float test_modff(float x, float* y) {
// DEFAULT-LABEL: @test_modf(
// DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5)
-// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
+// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract double @__ocml_modf_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
// DEFAULT-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA18:![0-9]+]]
// DEFAULT-NEXT: store double [[TMP0]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA18]]
-// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
+// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_modf(
@@ -2369,11 +2365,11 @@ extern "C" __device__ double test_remainder(double x, double y) {
// DEFAULT-LABEL: @test_remquof(
// DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca i32, align 4, addrspace(5)
-// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
+// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract float @__ocml_remquo_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
// DEFAULT-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA12]]
// DEFAULT-NEXT: store i32 [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA12]]
-// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
+// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
// DEFAULT-NEXT: ret float [[CALL_I]]
//
// FINITEONLY-LABEL: @test_remquof(
@@ -2393,11 +2389,11 @@ extern "C" __device__ float test_remquof(float x, float y, int* z) {
// DEFAULT-LABEL: @test_remquo(
// DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca i32, align 4, addrspace(5)
-// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
+// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract double @__ocml_remquo_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
// DEFAULT-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA12]]
// DEFAULT-NEXT: store i32 [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA12]]
-// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
+// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
// DEFAULT-NEXT: ret double [[CALL_I]]
//
// FINITEONLY-LABEL: @test_remquo(
@@ -2790,12 +2786,12 @@ extern "C" __device__ BOOL_TYPE test___signbit(double x) {
// DEFAULT-LABEL: @test_sincosf(
// DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5)
-// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
+// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract float @__ocml_sincos_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
// DEFAULT-NEXT: store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA16]]
// DEFAULT-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA16]]
// DEFAULT-NEXT: store float [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA16]]
-// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
+// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
// DEFAULT-NEXT: ret void
//
// FINITEONLY-LABEL: @test_sincosf(
@@ -2816,12 +2812,12 @@ extern "C" __device__ void test_sincosf(float x, float *y, float *z) {
// DEFAULT-LABEL: @test_sincos(
// DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5)
-// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
+// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract double @__ocml_sincos_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
// DEFAULT-NEXT: store double [[CALL_I]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA18]]
// DEFAULT-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA18]]
// DEFAULT-NEXT: store double [[TMP0]], ptr [[Z:%.*]], align 8, !tbaa [[TBAA18]]
-// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
+// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
// DEFAULT-NEXT: ret void
//
// FINITEONLY-LABEL: @test_sincos(
@@ -2842,12 +2838,12 @@ extern "C" __device__ void test_sincos(double x, double *y, double *z) {
// DEFAULT-LABEL: @test_sincospif(
// DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5)
-// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
+// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract float @__ocml_sincospi_f32(float noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
// DEFAULT-NEXT: store float [[CALL_I]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA16]]
// DEFAULT-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[__TMP_I]], align 4, !tbaa [[TBAA16]]
// DEFAULT-NEXT: store float [[TMP0]], ptr [[Z:%.*]], align 4, !tbaa [[TBAA16]]
-// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
+// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 4, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
// DEFAULT-NEXT: ret void
//
// FINITEONLY-LABEL: @test_sincospif(
@@ -2868,12 +2864,12 @@ extern "C" __device__ void test_sincospif(float x, float *y, float *z) {
// DEFAULT-LABEL: @test_sincospi(
// DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca double, align 8, addrspace(5)
-// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
+// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
// DEFAULT-NEXT: [[CALL_I:%.*]] = call contract double @__ocml_sincospi_f64(double noundef [[X:%.*]], ptr addrspace(5) noundef [[__TMP_I]]) #[[ATTR16]]
// DEFAULT-NEXT: store double [[CALL_I]], ptr [[Y:%.*]], align 8, !tbaa [[TBAA18]]
// DEFAULT-NEXT: [[TMP0:%.*]] = load double, ptr addrspace(5) [[__TMP_I]], align 8, !tbaa [[TBAA18]]
// DEFAULT-NEXT: store double [[TMP0]], ptr [[Z:%.*]], align 8, !tbaa [[TBAA18]]
-// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR18]]
+// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 8, ptr addrspace(5) [[__TMP_I]]) #[[ATTR17]]
// DEFAULT-NEXT: ret void
//
// FINITEONLY-LABEL: @test_sincospi(
diff --git a/clang/test/Headers/hip-header.hip b/clang/test/Headers/hip-header.hip
index 7d7b5e7176d9d0..b682bbba0c9d6e 100644
--- a/clang/test/Headers/hip-header.hip
+++ b/clang/test/Headers/hip-header.hip
@@ -120,12 +120,12 @@ __device__ double test_isnan() {
double d = 5.0;
float f = 5.0;
- // AMD_INT_RETURN: fcmp contract uno float
- // AMD_BOOL_RETURN: fcmp contract uno float
+ // AMD_INT_RETURN: call i1 @llvm.is.fpclass.f32(float {{.*}}, i32 3)
+ // AMD_BOOL_RETURN: call i1 @llvm.is.fpclass.f32(float {{.*}}, i32 3)
r += isnan(f);
- // AMD_INT_RETURN: fcmp contract uno double
- // AMD_BOOL_RETURN: fcmp contract uno double
+ // AMD_INT_RETURN: call i1 @llvm.is.fpclass.f64(double {{.*}}, i32 3)
+ // AMD_BOOL_RETURN: call i1 @llvm.is.fpclass.f64(double {{.*}}, i32 3)
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 d70370e50a48bf..0a632abcd8b06d 100644
--- a/clang/test/Headers/openmp_device_math_isnan.cpp
+++ b/clang/test/Headers/openmp_device_math_isnan.cpp
@@ -21,18 +21,18 @@
double math(float f, double d) {
double r = 0;
// INT_RETURN: call i32 @__nv_isnanf(float
- // AMD_INT_RETURN_SAFE: fcmp uno float
- // AMD_INT_RETURN_FAST: sitofp i32 0 to double
+ // AMD_INT_RETURN_SAFE: call i1 @llvm.is.fpclass.f32(float{{.*}}, i32 3)
+ // AMD_INT_RETURN_FAST: sitofp i32 {{.*}} to double
// BOOL_RETURN: call i32 @__nv_isnanf(float
- // AMD_BOOL_RETURN_SAFE: fcmp uno float
- // AMD_BOOL_RETURN_FAST: icmp ne i32 0, 0
+ // AMD_BOOL_RETURN_SAFE: call i1 @llvm.is.fpclass.f32(float{{.*}}, i32 3)
+ // AMD_BOOL_RETURN_FAST: icmp ne i32 {{.*}}, 0
r += std::isnan(f);
// INT_RETURN: call i32 @__nv_isnand(double
- // AMD_INT_RETURN_SAFE: fcmp uno double
- // AMD_INT_RETURN_FAST: sitofp i32 0 to double
+ // AMD_INT_RETURN_SAFE: call i1 @llvm.is.fpclass.f64(double{{.*}}, i32 3)
+ // AMD_INT_RETURN_FAST: sitofp i32 {{.*}} to double
// BOOL_RETURN: call i32 @__nv_isnand(double
- // AMD_BOOL_RETURN_SAFE: fcmp uno double
- // AMD_BOOL_RETURN_FAST: icmp ne i32 0, 0
+ // AMD_BOOL_RETURN_SAFE: call i1 @llvm.is.fpclass.f64(double{{.*}}, i32 3)
+ // AMD_BOOL_RETURN_FAST: icmp ne i32 {{.*}}, 0
r += std::isnan(d);
return r;
}
More information about the cfe-commits
mailing list