[clang] [llvm] Inliner: Handle nofpclass return attributes (PR #179776)
Matt Arsenault via llvm-commits
llvm-commits at lists.llvm.org
Thu Feb 5 00:26:23 PST 2026
https://github.com/arsenm updated https://github.com/llvm/llvm-project/pull/179776
>From 2831742578fe80c16fd6ee871a2efe255c7bffff Mon Sep 17 00:00:00 2001
From: Matt Arsenault <Matthew.Arsenault at amd.com>
Date: Wed, 4 Feb 2026 19:12:45 +0100
Subject: [PATCH 1/4] Inliner: Handle nofpclass return attributes
Follow along with how range is handled.
---
llvm/lib/Transforms/Utils/InlineFunction.cpp | 10 ++++
.../Transforms/Inline/ret_attr_nofpclass.ll | 59 +++++++++++++++++++
2 files changed, 69 insertions(+)
create mode 100644 llvm/test/Transforms/Inline/ret_attr_nofpclass.ll
diff --git a/llvm/lib/Transforms/Utils/InlineFunction.cpp b/llvm/lib/Transforms/Utils/InlineFunction.cpp
index aa902f687d8aa..3230b306f17d1 100644
--- a/llvm/lib/Transforms/Utils/InlineFunction.cpp
+++ b/llvm/lib/Transforms/Utils/InlineFunction.cpp
@@ -1543,6 +1543,8 @@ static AttrBuilder IdentifyValidPoisonGeneratingAttributes(CallBase &CB) {
Valid.addAlignmentAttr(CB.getRetAlign());
if (std::optional<ConstantRange> Range = CB.getRange())
Valid.addRangeAttr(*Range);
+ if (CB.hasRetAttr(Attribute::NoFPClass))
+ Valid.addNoFPClassAttr(CB.getRetNoFPClass());
return Valid;
}
@@ -1648,6 +1650,14 @@ static void AddReturnAttributes(CallBase &CB, ValueToValueMapTy &VMap,
CBRange.getRange().intersectWith(NewRange.getRange()));
}
}
+
+ Attribute CBNoFPClass = ValidPG.getAttribute(Attribute::NoFPClass);
+ if (CBNoFPClass.isValid() && AL.hasRetAttr(Attribute::NoFPClass)) {
+ ValidPG.addNoFPClassAttr(
+ CBNoFPClass.getNoFPClass() |
+ AL.getRetAttr(Attribute::NoFPClass).getNoFPClass());
+ }
+
// Three checks.
// If the callsite has `noundef`, then a poison due to violating the
// return attribute will create UB anyways so we can always propagate.
diff --git a/llvm/test/Transforms/Inline/ret_attr_nofpclass.ll b/llvm/test/Transforms/Inline/ret_attr_nofpclass.ll
new file mode 100644
index 0000000000000..08364f98f6a71
--- /dev/null
+++ b/llvm/test/Transforms/Inline/ret_attr_nofpclass.ll
@@ -0,0 +1,59 @@
+; NOTE: Assertions have been autogenerated by utils/update_test_checks.py UTC_ARGS: --version 6
+; RUN: opt -S -passes=inline %s | FileCheck %s
+; RUN: opt -S -passes='cgscc(inline)' %s | FileCheck %s
+; RUN: opt -S -passes='module-inline' %s | FileCheck %s
+
+declare float @val_f32()
+declare float @use.val(float) willreturn nounwind
+
+define float @callee() {
+; CHECK-LABEL: define float @callee() {
+; CHECK-NEXT: [[R:%.*]] = call float @val_f32()
+; CHECK-NEXT: ret float [[R]]
+;
+ %r = call float @val_f32()
+ ret float %r
+}
+
+define float @caller_okay_use_after_poison_anyways() {
+; CHECK-LABEL: define float @caller_okay_use_after_poison_anyways() {
+; CHECK-NEXT: [[R_I:%.*]] = call nofpclass(nan) float @val_f32()
+; CHECK-NEXT: call void @use.val(float [[R_I]])
+; CHECK-NEXT: ret float [[R_I]]
+;
+ %r = call nofpclass(nan) float @callee()
+ call void @use.val(float %r)
+ ret float %r
+}
+
+define float @callee_nofpclass_inf_nan() {
+; CHECK-LABEL: define float @callee_nofpclass_inf_nan() {
+; CHECK-NEXT: [[R:%.*]] = call nofpclass(nan inf) float @val_f32()
+; CHECK-NEXT: ret float [[R]]
+;
+ %r = call nofpclass(nan inf) float @val_f32()
+ ret float %r
+}
+
+define float @caller_okay_intersect_nofpclass() {
+; CHECK-LABEL: define float @caller_okay_intersect_nofpclass() {
+; CHECK-NEXT: [[R_I:%.*]] = call nofpclass(nan) float @val_f32()
+; CHECK-NEXT: call void @use.val(float [[R_I]])
+; CHECK-NEXT: ret float [[R_I]]
+;
+ %r = call nofpclass(nan) float @callee_nofpclass_inf_nan()
+ call void @use.val(float %r)
+ ret float %r
+}
+
+define float @caller_not_intersecting_nofpclass() {
+; CHECK-LABEL: define float @caller_not_intersecting_nofpclass() {
+; CHECK-NEXT: [[R_I:%.*]] = call nofpclass(zero) float @val_f32()
+; CHECK-NEXT: call void @use.val(float [[R_I]])
+; CHECK-NEXT: ret float [[R_I]]
+;
+ %r = call nofpclass(zero) float @callee_nofpclass_inf_nan()
+ call void @use.val(float %r)
+ ret float %r
+}
+
>From 8e2feccd65d6a564766d994c7e9fc45e23f279aa Mon Sep 17 00:00:00 2001
From: Matt Arsenault <Matthew.Arsenault at amd.com>
Date: Thu, 5 Feb 2026 07:39:59 +0100
Subject: [PATCH 2/4] Generate checks
---
llvm/test/Transforms/Inline/ret_attr_nofpclass.ll | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/llvm/test/Transforms/Inline/ret_attr_nofpclass.ll b/llvm/test/Transforms/Inline/ret_attr_nofpclass.ll
index 08364f98f6a71..38271f3d8eb30 100644
--- a/llvm/test/Transforms/Inline/ret_attr_nofpclass.ll
+++ b/llvm/test/Transforms/Inline/ret_attr_nofpclass.ll
@@ -37,7 +37,7 @@ define float @callee_nofpclass_inf_nan() {
define float @caller_okay_intersect_nofpclass() {
; CHECK-LABEL: define float @caller_okay_intersect_nofpclass() {
-; CHECK-NEXT: [[R_I:%.*]] = call nofpclass(nan) float @val_f32()
+; CHECK-NEXT: [[R_I:%.*]] = call nofpclass(nan inf) float @val_f32()
; CHECK-NEXT: call void @use.val(float [[R_I]])
; CHECK-NEXT: ret float [[R_I]]
;
@@ -48,7 +48,7 @@ define float @caller_okay_intersect_nofpclass() {
define float @caller_not_intersecting_nofpclass() {
; CHECK-LABEL: define float @caller_not_intersecting_nofpclass() {
-; CHECK-NEXT: [[R_I:%.*]] = call nofpclass(zero) float @val_f32()
+; CHECK-NEXT: [[R_I:%.*]] = call nofpclass(nan inf zero) float @val_f32()
; CHECK-NEXT: call void @use.val(float [[R_I]])
; CHECK-NEXT: ret float [[R_I]]
;
>From aba968e25609b17d7f0359088f588b0a399189ba Mon Sep 17 00:00:00 2001
From: Matt Arsenault <Matthew.Arsenault at amd.com>
Date: Thu, 5 Feb 2026 07:41:33 +0100
Subject: [PATCH 3/4] Remove extra run lines
---
llvm/test/Transforms/Inline/ret_attr_nofpclass.ll | 2 --
1 file changed, 2 deletions(-)
diff --git a/llvm/test/Transforms/Inline/ret_attr_nofpclass.ll b/llvm/test/Transforms/Inline/ret_attr_nofpclass.ll
index 38271f3d8eb30..16ad718d60349 100644
--- a/llvm/test/Transforms/Inline/ret_attr_nofpclass.ll
+++ b/llvm/test/Transforms/Inline/ret_attr_nofpclass.ll
@@ -1,7 +1,5 @@
; NOTE: Assertions have been autogenerated by utils/update_test_checks.py UTC_ARGS: --version 6
; RUN: opt -S -passes=inline %s | FileCheck %s
-; RUN: opt -S -passes='cgscc(inline)' %s | FileCheck %s
-; RUN: opt -S -passes='module-inline' %s | FileCheck %s
declare float @val_f32()
declare float @use.val(float) willreturn nounwind
>From 6315bb3362356b3f83b615d86e21306baf4a17b6 Mon Sep 17 00:00:00 2001
From: Matt Arsenault <Matthew.Arsenault at amd.com>
Date: Thu, 5 Feb 2026 09:19:16 +0100
Subject: [PATCH 4/4] Update clang tests
---
clang/test/CodeGenHLSL/builtins/distance.hlsl | 32 +++---
clang/test/CodeGenHLSL/builtins/length.hlsl | 32 +++---
.../CodeGenHLSL/builtins/max-overloads.hlsl | 10 +-
.../CodeGenHLSL/builtins/min-overloads.hlsl | 10 +-
.../CodeGenHLSL/builtins/pow-overloads.hlsl | 40 +++----
clang/test/CodeGenHLSL/builtins/reflect.hlsl | 12 +-
clang/test/CodeGenHLSL/builtins/refract.hlsl | 18 +--
.../test/CodeGenHLSL/builtins/smoothstep.hlsl | 16 +--
clang/test/Headers/__clang_hip_cmath.hip | 10 +-
clang/test/Headers/__clang_hip_math.hip | 104 +++++++++---------
10 files changed, 142 insertions(+), 142 deletions(-)
diff --git a/clang/test/CodeGenHLSL/builtins/distance.hlsl b/clang/test/CodeGenHLSL/builtins/distance.hlsl
index bf015415a7d2f..efe546bfdc5cf 100644
--- a/clang/test/CodeGenHLSL/builtins/distance.hlsl
+++ b/clang/test/CodeGenHLSL/builtins/distance.hlsl
@@ -10,14 +10,14 @@
// CHECK-SAME: half noundef nofpclass(nan inf) [[X:%.*]], half noundef nofpclass(nan inf) [[Y:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {
// CHECK-NEXT: [[ENTRY:.*:]]
// CHECK-NEXT: [[SUB_I:%.*]] = fsub reassoc nnan ninf nsz arcp afn half [[X]], [[Y]]
-// CHECK-NEXT: [[ELT_ABS_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef half @llvm.fabs.f16(half nofpclass(nan inf) [[SUB_I]])
+// CHECK-NEXT: [[ELT_ABS_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) half @llvm.fabs.f16(half nofpclass(nan inf) [[SUB_I]])
// CHECK-NEXT: ret half [[ELT_ABS_I]]
//
// SPVCHECK-LABEL: define hidden spir_func noundef nofpclass(nan inf) half @_Z18test_distance_halfDhDh(
// SPVCHECK-SAME: half noundef nofpclass(nan inf) [[X:%.*]], half noundef nofpclass(nan inf) [[Y:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {
// SPVCHECK-NEXT: [[ENTRY:.*:]]
// SPVCHECK-NEXT: [[SUB_I:%.*]] = fsub reassoc nnan ninf nsz arcp afn half [[X]], [[Y]]
-// SPVCHECK-NEXT: [[ELT_ABS_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef half @llvm.fabs.f16(half nofpclass(nan inf) [[SUB_I]])
+// SPVCHECK-NEXT: [[ELT_ABS_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) half @llvm.fabs.f16(half nofpclass(nan inf) [[SUB_I]])
// SPVCHECK-NEXT: ret half [[ELT_ABS_I]]
//
half test_distance_half(half X, half Y) { return distance(X, Y); }
@@ -27,14 +27,14 @@ half test_distance_half(half X, half Y) { return distance(X, Y); }
// CHECK-NEXT: [[ENTRY:.*:]]
// CHECK-NEXT: [[SUB_I:%.*]] = fsub reassoc nnan ninf nsz arcp afn <2 x half> [[X]], [[Y]]
// CHECK-NEXT: [[HLSL_DOT_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn half @llvm.dx.fdot.v2f16(<2 x half> nofpclass(nan inf) [[SUB_I]], <2 x half> nofpclass(nan inf) [[SUB_I]])
-// CHECK-NEXT: [[TMP0:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef half @llvm.sqrt.f16(half [[HLSL_DOT_I]])
+// CHECK-NEXT: [[TMP0:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) half @llvm.sqrt.f16(half [[HLSL_DOT_I]])
// CHECK-NEXT: ret half [[TMP0]]
//
// SPVCHECK-LABEL: define hidden spir_func noundef nofpclass(nan inf) half @_Z19test_distance_half2Dv2_DhS_(
// SPVCHECK-SAME: <2 x half> noundef nofpclass(nan inf) [[X:%.*]], <2 x half> noundef nofpclass(nan inf) [[Y:%.*]]) local_unnamed_addr #[[ATTR0]] {
// SPVCHECK-NEXT: [[ENTRY:.*:]]
// SPVCHECK-NEXT: [[SUB_I:%.*]] = fsub reassoc nnan ninf nsz arcp afn <2 x half> [[X]], [[Y]]
-// SPVCHECK-NEXT: [[SPV_LENGTH_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef half @llvm.spv.length.v2f16(<2 x half> nofpclass(nan inf) [[SUB_I]])
+// SPVCHECK-NEXT: [[SPV_LENGTH_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) half @llvm.spv.length.v2f16(<2 x half> nofpclass(nan inf) [[SUB_I]])
// SPVCHECK-NEXT: ret half [[SPV_LENGTH_I]]
//
half test_distance_half2(half2 X, half2 Y) { return distance(X, Y); }
@@ -44,14 +44,14 @@ half test_distance_half2(half2 X, half2 Y) { return distance(X, Y); }
// CHECK-NEXT: [[ENTRY:.*:]]
// CHECK-NEXT: [[SUB_I:%.*]] = fsub reassoc nnan ninf nsz arcp afn <3 x half> [[X]], [[Y]]
// CHECK-NEXT: [[HLSL_DOT_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn half @llvm.dx.fdot.v3f16(<3 x half> nofpclass(nan inf) [[SUB_I]], <3 x half> nofpclass(nan inf) [[SUB_I]])
-// CHECK-NEXT: [[TMP0:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef half @llvm.sqrt.f16(half [[HLSL_DOT_I]])
+// CHECK-NEXT: [[TMP0:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) half @llvm.sqrt.f16(half [[HLSL_DOT_I]])
// CHECK-NEXT: ret half [[TMP0]]
//
// SPVCHECK-LABEL: define hidden spir_func noundef nofpclass(nan inf) half @_Z19test_distance_half3Dv3_DhS_(
// SPVCHECK-SAME: <3 x half> noundef nofpclass(nan inf) [[X:%.*]], <3 x half> noundef nofpclass(nan inf) [[Y:%.*]]) local_unnamed_addr #[[ATTR0]] {
// SPVCHECK-NEXT: [[ENTRY:.*:]]
// SPVCHECK-NEXT: [[SUB_I:%.*]] = fsub reassoc nnan ninf nsz arcp afn <3 x half> [[X]], [[Y]]
-// SPVCHECK-NEXT: [[SPV_LENGTH_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef half @llvm.spv.length.v3f16(<3 x half> nofpclass(nan inf) [[SUB_I]])
+// SPVCHECK-NEXT: [[SPV_LENGTH_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) half @llvm.spv.length.v3f16(<3 x half> nofpclass(nan inf) [[SUB_I]])
// SPVCHECK-NEXT: ret half [[SPV_LENGTH_I]]
//
half test_distance_half3(half3 X, half3 Y) { return distance(X, Y); }
@@ -61,14 +61,14 @@ half test_distance_half3(half3 X, half3 Y) { return distance(X, Y); }
// CHECK-NEXT: [[ENTRY:.*:]]
// CHECK-NEXT: [[SUB_I:%.*]] = fsub reassoc nnan ninf nsz arcp afn <4 x half> [[X]], [[Y]]
// CHECK-NEXT: [[HLSL_DOT_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn half @llvm.dx.fdot.v4f16(<4 x half> nofpclass(nan inf) [[SUB_I]], <4 x half> nofpclass(nan inf) [[SUB_I]])
-// CHECK-NEXT: [[TMP0:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef half @llvm.sqrt.f16(half [[HLSL_DOT_I]])
+// CHECK-NEXT: [[TMP0:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) half @llvm.sqrt.f16(half [[HLSL_DOT_I]])
// CHECK-NEXT: ret half [[TMP0]]
//
// SPVCHECK-LABEL: define hidden spir_func noundef nofpclass(nan inf) half @_Z19test_distance_half4Dv4_DhS_(
// SPVCHECK-SAME: <4 x half> noundef nofpclass(nan inf) [[X:%.*]], <4 x half> noundef nofpclass(nan inf) [[Y:%.*]]) local_unnamed_addr #[[ATTR0]] {
// SPVCHECK-NEXT: [[ENTRY:.*:]]
// SPVCHECK-NEXT: [[SUB_I:%.*]] = fsub reassoc nnan ninf nsz arcp afn <4 x half> [[X]], [[Y]]
-// SPVCHECK-NEXT: [[SPV_LENGTH_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef half @llvm.spv.length.v4f16(<4 x half> nofpclass(nan inf) [[SUB_I]])
+// SPVCHECK-NEXT: [[SPV_LENGTH_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) half @llvm.spv.length.v4f16(<4 x half> nofpclass(nan inf) [[SUB_I]])
// SPVCHECK-NEXT: ret half [[SPV_LENGTH_I]]
//
half test_distance_half4(half4 X, half4 Y) { return distance(X, Y); }
@@ -77,14 +77,14 @@ half test_distance_half4(half4 X, half4 Y) { return distance(X, Y); }
// CHECK-SAME: float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) local_unnamed_addr #[[ATTR0]] {
// CHECK-NEXT: [[ENTRY:.*:]]
// CHECK-NEXT: [[SUB_I:%.*]] = fsub reassoc nnan ninf nsz arcp afn float [[X]], [[Y]]
-// CHECK-NEXT: [[ELT_ABS_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef float @llvm.fabs.f32(float nofpclass(nan inf) [[SUB_I]])
+// CHECK-NEXT: [[ELT_ABS_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) float @llvm.fabs.f32(float nofpclass(nan inf) [[SUB_I]])
// CHECK-NEXT: ret float [[ELT_ABS_I]]
//
// SPVCHECK-LABEL: define hidden spir_func noundef nofpclass(nan inf) float @_Z19test_distance_floatff(
// SPVCHECK-SAME: float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) local_unnamed_addr #[[ATTR0]] {
// SPVCHECK-NEXT: [[ENTRY:.*:]]
// SPVCHECK-NEXT: [[SUB_I:%.*]] = fsub reassoc nnan ninf nsz arcp afn float [[X]], [[Y]]
-// SPVCHECK-NEXT: [[ELT_ABS_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef float @llvm.fabs.f32(float nofpclass(nan inf) [[SUB_I]])
+// SPVCHECK-NEXT: [[ELT_ABS_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) float @llvm.fabs.f32(float nofpclass(nan inf) [[SUB_I]])
// SPVCHECK-NEXT: ret float [[ELT_ABS_I]]
//
float test_distance_float(float X, float Y) { return distance(X, Y); }
@@ -94,14 +94,14 @@ float test_distance_float(float X, float Y) { return distance(X, Y); }
// CHECK-NEXT: [[ENTRY:.*:]]
// CHECK-NEXT: [[SUB_I:%.*]] = fsub reassoc nnan ninf nsz arcp afn <2 x float> [[X]], [[Y]]
// CHECK-NEXT: [[HLSL_DOT_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn float @llvm.dx.fdot.v2f32(<2 x float> nofpclass(nan inf) [[SUB_I]], <2 x float> nofpclass(nan inf) [[SUB_I]])
-// CHECK-NEXT: [[TMP0:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef float @llvm.sqrt.f32(float [[HLSL_DOT_I]])
+// CHECK-NEXT: [[TMP0:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) float @llvm.sqrt.f32(float [[HLSL_DOT_I]])
// CHECK-NEXT: ret float [[TMP0]]
//
// SPVCHECK-LABEL: define hidden spir_func noundef nofpclass(nan inf) float @_Z20test_distance_float2Dv2_fS_(
// SPVCHECK-SAME: <2 x float> noundef nofpclass(nan inf) [[X:%.*]], <2 x float> noundef nofpclass(nan inf) [[Y:%.*]]) local_unnamed_addr #[[ATTR0]] {
// SPVCHECK-NEXT: [[ENTRY:.*:]]
// SPVCHECK-NEXT: [[SUB_I:%.*]] = fsub reassoc nnan ninf nsz arcp afn <2 x float> [[X]], [[Y]]
-// SPVCHECK-NEXT: [[SPV_LENGTH_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef float @llvm.spv.length.v2f32(<2 x float> nofpclass(nan inf) [[SUB_I]])
+// SPVCHECK-NEXT: [[SPV_LENGTH_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) float @llvm.spv.length.v2f32(<2 x float> nofpclass(nan inf) [[SUB_I]])
// SPVCHECK-NEXT: ret float [[SPV_LENGTH_I]]
//
float test_distance_float2(float2 X, float2 Y) { return distance(X, Y); }
@@ -111,14 +111,14 @@ float test_distance_float2(float2 X, float2 Y) { return distance(X, Y); }
// CHECK-NEXT: [[ENTRY:.*:]]
// CHECK-NEXT: [[SUB_I:%.*]] = fsub reassoc nnan ninf nsz arcp afn <3 x float> [[X]], [[Y]]
// CHECK-NEXT: [[HLSL_DOT_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn float @llvm.dx.fdot.v3f32(<3 x float> nofpclass(nan inf) [[SUB_I]], <3 x float> nofpclass(nan inf) [[SUB_I]])
-// CHECK-NEXT: [[TMP0:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef float @llvm.sqrt.f32(float [[HLSL_DOT_I]])
+// CHECK-NEXT: [[TMP0:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) float @llvm.sqrt.f32(float [[HLSL_DOT_I]])
// CHECK-NEXT: ret float [[TMP0]]
//
// SPVCHECK-LABEL: define hidden spir_func noundef nofpclass(nan inf) float @_Z20test_distance_float3Dv3_fS_(
// SPVCHECK-SAME: <3 x float> noundef nofpclass(nan inf) [[X:%.*]], <3 x float> noundef nofpclass(nan inf) [[Y:%.*]]) local_unnamed_addr #[[ATTR0]] {
// SPVCHECK-NEXT: [[ENTRY:.*:]]
// SPVCHECK-NEXT: [[SUB_I:%.*]] = fsub reassoc nnan ninf nsz arcp afn <3 x float> [[X]], [[Y]]
-// SPVCHECK-NEXT: [[SPV_LENGTH_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef float @llvm.spv.length.v3f32(<3 x float> nofpclass(nan inf) [[SUB_I]])
+// SPVCHECK-NEXT: [[SPV_LENGTH_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) float @llvm.spv.length.v3f32(<3 x float> nofpclass(nan inf) [[SUB_I]])
// SPVCHECK-NEXT: ret float [[SPV_LENGTH_I]]
//
float test_distance_float3(float3 X, float3 Y) { return distance(X, Y); }
@@ -128,14 +128,14 @@ float test_distance_float3(float3 X, float3 Y) { return distance(X, Y); }
// CHECK-NEXT: [[ENTRY:.*:]]
// CHECK-NEXT: [[SUB_I:%.*]] = fsub reassoc nnan ninf nsz arcp afn <4 x float> [[X]], [[Y]]
// CHECK-NEXT: [[HLSL_DOT_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn float @llvm.dx.fdot.v4f32(<4 x float> nofpclass(nan inf) [[SUB_I]], <4 x float> nofpclass(nan inf) [[SUB_I]])
-// CHECK-NEXT: [[TMP0:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef float @llvm.sqrt.f32(float [[HLSL_DOT_I]])
+// CHECK-NEXT: [[TMP0:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) float @llvm.sqrt.f32(float [[HLSL_DOT_I]])
// CHECK-NEXT: ret float [[TMP0]]
//
// SPVCHECK-LABEL: define hidden spir_func noundef nofpclass(nan inf) float @_Z20test_distance_float4Dv4_fS_(
// SPVCHECK-SAME: <4 x float> noundef nofpclass(nan inf) [[X:%.*]], <4 x float> noundef nofpclass(nan inf) [[Y:%.*]]) local_unnamed_addr #[[ATTR0]] {
// SPVCHECK-NEXT: [[ENTRY:.*:]]
// SPVCHECK-NEXT: [[SUB_I:%.*]] = fsub reassoc nnan ninf nsz arcp afn <4 x float> [[X]], [[Y]]
-// SPVCHECK-NEXT: [[SPV_LENGTH_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef float @llvm.spv.length.v4f32(<4 x float> nofpclass(nan inf) [[SUB_I]])
+// SPVCHECK-NEXT: [[SPV_LENGTH_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) float @llvm.spv.length.v4f32(<4 x float> nofpclass(nan inf) [[SUB_I]])
// SPVCHECK-NEXT: ret float [[SPV_LENGTH_I]]
//
float test_distance_float4(float4 X, float4 Y) { return distance(X, Y); }
diff --git a/clang/test/CodeGenHLSL/builtins/length.hlsl b/clang/test/CodeGenHLSL/builtins/length.hlsl
index 95edb20dacdac..05f3fa4233ff5 100644
--- a/clang/test/CodeGenHLSL/builtins/length.hlsl
+++ b/clang/test/CodeGenHLSL/builtins/length.hlsl
@@ -11,13 +11,13 @@
// CHECK-LABEL: define hidden noundef nofpclass(nan inf) half @_Z16test_length_halfDh(
// CHECK-SAME: half noundef nofpclass(nan inf) [[P0:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {
// CHECK-NEXT: [[ENTRY:.*:]]
-// CHECK-NEXT: [[ELT_ABS_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef half @llvm.fabs.f16(half nofpclass(nan inf) [[P0]])
+// CHECK-NEXT: [[ELT_ABS_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) half @llvm.fabs.f16(half nofpclass(nan inf) [[P0]])
// CHECK-NEXT: ret half [[ELT_ABS_I]]
//
// SPVCHECK-LABEL: define hidden spir_func noundef nofpclass(nan inf) half @_Z16test_length_halfDh(
// SPVCHECK-SAME: half noundef nofpclass(nan inf) [[P0:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {
// SPVCHECK-NEXT: [[ENTRY:.*:]]
-// SPVCHECK-NEXT: [[ELT_ABS_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef half @llvm.fabs.f16(half nofpclass(nan inf) [[P0]])
+// SPVCHECK-NEXT: [[ELT_ABS_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) half @llvm.fabs.f16(half nofpclass(nan inf) [[P0]])
// SPVCHECK-NEXT: ret half [[ELT_ABS_I]]
//
half test_length_half(half p0)
@@ -29,13 +29,13 @@ half test_length_half(half p0)
// CHECK-SAME: <2 x half> noundef nofpclass(nan inf) [[P0:%.*]]) local_unnamed_addr #[[ATTR0]] {
// CHECK-NEXT: [[ENTRY:.*:]]
// CHECK-NEXT: [[HLSL_DOT_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn half @llvm.dx.fdot.v2f16(<2 x half> nofpclass(nan inf) [[P0]], <2 x half> nofpclass(nan inf) [[P0]])
-// CHECK-NEXT: [[TMP0:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef half @llvm.sqrt.f16(half [[HLSL_DOT_I]])
+// CHECK-NEXT: [[TMP0:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) half @llvm.sqrt.f16(half [[HLSL_DOT_I]])
// CHECK-NEXT: ret half [[TMP0]]
//
// SPVCHECK-LABEL: define hidden spir_func noundef nofpclass(nan inf) half @_Z17test_length_half2Dv2_Dh(
// SPVCHECK-SAME: <2 x half> noundef nofpclass(nan inf) [[P0:%.*]]) local_unnamed_addr #[[ATTR0]] {
// SPVCHECK-NEXT: [[ENTRY:.*:]]
-// SPVCHECK-NEXT: [[SPV_LENGTH_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef half @llvm.spv.length.v2f16(<2 x half> nofpclass(nan inf) [[P0]])
+// SPVCHECK-NEXT: [[SPV_LENGTH_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) half @llvm.spv.length.v2f16(<2 x half> nofpclass(nan inf) [[P0]])
// SPVCHECK-NEXT: ret half [[SPV_LENGTH_I]]
//
half test_length_half2(half2 p0)
@@ -47,13 +47,13 @@ half test_length_half2(half2 p0)
// CHECK-SAME: <3 x half> noundef nofpclass(nan inf) [[P0:%.*]]) local_unnamed_addr #[[ATTR0]] {
// CHECK-NEXT: [[ENTRY:.*:]]
// CHECK-NEXT: [[HLSL_DOT_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn half @llvm.dx.fdot.v3f16(<3 x half> nofpclass(nan inf) [[P0]], <3 x half> nofpclass(nan inf) [[P0]])
-// CHECK-NEXT: [[TMP0:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef half @llvm.sqrt.f16(half [[HLSL_DOT_I]])
+// CHECK-NEXT: [[TMP0:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) half @llvm.sqrt.f16(half [[HLSL_DOT_I]])
// CHECK-NEXT: ret half [[TMP0]]
//
// SPVCHECK-LABEL: define hidden spir_func noundef nofpclass(nan inf) half @_Z17test_length_half3Dv3_Dh(
// SPVCHECK-SAME: <3 x half> noundef nofpclass(nan inf) [[P0:%.*]]) local_unnamed_addr #[[ATTR0]] {
// SPVCHECK-NEXT: [[ENTRY:.*:]]
-// SPVCHECK-NEXT: [[SPV_LENGTH_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef half @llvm.spv.length.v3f16(<3 x half> nofpclass(nan inf) [[P0]])
+// SPVCHECK-NEXT: [[SPV_LENGTH_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) half @llvm.spv.length.v3f16(<3 x half> nofpclass(nan inf) [[P0]])
// SPVCHECK-NEXT: ret half [[SPV_LENGTH_I]]
//
half test_length_half3(half3 p0)
@@ -65,13 +65,13 @@ half test_length_half3(half3 p0)
// CHECK-SAME: <4 x half> noundef nofpclass(nan inf) [[P0:%.*]]) local_unnamed_addr #[[ATTR0]] {
// CHECK-NEXT: [[ENTRY:.*:]]
// CHECK-NEXT: [[HLSL_DOT_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn half @llvm.dx.fdot.v4f16(<4 x half> nofpclass(nan inf) [[P0]], <4 x half> nofpclass(nan inf) [[P0]])
-// CHECK-NEXT: [[TMP0:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef half @llvm.sqrt.f16(half [[HLSL_DOT_I]])
+// CHECK-NEXT: [[TMP0:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) half @llvm.sqrt.f16(half [[HLSL_DOT_I]])
// CHECK-NEXT: ret half [[TMP0]]
//
// SPVCHECK-LABEL: define hidden spir_func noundef nofpclass(nan inf) half @_Z17test_length_half4Dv4_Dh(
// SPVCHECK-SAME: <4 x half> noundef nofpclass(nan inf) [[P0:%.*]]) local_unnamed_addr #[[ATTR0]] {
// SPVCHECK-NEXT: [[ENTRY:.*:]]
-// SPVCHECK-NEXT: [[SPV_LENGTH_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef half @llvm.spv.length.v4f16(<4 x half> nofpclass(nan inf) [[P0]])
+// SPVCHECK-NEXT: [[SPV_LENGTH_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) half @llvm.spv.length.v4f16(<4 x half> nofpclass(nan inf) [[P0]])
// SPVCHECK-NEXT: ret half [[SPV_LENGTH_I]]
//
half test_length_half4(half4 p0)
@@ -82,13 +82,13 @@ half test_length_half4(half4 p0)
// CHECK-LABEL: define hidden noundef nofpclass(nan inf) float @_Z17test_length_floatf(
// CHECK-SAME: float noundef nofpclass(nan inf) [[P0:%.*]]) local_unnamed_addr #[[ATTR0]] {
// CHECK-NEXT: [[ENTRY:.*:]]
-// CHECK-NEXT: [[ELT_ABS_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef float @llvm.fabs.f32(float nofpclass(nan inf) [[P0]])
+// CHECK-NEXT: [[ELT_ABS_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) float @llvm.fabs.f32(float nofpclass(nan inf) [[P0]])
// CHECK-NEXT: ret float [[ELT_ABS_I]]
//
// SPVCHECK-LABEL: define hidden spir_func noundef nofpclass(nan inf) float @_Z17test_length_floatf(
// SPVCHECK-SAME: float noundef nofpclass(nan inf) [[P0:%.*]]) local_unnamed_addr #[[ATTR0]] {
// SPVCHECK-NEXT: [[ENTRY:.*:]]
-// SPVCHECK-NEXT: [[ELT_ABS_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef float @llvm.fabs.f32(float nofpclass(nan inf) [[P0]])
+// SPVCHECK-NEXT: [[ELT_ABS_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) float @llvm.fabs.f32(float nofpclass(nan inf) [[P0]])
// SPVCHECK-NEXT: ret float [[ELT_ABS_I]]
//
float test_length_float(float p0)
@@ -100,13 +100,13 @@ float test_length_float(float p0)
// CHECK-SAME: <2 x float> noundef nofpclass(nan inf) [[P0:%.*]]) local_unnamed_addr #[[ATTR0]] {
// CHECK-NEXT: [[ENTRY:.*:]]
// CHECK-NEXT: [[HLSL_DOT_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn float @llvm.dx.fdot.v2f32(<2 x float> nofpclass(nan inf) [[P0]], <2 x float> nofpclass(nan inf) [[P0]])
-// CHECK-NEXT: [[TMP0:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef float @llvm.sqrt.f32(float [[HLSL_DOT_I]])
+// CHECK-NEXT: [[TMP0:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) float @llvm.sqrt.f32(float [[HLSL_DOT_I]])
// CHECK-NEXT: ret float [[TMP0]]
//
// SPVCHECK-LABEL: define hidden spir_func noundef nofpclass(nan inf) float @_Z18test_length_float2Dv2_f(
// SPVCHECK-SAME: <2 x float> noundef nofpclass(nan inf) [[P0:%.*]]) local_unnamed_addr #[[ATTR0]] {
// SPVCHECK-NEXT: [[ENTRY:.*:]]
-// SPVCHECK-NEXT: [[SPV_LENGTH_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef float @llvm.spv.length.v2f32(<2 x float> nofpclass(nan inf) [[P0]])
+// SPVCHECK-NEXT: [[SPV_LENGTH_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) float @llvm.spv.length.v2f32(<2 x float> nofpclass(nan inf) [[P0]])
// SPVCHECK-NEXT: ret float [[SPV_LENGTH_I]]
//
float test_length_float2(float2 p0)
@@ -118,13 +118,13 @@ float test_length_float2(float2 p0)
// CHECK-SAME: <3 x float> noundef nofpclass(nan inf) [[P0:%.*]]) local_unnamed_addr #[[ATTR0]] {
// CHECK-NEXT: [[ENTRY:.*:]]
// CHECK-NEXT: [[HLSL_DOT_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn float @llvm.dx.fdot.v3f32(<3 x float> nofpclass(nan inf) [[P0]], <3 x float> nofpclass(nan inf) [[P0]])
-// CHECK-NEXT: [[TMP0:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef float @llvm.sqrt.f32(float [[HLSL_DOT_I]])
+// CHECK-NEXT: [[TMP0:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) float @llvm.sqrt.f32(float [[HLSL_DOT_I]])
// CHECK-NEXT: ret float [[TMP0]]
//
// SPVCHECK-LABEL: define hidden spir_func noundef nofpclass(nan inf) float @_Z18test_length_float3Dv3_f(
// SPVCHECK-SAME: <3 x float> noundef nofpclass(nan inf) [[P0:%.*]]) local_unnamed_addr #[[ATTR0]] {
// SPVCHECK-NEXT: [[ENTRY:.*:]]
-// SPVCHECK-NEXT: [[SPV_LENGTH_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef float @llvm.spv.length.v3f32(<3 x float> nofpclass(nan inf) [[P0]])
+// SPVCHECK-NEXT: [[SPV_LENGTH_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) float @llvm.spv.length.v3f32(<3 x float> nofpclass(nan inf) [[P0]])
// SPVCHECK-NEXT: ret float [[SPV_LENGTH_I]]
//
float test_length_float3(float3 p0)
@@ -136,13 +136,13 @@ float test_length_float3(float3 p0)
// CHECK-SAME: <4 x float> noundef nofpclass(nan inf) [[P0:%.*]]) local_unnamed_addr #[[ATTR0]] {
// CHECK-NEXT: [[ENTRY:.*:]]
// CHECK-NEXT: [[HLSL_DOT_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn float @llvm.dx.fdot.v4f32(<4 x float> nofpclass(nan inf) [[P0]], <4 x float> nofpclass(nan inf) [[P0]])
-// CHECK-NEXT: [[TMP0:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef float @llvm.sqrt.f32(float [[HLSL_DOT_I]])
+// CHECK-NEXT: [[TMP0:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) float @llvm.sqrt.f32(float [[HLSL_DOT_I]])
// CHECK-NEXT: ret float [[TMP0]]
//
// SPVCHECK-LABEL: define hidden spir_func noundef nofpclass(nan inf) float @_Z18test_length_float4Dv4_f(
// SPVCHECK-SAME: <4 x float> noundef nofpclass(nan inf) [[P0:%.*]]) local_unnamed_addr #[[ATTR0]] {
// SPVCHECK-NEXT: [[ENTRY:.*:]]
-// SPVCHECK-NEXT: [[SPV_LENGTH_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef float @llvm.spv.length.v4f32(<4 x float> nofpclass(nan inf) [[P0]])
+// SPVCHECK-NEXT: [[SPV_LENGTH_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) float @llvm.spv.length.v4f32(<4 x float> nofpclass(nan inf) [[P0]])
// SPVCHECK-NEXT: ret float [[SPV_LENGTH_I]]
//
float test_length_float4(float4 p0)
diff --git a/clang/test/CodeGenHLSL/builtins/max-overloads.hlsl b/clang/test/CodeGenHLSL/builtins/max-overloads.hlsl
index a5ef87a822dd5..da5cd8ff37510 100644
--- a/clang/test/CodeGenHLSL/builtins/max-overloads.hlsl
+++ b/clang/test/CodeGenHLSL/builtins/max-overloads.hlsl
@@ -50,32 +50,32 @@ uint64_t4 test_max_ulong4_mismatch(uint64_t4 p0, uint64_t p1) { return max(p0, p
// NATIVE_HALF-LABEL: define hidden noundef nofpclass(nan inf) <4 x half> {{.*}}test_max_half4_mismatch
// NATIVE_HALF: [[CONV0:%.*]] = insertelement <4 x half> poison, half %{{.*}}, i64 0
// NATIVE_HALF: [[CONV1:%.*]] = shufflevector <4 x half> [[CONV0]], <4 x half> poison, <4 x i32> zeroinitializer
-// NATIVE_HALF: [[MAX:%.*]] = call reassoc nnan ninf nsz arcp afn noundef <4 x half> @llvm.maxnum.v4f16(<4 x half> %{{.*}}, <4 x half> [[CONV1]])
+// NATIVE_HALF: [[MAX:%.*]] = call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) <4 x half> @llvm.maxnum.v4f16(<4 x half> %{{.*}}, <4 x half> [[CONV1]])
// NATIVE_HALF: ret <4 x half> [[MAX]]
// NO_HALF-LABEL: define hidden noundef nofpclass(nan inf) <4 x float> {{.*}}test_max_half4_mismatch
// NO_HALF: [[CONV0:%.*]] = insertelement <4 x float> poison, float %{{.*}}, i64 0
// NO_HALF: [[CONV1:%.*]] = shufflevector <4 x float> [[CONV0]], <4 x float> poison, <4 x i32> zeroinitializer
-// NO_HALF: [[MAX:%.*]] = call reassoc nnan ninf nsz arcp afn noundef <4 x float> @llvm.maxnum.v4f32(<4 x float> %{{.*}}, <4 x float> [[CONV1]])
+// NO_HALF: [[MAX:%.*]] = call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) <4 x float> @llvm.maxnum.v4f32(<4 x float> %{{.*}}, <4 x float> [[CONV1]])
// NO_HALF: ret <4 x float> [[MAX]]
half4 test_max_half4_mismatch(half4 p0, half p1) { return max(p0, p1); }
// CHECK-LABEL: define hidden noundef nofpclass(nan inf) <4 x float> {{.*}}test_max_float4_mismatch
// CHECK: [[CONV0:%.*]] = insertelement <4 x float> poison, float %{{.*}}, i64 0
// CHECK: [[CONV1:%.*]] = shufflevector <4 x float> [[CONV0]], <4 x float> poison, <4 x i32> zeroinitializer
-// CHECK: [[MAX:%.*]] = call reassoc nnan ninf nsz arcp afn noundef <4 x float> @llvm.maxnum.v4f32(<4 x float> %{{.*}}, <4 x float> [[CONV1]])
+// CHECK: [[MAX:%.*]] = call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) <4 x float> @llvm.maxnum.v4f32(<4 x float> %{{.*}}, <4 x float> [[CONV1]])
// CHECK: ret <4 x float> [[MAX]]
float4 test_max_float4_mismatch(float4 p0, float p1) { return max(p0, p1); }
// CHECK-LABEL: define hidden noundef nofpclass(nan inf) <4 x double> {{.*}}test_max_double4_mismatch
// CHECK: [[CONV0:%.*]] = insertelement <4 x double> poison, double %{{.*}}, i64 0
// CHECK: [[CONV1:%.*]] = shufflevector <4 x double> [[CONV0]], <4 x double> poison, <4 x i32> zeroinitializer
-// CHECK: [[MAX:%.*]] = call reassoc nnan ninf nsz arcp afn noundef <4 x double> @llvm.maxnum.v4f64(<4 x double> %{{.*}}, <4 x double> [[CONV1]])
+// CHECK: [[MAX:%.*]] = call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) <4 x double> @llvm.maxnum.v4f64(<4 x double> %{{.*}}, <4 x double> [[CONV1]])
// CHECK: ret <4 x double> [[MAX]]
double4 test_max_double4_mismatch(double4 p0, double p1) { return max(p0, p1); }
// CHECK-LABEL: define hidden noundef nofpclass(nan inf) <4 x double> {{.*}}test_max_double4_mismatch2
// CHECK: [[CONV0:%.*]] = insertelement <4 x double> poison, double %{{.*}}, i64 0
// CHECK: [[CONV1:%.*]] = shufflevector <4 x double> [[CONV0]], <4 x double> poison, <4 x i32> zeroinitializer
-// CHECK: [[MAX:%.*]] = call reassoc nnan ninf nsz arcp afn noundef <4 x double> @llvm.maxnum.v4f64(<4 x double> [[CONV1]], <4 x double> %{{.*}})
+// CHECK: [[MAX:%.*]] = call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) <4 x double> @llvm.maxnum.v4f64(<4 x double> [[CONV1]], <4 x double> %{{.*}})
// CHECK: ret <4 x double> [[MAX]]
double4 test_max_double4_mismatch2(double4 p0, double p1) { return max(p1, p0); }
diff --git a/clang/test/CodeGenHLSL/builtins/min-overloads.hlsl b/clang/test/CodeGenHLSL/builtins/min-overloads.hlsl
index c0e06b0d204b3..ee3455a07c8e1 100644
--- a/clang/test/CodeGenHLSL/builtins/min-overloads.hlsl
+++ b/clang/test/CodeGenHLSL/builtins/min-overloads.hlsl
@@ -50,32 +50,32 @@ uint64_t4 test_min_ulong4_mismatch(uint64_t4 p0, uint64_t p1) { return min(p0, p
// NATIVE_HALF-LABEL: define hidden noundef nofpclass(nan inf) <4 x half> {{.*}}test_min_half4_mismatch
// NATIVE_HALF: [[CONV0:%.*]] = insertelement <4 x half> poison, half %{{.*}}, i64 0
// NATIVE_HALF: [[CONV1:%.*]] = shufflevector <4 x half> [[CONV0]], <4 x half> poison, <4 x i32> zeroinitializer
-// NATIVE_HALF: [[MIN:%.*]] = call reassoc nnan ninf nsz arcp afn noundef <4 x half> @llvm.minnum.v4f16(<4 x half> %{{.*}}, <4 x half> [[CONV1]])
+// NATIVE_HALF: [[MIN:%.*]] = call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) <4 x half> @llvm.minnum.v4f16(<4 x half> %{{.*}}, <4 x half> [[CONV1]])
// NATIVE_HALF: ret <4 x half> [[MIN]]
// NO_HALF-LABEL: define hidden noundef nofpclass(nan inf) <4 x float> {{.*}}test_min_half4_mismatch
// NO_HALF: [[CONV0:%.*]] = insertelement <4 x float> poison, float %{{.*}}, i64 0
// NO_HALF: [[CONV1:%.*]] = shufflevector <4 x float> [[CONV0]], <4 x float> poison, <4 x i32> zeroinitializer
-// NO_HALF: [[MIN:%.*]] = call reassoc nnan ninf nsz arcp afn noundef <4 x float> @llvm.minnum.v4f32(<4 x float> %{{.*}}, <4 x float> [[CONV1]])
+// NO_HALF: [[MIN:%.*]] = call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) <4 x float> @llvm.minnum.v4f32(<4 x float> %{{.*}}, <4 x float> [[CONV1]])
// NO_HALF: ret <4 x float> [[MIN]]
half4 test_min_half4_mismatch(half4 p0, half p1) { return min(p0, p1); }
// CHECK-LABEL: define hidden noundef nofpclass(nan inf) <4 x float> {{.*}}test_min_float4_mismatch
// CHECK: [[CONV0:%.*]] = insertelement <4 x float> poison, float %{{.*}}, i64 0
// CHECK: [[CONV1:%.*]] = shufflevector <4 x float> [[CONV0]], <4 x float> poison, <4 x i32> zeroinitializer
-// CHECK: [[MIN:%.*]] = call reassoc nnan ninf nsz arcp afn noundef <4 x float> @llvm.minnum.v4f32(<4 x float> %{{.*}}, <4 x float> [[CONV1]])
+// CHECK: [[MIN:%.*]] = call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) <4 x float> @llvm.minnum.v4f32(<4 x float> %{{.*}}, <4 x float> [[CONV1]])
// CHECK: ret <4 x float> [[MIN]]
float4 test_min_float4_mismatch(float4 p0, float p1) { return min(p0, p1); }
// CHECK-LABEL: define hidden noundef nofpclass(nan inf) <4 x double> {{.*}}test_min_double4_mismatch
// CHECK: [[CONV0:%.*]] = insertelement <4 x double> poison, double %{{.*}}, i64 0
// CHECK: [[CONV1:%.*]] = shufflevector <4 x double> [[CONV0]], <4 x double> poison, <4 x i32> zeroinitializer
-// CHECK: [[MIN:%.*]] = call reassoc nnan ninf nsz arcp afn noundef <4 x double> @llvm.minnum.v4f64(<4 x double> %{{.*}}, <4 x double> [[CONV1]])
+// CHECK: [[MIN:%.*]] = call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) <4 x double> @llvm.minnum.v4f64(<4 x double> %{{.*}}, <4 x double> [[CONV1]])
// CHECK: ret <4 x double> [[MIN]]
double4 test_min_double4_mismatch(double4 p0, double p1) { return min(p0, p1); }
// CHECK-LABEL: define hidden noundef nofpclass(nan inf) <4 x double> {{.*}}test_min_double4_mismatch2
// CHECK: [[CONV0:%.*]] = insertelement <4 x double> poison, double %{{.*}}, i64 0
// CHECK: [[CONV1:%.*]] = shufflevector <4 x double> [[CONV0]], <4 x double> poison, <4 x i32> zeroinitializer
-// CHECK: [[MIN:%.*]] = call reassoc nnan ninf nsz arcp afn noundef <4 x double> @llvm.minnum.v4f64(<4 x double> [[CONV1]], <4 x double> %{{.*}})
+// CHECK: [[MIN:%.*]] = call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) <4 x double> @llvm.minnum.v4f64(<4 x double> [[CONV1]], <4 x double> %{{.*}})
// CHECK: ret <4 x double> [[MIN]]
double4 test_min_double4_mismatch2(double4 p0, double p1) { return min(p1, p0); }
diff --git a/clang/test/CodeGenHLSL/builtins/pow-overloads.hlsl b/clang/test/CodeGenHLSL/builtins/pow-overloads.hlsl
index 0d1f3d3546a33..6a52051aefc62 100644
--- a/clang/test/CodeGenHLSL/builtins/pow-overloads.hlsl
+++ b/clang/test/CodeGenHLSL/builtins/pow-overloads.hlsl
@@ -5,124 +5,124 @@
// CHECK-LABEL: define hidden noundef nofpclass(nan inf) float {{.*}}test_pow_double
// CHECK: [[CONV0:%.*]] = fptrunc [[FLOATATTRS]] double %{{.*}} to float
// CHECK: [[CONV1:%.*]] = fptrunc [[FLOATATTRS]] double %{{.*}} to float
-// CHECK: [[POW:%.*]] = call [[FLOATATTRS]] noundef float @llvm.pow.f32(float [[CONV0]], float [[CONV1]])
+// CHECK: [[POW:%.*]] = call [[FLOATATTRS]] noundef nofpclass(nan inf) float @llvm.pow.f32(float [[CONV0]], float [[CONV1]])
// CHECK: ret float [[POW]]
float test_pow_double(double p0, double p1) { return pow(p0, p1); }
// CHECK-LABEL: define hidden noundef nofpclass(nan inf) <2 x float> {{.*}}test_pow_double2
// CHECK: [[CONV0:%.*]] = fptrunc [[FLOATATTRS]] <2 x double> %{{.*}} to <2 x float>
// CHECK: [[CONV1:%.*]] = fptrunc [[FLOATATTRS]] <2 x double> %{{.*}} to <2 x float>
-// CHECK: [[POW:%.*]] = call [[FLOATATTRS]] noundef <2 x float> @llvm.pow.v2f32(<2 x float> [[CONV0]], <2 x float> [[CONV1]])
+// CHECK: [[POW:%.*]] = call [[FLOATATTRS]] noundef nofpclass(nan inf) <2 x float> @llvm.pow.v2f32(<2 x float> [[CONV0]], <2 x float> [[CONV1]])
// CHECK: ret <2 x float> [[POW]]
float2 test_pow_double2(double2 p0, double2 p1) { return pow(p0, p1); }
// CHECK-LABEL: define hidden noundef nofpclass(nan inf) <3 x float> {{.*}}test_pow_double3
// CHECK: [[CONV0:%.*]] = fptrunc [[FLOATATTRS]] <3 x double> %{{.*}} to <3 x float>
// CHECK: [[CONV1:%.*]] = fptrunc [[FLOATATTRS]] <3 x double> %{{.*}} to <3 x float>
-// CHECK: [[POW:%.*]] = call [[FLOATATTRS]] noundef <3 x float> @llvm.pow.v3f32(<3 x float> [[CONV0]], <3 x float> [[CONV1]])
+// CHECK: [[POW:%.*]] = call [[FLOATATTRS]] noundef nofpclass(nan inf) <3 x float> @llvm.pow.v3f32(<3 x float> [[CONV0]], <3 x float> [[CONV1]])
// CHECK: ret <3 x float> [[POW]]
float3 test_pow_double3(double3 p0, double3 p1) { return pow(p0, p1); }
// CHECK-LABEL: define hidden noundef nofpclass(nan inf) <4 x float> {{.*}}test_pow_double4
// CHECK: [[CONV0:%.*]] = fptrunc [[FLOATATTRS]] <4 x double> %{{.*}} to <4 x float>
// CHECK: [[CONV1:%.*]] = fptrunc [[FLOATATTRS]] <4 x double> %{{.*}} to <4 x float>
-// CHECK: [[POW:%.*]] = call [[FLOATATTRS]] noundef <4 x float> @llvm.pow.v4f32(<4 x float> [[CONV0]], <4 x float> [[CONV1]])
+// CHECK: [[POW:%.*]] = call [[FLOATATTRS]] noundef nofpclass(nan inf) <4 x float> @llvm.pow.v4f32(<4 x float> [[CONV0]], <4 x float> [[CONV1]])
// CHECK: ret <4 x float> [[POW]]
float4 test_pow_double4(double4 p0, double4 p1) { return pow(p0, p1); }
// CHECK-LABEL: define hidden noundef nofpclass(nan inf) float {{.*}}test_pow_int
// CHECK: [[CONV0:%.*]] = sitofp i32 %{{.*}} to float
// CHECK: [[CONV1:%.*]] = sitofp i32 %{{.*}} to float
-// CHECK: [[POW:%.*]] = call [[FLOATATTRS]] noundef float @llvm.pow.f32(float [[CONV0]], float [[CONV1]])
+// CHECK: [[POW:%.*]] = call [[FLOATATTRS]] noundef nofpclass(nan inf) float @llvm.pow.f32(float [[CONV0]], float [[CONV1]])
// CHECK: ret float [[POW]]
float test_pow_int(int p0, int p1) { return pow(p0, p1); }
// CHECK-LABEL: define hidden noundef nofpclass(nan inf) <2 x float> {{.*}}test_pow_int2
// CHECK: [[CONV0:%.*]] = sitofp <2 x i32> %{{.*}} to <2 x float>
// CHECK: [[CONV1:%.*]] = sitofp <2 x i32> %{{.*}} to <2 x float>
-// CHECK: [[POW:%.*]] = call [[FLOATATTRS]] noundef <2 x float> @llvm.pow.v2f32(<2 x float> [[CONV0]], <2 x float> [[CONV1]])
+// CHECK: [[POW:%.*]] = call [[FLOATATTRS]] noundef nofpclass(nan inf) <2 x float> @llvm.pow.v2f32(<2 x float> [[CONV0]], <2 x float> [[CONV1]])
// CHECK: ret <2 x float> [[POW]]
float2 test_pow_int2(int2 p0, int2 p1) { return pow(p0, p1); }
// CHECK-LABEL: define hidden noundef nofpclass(nan inf) <3 x float> {{.*}}test_pow_int3
// CHECK: [[CONV0:%.*]] = sitofp <3 x i32> %{{.*}} to <3 x float>
// CHECK: [[CONV1:%.*]] = sitofp <3 x i32> %{{.*}} to <3 x float>
-// CHECK: [[POW:%.*]] = call [[FLOATATTRS]] noundef <3 x float> @llvm.pow.v3f32(<3 x float> [[CONV0]], <3 x float> [[CONV1]])
+// CHECK: [[POW:%.*]] = call [[FLOATATTRS]] noundef nofpclass(nan inf) <3 x float> @llvm.pow.v3f32(<3 x float> [[CONV0]], <3 x float> [[CONV1]])
// CHECK: ret <3 x float> [[POW]]
float3 test_pow_int3(int3 p0, int3 p1) { return pow(p0, p1); }
// CHECK-LABEL: define hidden noundef nofpclass(nan inf) <4 x float> {{.*}}test_pow_int4
// CHECK: [[CONV0:%.*]] = sitofp <4 x i32> %{{.*}} to <4 x float>
// CHECK: [[CONV1:%.*]] = sitofp <4 x i32> %{{.*}} to <4 x float>
-// CHECK: [[POW:%.*]] = call [[FLOATATTRS]] noundef <4 x float> @llvm.pow.v4f32(<4 x float> [[CONV0]], <4 x float> [[CONV1]])
+// CHECK: [[POW:%.*]] = call [[FLOATATTRS]] noundef nofpclass(nan inf) <4 x float> @llvm.pow.v4f32(<4 x float> [[CONV0]], <4 x float> [[CONV1]])
// CHECK: ret <4 x float> [[POW]]
float4 test_pow_int4(int4 p0, int4 p1) { return pow(p0, p1); }
// CHECK-LABEL: define hidden noundef nofpclass(nan inf) float {{.*}}test_pow_uint
// CHECK: [[CONV0:%.*]] = uitofp i32 %{{.*}} to float
// CHECK: [[CONV1:%.*]] = uitofp i32 %{{.*}} to float
-// CHECK: [[POW:%.*]] = call [[FLOATATTRS]] noundef float @llvm.pow.f32(float [[CONV0]], float [[CONV1]])
+// CHECK: [[POW:%.*]] = call [[FLOATATTRS]] noundef nofpclass(nan inf) float @llvm.pow.f32(float [[CONV0]], float [[CONV1]])
// CHECK: ret float [[POW]]
float test_pow_uint(uint p0, uint p1) { return pow(p0, p1); }
// CHECK-LABEL: define hidden noundef nofpclass(nan inf) <2 x float> {{.*}}test_pow_uint2
// CHECK: [[CONV0:%.*]] = uitofp <2 x i32> %{{.*}} to <2 x float>
// CHECK: [[CONV1:%.*]] = uitofp <2 x i32> %{{.*}} to <2 x float>
-// CHECK: [[POW:%.*]] = call [[FLOATATTRS]] noundef <2 x float> @llvm.pow.v2f32(<2 x float> [[CONV0]], <2 x float> [[CONV1]])
+// CHECK: [[POW:%.*]] = call [[FLOATATTRS]] noundef nofpclass(nan inf) <2 x float> @llvm.pow.v2f32(<2 x float> [[CONV0]], <2 x float> [[CONV1]])
// CHECK: ret <2 x float> [[POW]]
float2 test_pow_uint2(uint2 p0, uint2 p1) { return pow(p0, p1); }
// CHECK-LABEL: define hidden noundef nofpclass(nan inf) <3 x float> {{.*}}test_pow_uint3
// CHECK: [[CONV0:%.*]] = uitofp <3 x i32> %{{.*}} to <3 x float>
// CHECK: [[CONV1:%.*]] = uitofp <3 x i32> %{{.*}} to <3 x float>
-// CHECK: [[POW:%.*]] = call [[FLOATATTRS]] noundef <3 x float> @llvm.pow.v3f32(<3 x float> [[CONV0]], <3 x float> [[CONV1]])
+// CHECK: [[POW:%.*]] = call [[FLOATATTRS]] noundef nofpclass(nan inf) <3 x float> @llvm.pow.v3f32(<3 x float> [[CONV0]], <3 x float> [[CONV1]])
// CHECK: ret <3 x float> [[POW]]
float3 test_pow_uint3(uint3 p0, uint3 p1) { return pow(p0, p1); }
// CHECK-LABEL: define hidden noundef nofpclass(nan inf) <4 x float> {{.*}}test_pow_uint4
// CHECK: [[CONV0:%.*]] = uitofp <4 x i32> %{{.*}} to <4 x float>
// CHECK: [[CONV1:%.*]] = uitofp <4 x i32> %{{.*}} to <4 x float>
-// CHECK: [[POW:%.*]] = call [[FLOATATTRS]] noundef <4 x float> @llvm.pow.v4f32(<4 x float> [[CONV0]], <4 x float> [[CONV1]])
+// CHECK: [[POW:%.*]] = call [[FLOATATTRS]] noundef nofpclass(nan inf) <4 x float> @llvm.pow.v4f32(<4 x float> [[CONV0]], <4 x float> [[CONV1]])
// CHECK: ret <4 x float> [[POW]]
float4 test_pow_uint4(uint4 p0, uint4 p1) { return pow(p0, p1); }
// CHECK-LABEL: define hidden noundef nofpclass(nan inf) float {{.*}}test_pow_int64_t
// CHECK: [[CONV0:%.*]] = sitofp i64 %{{.*}} to float
// CHECK: [[CONV1:%.*]] = sitofp i64 %{{.*}} to float
-// CHECK: [[POW:%.*]] = call [[FLOATATTRS]] noundef float @llvm.pow.f32(float [[CONV0]], float [[CONV1]])
+// CHECK: [[POW:%.*]] = call [[FLOATATTRS]] noundef nofpclass(nan inf) float @llvm.pow.f32(float [[CONV0]], float [[CONV1]])
// CHECK: ret float [[POW]]
float test_pow_int64_t(int64_t p0, int64_t p1) { return pow(p0, p1); }
// CHECK-LABEL: define hidden noundef nofpclass(nan inf) <2 x float> {{.*}}test_pow_int64_t2
// CHECK: [[CONV0:%.*]] = sitofp <2 x i64> %{{.*}} to <2 x float>
// CHECK: [[CONV1:%.*]] = sitofp <2 x i64> %{{.*}} to <2 x float>
-// CHECK: [[POW:%.*]] = call [[FLOATATTRS]] noundef <2 x float> @llvm.pow.v2f32(<2 x float> [[CONV0]], <2 x float> [[CONV1]])
+// CHECK: [[POW:%.*]] = call [[FLOATATTRS]] noundef nofpclass(nan inf) <2 x float> @llvm.pow.v2f32(<2 x float> [[CONV0]], <2 x float> [[CONV1]])
// CHECK: ret <2 x float> [[POW]]
float2 test_pow_int64_t2(int64_t2 p0, int64_t2 p1) { return pow(p0, p1); }
// CHECK-LABEL: define hidden noundef nofpclass(nan inf) <3 x float> {{.*}}test_pow_int64_t3
// CHECK: [[CONV0:%.*]] = sitofp <3 x i64> %{{.*}} to <3 x float>
// CHECK: [[CONV1:%.*]] = sitofp <3 x i64> %{{.*}} to <3 x float>
-// CHECK: [[POW:%.*]] = call [[FLOATATTRS]] noundef <3 x float> @llvm.pow.v3f32(<3 x float> [[CONV0]], <3 x float> [[CONV1]])
+// CHECK: [[POW:%.*]] = call [[FLOATATTRS]] noundef nofpclass(nan inf) <3 x float> @llvm.pow.v3f32(<3 x float> [[CONV0]], <3 x float> [[CONV1]])
// CHECK: ret <3 x float> [[POW]]
float3 test_pow_int64_t3(int64_t3 p0, int64_t3 p1) { return pow(p0, p1); }
// CHECK-LABEL: define hidden noundef nofpclass(nan inf) <4 x float> {{.*}}test_pow_int64_t4
// CHECK: [[CONV0:%.*]] = sitofp <4 x i64> %{{.*}} to <4 x float>
// CHECK: [[CONV1:%.*]] = sitofp <4 x i64> %{{.*}} to <4 x float>
-// CHECK: [[POW:%.*]] = call [[FLOATATTRS]] noundef <4 x float> @llvm.pow.v4f32(<4 x float> [[CONV0]], <4 x float> [[CONV1]])
+// CHECK: [[POW:%.*]] = call [[FLOATATTRS]] noundef nofpclass(nan inf) <4 x float> @llvm.pow.v4f32(<4 x float> [[CONV0]], <4 x float> [[CONV1]])
// CHECK: ret <4 x float> [[POW]]
float4 test_pow_int64_t4(int64_t4 p0, int64_t4 p1) { return pow(p0, p1); }
// CHECK-LABEL: define hidden noundef nofpclass(nan inf) float {{.*}}test_pow_uint64_t
// CHECK: [[CONV0:%.*]] = uitofp i64 %{{.*}} to float
// CHECK: [[CONV1:%.*]] = uitofp i64 %{{.*}} to float
-// CHECK: [[POW:%.*]] = call [[FLOATATTRS]] noundef float @llvm.pow.f32(float [[CONV0]], float [[CONV1]])
+// CHECK: [[POW:%.*]] = call [[FLOATATTRS]] noundef nofpclass(nan inf) float @llvm.pow.f32(float [[CONV0]], float [[CONV1]])
// CHECK: ret float [[POW]]
float test_pow_uint64_t(uint64_t p0, uint64_t p1) { return pow(p0, p1); }
// CHECK-LABEL: define hidden noundef nofpclass(nan inf) <2 x float> {{.*}}test_pow_uint64_t2
// CHECK: [[CONV0:%.*]] = uitofp <2 x i64> %{{.*}} to <2 x float>
// CHECK: [[CONV1:%.*]] = uitofp <2 x i64> %{{.*}} to <2 x float>
-// CHECK: [[POW:%.*]] = call [[FLOATATTRS]] noundef <2 x float> @llvm.pow.v2f32(<2 x float> [[CONV0]], <2 x float> [[CONV1]])
+// CHECK: [[POW:%.*]] = call [[FLOATATTRS]] noundef nofpclass(nan inf) <2 x float> @llvm.pow.v2f32(<2 x float> [[CONV0]], <2 x float> [[CONV1]])
// CHECK: ret <2 x float> [[POW]]
float2 test_pow_uint64_t2(uint64_t2 p0, uint64_t2 p1) { return pow(p0, p1); }
// CHECK-LABEL: define hidden noundef nofpclass(nan inf) <3 x float> {{.*}}test_pow_uint64_t3
// CHECK: [[CONV0:%.*]] = uitofp <3 x i64> %{{.*}} to <3 x float>
// CHECK: [[CONV1:%.*]] = uitofp <3 x i64> %{{.*}} to <3 x float>
-// CHECK: [[POW:%.*]] = call [[FLOATATTRS]] noundef <3 x float> @llvm.pow.v3f32(<3 x float> [[CONV0]], <3 x float> [[CONV1]])
+// CHECK: [[POW:%.*]] = call [[FLOATATTRS]] noundef nofpclass(nan inf) <3 x float> @llvm.pow.v3f32(<3 x float> [[CONV0]], <3 x float> [[CONV1]])
// CHECK: ret <3 x float> [[POW]]
float3 test_pow_uint64_t3(uint64_t3 p0, uint64_t3 p1) { return pow(p0, p1); }
// CHECK-LABEL: define hidden noundef nofpclass(nan inf) <4 x float> {{.*}}test_pow_uint64_t4
// CHECK: [[CONV0:%.*]] = uitofp <4 x i64> %{{.*}} to <4 x float>
// CHECK: [[CONV1:%.*]] = uitofp <4 x i64> %{{.*}} to <4 x float>
-// CHECK: [[POW:%.*]] = call [[FLOATATTRS]] noundef <4 x float> @llvm.pow.v4f32(<4 x float> [[CONV0]], <4 x float> [[CONV1]])
+// CHECK: [[POW:%.*]] = call [[FLOATATTRS]] noundef nofpclass(nan inf) <4 x float> @llvm.pow.v4f32(<4 x float> [[CONV0]], <4 x float> [[CONV1]])
// CHECK: ret <4 x float> [[POW]]
float4 test_pow_uint64_t4(uint64_t4 p0, uint64_t4 p1) { return pow(p0, p1); }
diff --git a/clang/test/CodeGenHLSL/builtins/reflect.hlsl b/clang/test/CodeGenHLSL/builtins/reflect.hlsl
index feb5a5b2ea78f..259490afb1c78 100644
--- a/clang/test/CodeGenHLSL/builtins/reflect.hlsl
+++ b/clang/test/CodeGenHLSL/builtins/reflect.hlsl
@@ -42,7 +42,7 @@ half test_reflect_half(half I, half N) {
// SPVCHECK-LABEL: define hidden spir_func noundef nofpclass(nan inf) <2 x half> @_Z18test_reflect_half2Dv2_DhS_(
// SPVCHECK-SAME: <2 x half> noundef nofpclass(nan inf) [[I:%.*]], <2 x half> noundef nofpclass(nan inf) [[N:%.*]]) local_unnamed_addr #[[ATTR0]] {
// SPVCHECK-NEXT: [[ENTRY:.*:]]
-// SPVCHECK-NEXT: [[SPV_REFLECT_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef <2 x half> @llvm.spv.reflect.v2f16(<2 x half> nofpclass(nan inf) [[I]], <2 x half> nofpclass(nan inf) [[N]])
+// SPVCHECK-NEXT: [[SPV_REFLECT_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) <2 x half> @llvm.spv.reflect.v2f16(<2 x half> nofpclass(nan inf) [[I]], <2 x half> nofpclass(nan inf) [[N]])
// SPVCHECK-NEXT: ret <2 x half> [[SPV_REFLECT_I]]
//
half2 test_reflect_half2(half2 I, half2 N) {
@@ -63,7 +63,7 @@ half2 test_reflect_half2(half2 I, half2 N) {
// SPVCHECK-LABEL: define hidden spir_func noundef nofpclass(nan inf) <3 x half> @_Z18test_reflect_half3Dv3_DhS_(
// SPVCHECK-SAME: <3 x half> noundef nofpclass(nan inf) [[I:%.*]], <3 x half> noundef nofpclass(nan inf) [[N:%.*]]) local_unnamed_addr #[[ATTR0]] {
// SPVCHECK-NEXT: [[ENTRY:.*:]]
-// SPVCHECK-NEXT: [[SPV_REFLECT_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef <3 x half> @llvm.spv.reflect.v3f16(<3 x half> nofpclass(nan inf) [[I]], <3 x half> nofpclass(nan inf) [[N]])
+// SPVCHECK-NEXT: [[SPV_REFLECT_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) <3 x half> @llvm.spv.reflect.v3f16(<3 x half> nofpclass(nan inf) [[I]], <3 x half> nofpclass(nan inf) [[N]])
// SPVCHECK-NEXT: ret <3 x half> [[SPV_REFLECT_I]]
//
half3 test_reflect_half3(half3 I, half3 N) {
@@ -84,7 +84,7 @@ half3 test_reflect_half3(half3 I, half3 N) {
// SPVCHECK-LABEL: define hidden spir_func noundef nofpclass(nan inf) <4 x half> @_Z18test_reflect_half4Dv4_DhS_(
// SPVCHECK-SAME: <4 x half> noundef nofpclass(nan inf) [[I:%.*]], <4 x half> noundef nofpclass(nan inf) [[N:%.*]]) local_unnamed_addr #[[ATTR0]] {
// SPVCHECK-NEXT: [[ENTRY:.*:]]
-// SPVCHECK-NEXT: [[SPV_REFLECT_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef <4 x half> @llvm.spv.reflect.v4f16(<4 x half> nofpclass(nan inf) [[I]], <4 x half> nofpclass(nan inf) [[N]])
+// SPVCHECK-NEXT: [[SPV_REFLECT_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) <4 x half> @llvm.spv.reflect.v4f16(<4 x half> nofpclass(nan inf) [[I]], <4 x half> nofpclass(nan inf) [[N]])
// SPVCHECK-NEXT: ret <4 x half> [[SPV_REFLECT_I]]
//
half4 test_reflect_half4(half4 I, half4 N) {
@@ -127,7 +127,7 @@ float test_reflect_float(float I, float N) {
// SPVCHECK-LABEL: define hidden spir_func noundef nofpclass(nan inf) <2 x float> @_Z19test_reflect_float2Dv2_fS_(
// SPVCHECK-SAME: <2 x float> noundef nofpclass(nan inf) [[I:%.*]], <2 x float> noundef nofpclass(nan inf) [[N:%.*]]) local_unnamed_addr #[[ATTR0]] {
// SPVCHECK-NEXT: [[ENTRY:.*:]]
-// SPVCHECK-NEXT: [[SPV_REFLECT_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef <2 x float> @llvm.spv.reflect.v2f32(<2 x float> nofpclass(nan inf) [[I]], <2 x float> nofpclass(nan inf) [[N]])
+// SPVCHECK-NEXT: [[SPV_REFLECT_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) <2 x float> @llvm.spv.reflect.v2f32(<2 x float> nofpclass(nan inf) [[I]], <2 x float> nofpclass(nan inf) [[N]])
// SPVCHECK-NEXT: ret <2 x float> [[SPV_REFLECT_I]]
//
float2 test_reflect_float2(float2 I, float2 N) {
@@ -148,7 +148,7 @@ float2 test_reflect_float2(float2 I, float2 N) {
// SPVCHECK-LABEL: define hidden spir_func noundef nofpclass(nan inf) <3 x float> @_Z19test_reflect_float3Dv3_fS_(
// SPVCHECK-SAME: <3 x float> noundef nofpclass(nan inf) [[I:%.*]], <3 x float> noundef nofpclass(nan inf) [[N:%.*]]) local_unnamed_addr #[[ATTR0]] {
// SPVCHECK-NEXT: [[ENTRY:.*:]]
-// SPVCHECK-NEXT: [[SPV_REFLECT_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef <3 x float> @llvm.spv.reflect.v3f32(<3 x float> nofpclass(nan inf) [[I]], <3 x float> nofpclass(nan inf) [[N]])
+// SPVCHECK-NEXT: [[SPV_REFLECT_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) <3 x float> @llvm.spv.reflect.v3f32(<3 x float> nofpclass(nan inf) [[I]], <3 x float> nofpclass(nan inf) [[N]])
// SPVCHECK-NEXT: ret <3 x float> [[SPV_REFLECT_I]]
//
float3 test_reflect_float3(float3 I, float3 N) {
@@ -169,7 +169,7 @@ float3 test_reflect_float3(float3 I, float3 N) {
// SPVCHECK-LABEL: define hidden spir_func noundef nofpclass(nan inf) <4 x float> @_Z19test_reflect_float4Dv4_fS_(
// SPVCHECK-SAME: <4 x float> noundef nofpclass(nan inf) [[I:%.*]], <4 x float> noundef nofpclass(nan inf) [[N:%.*]]) local_unnamed_addr #[[ATTR0]] {
// SPVCHECK-NEXT: [[ENTRY:.*:]]
-// SPVCHECK-NEXT: [[SPV_REFLECT_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef <4 x float> @llvm.spv.reflect.v4f32(<4 x float> nofpclass(nan inf) [[I]], <4 x float> nofpclass(nan inf) [[N]])
+// SPVCHECK-NEXT: [[SPV_REFLECT_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) <4 x float> @llvm.spv.reflect.v4f32(<4 x float> nofpclass(nan inf) [[I]], <4 x float> nofpclass(nan inf) [[N]])
// SPVCHECK-NEXT: ret <4 x float> [[SPV_REFLECT_I]]
//
float4 test_reflect_float4(float4 I, float4 N) {
diff --git a/clang/test/CodeGenHLSL/builtins/refract.hlsl b/clang/test/CodeGenHLSL/builtins/refract.hlsl
index ffeb2a78b2517..ecc36d5c7aadf 100644
--- a/clang/test/CodeGenHLSL/builtins/refract.hlsl
+++ b/clang/test/CodeGenHLSL/builtins/refract.hlsl
@@ -25,9 +25,9 @@
// CHECK: ret half [[HLSL_SELECT_I]]
//
// SPVCHECK-LABEL: define hidden spir_func noundef nofpclass(nan inf) half @_Z17test_refract_halfDhDhDh(
-// SPVCHECK-SAME: half noundef nofpclass(nan inf) [[I:%.*]], half noundef nofpclass(nan inf) [[N:%.*]], half noundef nofpclass(nan inf) [[ETA:%.*]]) #[[ATTR0:[0-9]+]]
+// SPVCHECK-SAME: half noundef nofpclass(nan inf) [[I:%.*]], half noundef nofpclass(nan inf) [[N:%.*]], half noundef nofpclass(nan inf) [[ETA:%.*]]) #[[ATTR0:[0-9]+]]
// SPVCHECK: [[ENTRY:.*:]]
-// SPVCHECK: [[SPV_REFRACT_I:%.*]] = call reassoc nnan ninf nsz arcp afn noundef half @llvm.spv.refract.f16.f16(half %{{.*}}, half %{{.*}}, half %{{.*}})
+// SPVCHECK: [[SPV_REFRACT_I:%.*]] = call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) half @llvm.spv.refract.f16.f16(half %{{.*}}, half %{{.*}}, half %{{.*}})
// SPVCHECK: ret half [[SPV_REFRACT_I]]
//
half test_refract_half(half I, half N, half ETA) {
@@ -57,7 +57,7 @@ half test_refract_half(half I, half N, half ETA) {
// SPVCHECK-LABEL: define hidden spir_func noundef nofpclass(nan inf) <2 x half> @_Z18test_refract_half2Dv2_DhS_Dh(
// SPVCHECK-SAME: <2 x half> noundef nofpclass(nan inf) [[I:%.*]], <2 x half> noundef nofpclass(nan inf) [[N:%.*]], half noundef nofpclass(nan inf) [[ETA:%.*]]) #[[ATTR0:[0-9]+]] {
// SPVCHECK: [[ENTRY:.*:]]
-// SPVCHECK: [[SPV_REFRACT_I:%.*]] = call reassoc nnan ninf nsz arcp afn noundef <2 x half> @llvm.spv.refract.v2f16.f16(<2 x half> %{{.*}}, <2 x half> %{{.*}}, half %{{.*}})
+// SPVCHECK: [[SPV_REFRACT_I:%.*]] = call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) <2 x half> @llvm.spv.refract.v2f16.f16(<2 x half> %{{.*}}, <2 x half> %{{.*}}, half %{{.*}})
// SPVCHECK: ret <2 x half> [[SPV_REFRACT_I]]
//
half2 test_refract_half2(half2 I, half2 N, half ETA) {
@@ -87,7 +87,7 @@ half2 test_refract_half2(half2 I, half2 N, half ETA) {
// SPVCHECK-LABEL: define hidden spir_func noundef nofpclass(nan inf) <3 x half> @_Z18test_refract_half3Dv3_DhS_Dh(
// SPVCHECK-SAME: <3 x half> noundef nofpclass(nan inf) [[I:%.*]], <3 x half> noundef nofpclass(nan inf) [[N:%.*]], half noundef nofpclass(nan inf) [[ETA:%.*]]) #[[ATTR0:[0-9]+]] {
// SPVCHECK: [[ENTRY:.*:]]
-// SPVCHECK: [[SPV_REFRACT_I:%.*]] = call reassoc nnan ninf nsz arcp afn noundef <3 x half> @llvm.spv.refract.v3f16.f16(<3 x half> %{{.*}}, <3 x half> %{{.*}}, half %{{.*}})
+// SPVCHECK: [[SPV_REFRACT_I:%.*]] = call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) <3 x half> @llvm.spv.refract.v3f16.f16(<3 x half> %{{.*}}, <3 x half> %{{.*}}, half %{{.*}})
// SPVCHECK: ret <3 x half> [[SPV_REFRACT_I]]
//
half3 test_refract_half3(half3 I, half3 N, half ETA) {
@@ -117,7 +117,7 @@ half3 test_refract_half3(half3 I, half3 N, half ETA) {
// SPVCHECK-LABEL: define hidden spir_func noundef nofpclass(nan inf) <4 x half> @_Z18test_refract_half4Dv4_DhS_Dh(
// SPVCHECK-SAME: <4 x half> noundef nofpclass(nan inf) [[I:%.*]], <4 x half> noundef nofpclass(nan inf) [[N:%.*]], half noundef nofpclass(nan inf) [[ETA:%.*]]) #[[ATTR0:[0-9]+]] {
// SPVCHECK: [[ENTRY:.*:]]
-// SPVCHECK: [[SPV_REFRACT_I:%.*]] = call reassoc nnan ninf nsz arcp afn noundef <4 x half> @llvm.spv.refract.v4f16.f16(<4 x half> %{{.*}}, <4 x half> %{{.*}}, half %{{.*}})
+// SPVCHECK: [[SPV_REFRACT_I:%.*]] = call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) <4 x half> @llvm.spv.refract.v4f16.f16(<4 x half> %{{.*}}, <4 x half> %{{.*}}, half %{{.*}})
// SPVCHECK: ret <4 x half> [[SPV_REFRACT_I]]
//
half4 test_refract_half4(half4 I, half4 N, half ETA) {
@@ -146,7 +146,7 @@ half4 test_refract_half4(half4 I, half4 N, half ETA) {
// SPVCHECK-LABEL: define hidden spir_func noundef nofpclass(nan inf) float @_Z18test_refract_floatfff(
// SPVCHECK-SAME: float noundef nofpclass(nan inf) [[I:%.*]], float noundef nofpclass(nan inf) [[N:%.*]], float noundef nofpclass(nan inf) [[ETA:%.*]]) #[[ATTR0:[0-9]+]] {
// SPVCHECK: [[ENTRY:.*:]]
-// SPVCHECK: [[SPV_REFRACT_I:%.*]] = call reassoc nnan ninf nsz arcp afn noundef float @llvm.spv.refract.f32.f32(float %{{.*}}, float %{{.*}}, float %{{.*}})
+// SPVCHECK: [[SPV_REFRACT_I:%.*]] = call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) float @llvm.spv.refract.f32.f32(float %{{.*}}, float %{{.*}}, float %{{.*}})
// SPVCHECK: ret float [[SPV_REFRACT_I]]
//
float test_refract_float(float I, float N, float ETA) {
@@ -176,7 +176,7 @@ float test_refract_float(float I, float N, float ETA) {
// SPVCHECK-LABEL: define hidden spir_func noundef nofpclass(nan inf) <2 x float> @_Z19test_refract_float2Dv2_fS_f(
// SPVCHECK-SAME: <2 x float> noundef nofpclass(nan inf) [[I:%.*]], <2 x float> noundef nofpclass(nan inf) [[N:%.*]], float noundef nofpclass(nan inf) [[ETA:%.*]]) #[[ATTR0:[0-9]+]] {
// SPVCHECK: [[ENTRY:.*:]]
-// SPVCHECK: [[SPV_REFRACT_I:%.*]] = call reassoc nnan ninf nsz arcp afn noundef <2 x float> @llvm.spv.refract.v2f32.f32(<2 x float> %{{.*}}, <2 x float> %{{.*}}, float %{{.*}})
+// SPVCHECK: [[SPV_REFRACT_I:%.*]] = call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) <2 x float> @llvm.spv.refract.v2f32.f32(<2 x float> %{{.*}}, <2 x float> %{{.*}}, float %{{.*}})
// SPVCHECK: ret <2 x float> [[SPV_REFRACT_I]]
//
float2 test_refract_float2(float2 I, float2 N, float ETA) {
@@ -206,7 +206,7 @@ float2 test_refract_float2(float2 I, float2 N, float ETA) {
// SPVCHECK-LABEL: define hidden spir_func noundef nofpclass(nan inf) <3 x float> @_Z19test_refract_float3Dv3_fS_f(
// SPVCHECK-SAME: <3 x float> noundef nofpclass(nan inf) [[I:%.*]], <3 x float> noundef nofpclass(nan inf) [[N:%.*]], float noundef nofpclass(nan inf) [[ETA:%.*]]) #[[ATTR0:[0-9]+]] {
// SPVCHECK: [[ENTRY:.*:]]
-// SPVCHECK: [[SPV_REFRACT_I:%.*]] = call reassoc nnan ninf nsz arcp afn noundef <3 x float> @llvm.spv.refract.v3f32.f32(<3 x float> %{{.*}}, <3 x float> %{{.*}}, float %{{.*}})
+// SPVCHECK: [[SPV_REFRACT_I:%.*]] = call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) <3 x float> @llvm.spv.refract.v3f32.f32(<3 x float> %{{.*}}, <3 x float> %{{.*}}, float %{{.*}})
// SPVCHECK: ret <3 x float> [[SPV_REFRACT_I]]
//
float3 test_refract_float3(float3 I, float3 N, float ETA) {
@@ -236,7 +236,7 @@ float3 test_refract_float3(float3 I, float3 N, float ETA) {
// SPVCHECK-LABEL: define hidden spir_func noundef nofpclass(nan inf) <4 x float> @_Z19test_refract_float4Dv4_fS_f(
// SPVCHECK-SAME: <4 x float> noundef nofpclass(nan inf) %{{.*}}, <4 x float> noundef nofpclass(nan inf) %{{.*}}, float noundef nofpclass(nan inf) %{{.*}}) #[[ATTR0:[0-9]+]] {
// SPVCHECK: [[ENTRY:.*:]]
-// SPVCHECK: [[SPV_REFRACT_I:%.*]] = call reassoc nnan ninf nsz arcp afn noundef <4 x float> @llvm.spv.refract.v4f32.f32(<4 x float> %{{.*}}, <4 x float> %{{.*}}, float %{{.*}})
+// SPVCHECK: [[SPV_REFRACT_I:%.*]] = call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) <4 x float> @llvm.spv.refract.v4f32.f32(<4 x float> %{{.*}}, <4 x float> %{{.*}}, float %{{.*}})
// SPVCHECK: ret <4 x float> [[SPV_REFRACT_I]]
//
float4 test_refract_float4(float4 I, float4 N, float ETA) {
diff --git a/clang/test/CodeGenHLSL/builtins/smoothstep.hlsl b/clang/test/CodeGenHLSL/builtins/smoothstep.hlsl
index dcf9013045c07..c39261cfb7915 100644
--- a/clang/test/CodeGenHLSL/builtins/smoothstep.hlsl
+++ b/clang/test/CodeGenHLSL/builtins/smoothstep.hlsl
@@ -22,7 +22,7 @@
// SPVCHECK-LABEL: define hidden spir_func noundef nofpclass(nan inf) half @_Z20test_smoothstep_halfDhDhDh(
// SPVCHECK-SAME: half noundef nofpclass(nan inf) [[MIN:%.*]], half noundef nofpclass(nan inf) [[MAX:%.*]], half noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {
// SPVCHECK-NEXT: [[ENTRY:.*:]]
-// SPVCHECK-NEXT: [[SPV_SMOOTHSTEP_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef half @llvm.spv.smoothstep.f16(half nofpclass(nan inf) [[MIN]], half nofpclass(nan inf) [[MAX]], half nofpclass(nan inf) [[X]])
+// SPVCHECK-NEXT: [[SPV_SMOOTHSTEP_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) half @llvm.spv.smoothstep.f16(half nofpclass(nan inf) [[MIN]], half nofpclass(nan inf) [[MAX]], half nofpclass(nan inf) [[X]])
// SPVCHECK-NEXT: ret half [[SPV_SMOOTHSTEP_I]]
//
half test_smoothstep_half(half Min, half Max, half X) { return smoothstep(Min, Max, X); }
@@ -43,7 +43,7 @@ half test_smoothstep_half(half Min, half Max, half X) { return smoothstep(Min, M
// SPVCHECK-LABEL: define hidden spir_func noundef nofpclass(nan inf) <2 x half> @_Z21test_smoothstep_half2Dv2_DhS_S_(
// SPVCHECK-SAME: <2 x half> noundef nofpclass(nan inf) [[MIN:%.*]], <2 x half> noundef nofpclass(nan inf) [[MAX:%.*]], <2 x half> noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR0]] {
// SPVCHECK-NEXT: [[ENTRY:.*:]]
-// SPVCHECK-NEXT: [[SPV_SMOOTHSTEP_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef <2 x half> @llvm.spv.smoothstep.v2f16(<2 x half> nofpclass(nan inf) [[MIN]], <2 x half> nofpclass(nan inf) [[MAX]], <2 x half> nofpclass(nan inf) [[X]])
+// SPVCHECK-NEXT: [[SPV_SMOOTHSTEP_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) <2 x half> @llvm.spv.smoothstep.v2f16(<2 x half> nofpclass(nan inf) [[MIN]], <2 x half> nofpclass(nan inf) [[MAX]], <2 x half> nofpclass(nan inf) [[X]])
// SPVCHECK-NEXT: ret <2 x half> [[SPV_SMOOTHSTEP_I]]
//
half2 test_smoothstep_half2(half2 Min, half2 Max, half2 X) { return smoothstep(Min, Max, X); }
@@ -64,7 +64,7 @@ half2 test_smoothstep_half2(half2 Min, half2 Max, half2 X) { return smoothstep(M
// SPVCHECK-LABEL: define hidden spir_func noundef nofpclass(nan inf) <3 x half> @_Z21test_smoothstep_half3Dv3_DhS_S_(
// SPVCHECK-SAME: <3 x half> noundef nofpclass(nan inf) [[MIN:%.*]], <3 x half> noundef nofpclass(nan inf) [[MAX:%.*]], <3 x half> noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR0]] {
// SPVCHECK-NEXT: [[ENTRY:.*:]]
-// SPVCHECK-NEXT: [[SPV_SMOOTHSTEP_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef <3 x half> @llvm.spv.smoothstep.v3f16(<3 x half> nofpclass(nan inf) [[MIN]], <3 x half> nofpclass(nan inf) [[MAX]], <3 x half> nofpclass(nan inf) [[X]])
+// SPVCHECK-NEXT: [[SPV_SMOOTHSTEP_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) <3 x half> @llvm.spv.smoothstep.v3f16(<3 x half> nofpclass(nan inf) [[MIN]], <3 x half> nofpclass(nan inf) [[MAX]], <3 x half> nofpclass(nan inf) [[X]])
// SPVCHECK-NEXT: ret <3 x half> [[SPV_SMOOTHSTEP_I]]
//
half3 test_smoothstep_half3(half3 Min, half3 Max, half3 X) { return smoothstep(Min, Max, X); }
@@ -85,7 +85,7 @@ half3 test_smoothstep_half3(half3 Min, half3 Max, half3 X) { return smoothstep(M
// SPVCHECK-LABEL: define hidden spir_func noundef nofpclass(nan inf) <4 x half> @_Z21test_smoothstep_half4Dv4_DhS_S_(
// SPVCHECK-SAME: <4 x half> noundef nofpclass(nan inf) [[MIN:%.*]], <4 x half> noundef nofpclass(nan inf) [[MAX:%.*]], <4 x half> noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR0]] {
// SPVCHECK-NEXT: [[ENTRY:.*:]]
-// SPVCHECK-NEXT: [[SPV_SMOOTHSTEP_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef <4 x half> @llvm.spv.smoothstep.v4f16(<4 x half> nofpclass(nan inf) [[MIN]], <4 x half> nofpclass(nan inf) [[MAX]], <4 x half> nofpclass(nan inf) [[X]])
+// SPVCHECK-NEXT: [[SPV_SMOOTHSTEP_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) <4 x half> @llvm.spv.smoothstep.v4f16(<4 x half> nofpclass(nan inf) [[MIN]], <4 x half> nofpclass(nan inf) [[MAX]], <4 x half> nofpclass(nan inf) [[X]])
// SPVCHECK-NEXT: ret <4 x half> [[SPV_SMOOTHSTEP_I]]
//
half4 test_smoothstep_half4(half4 Min, half4 Max, half4 X) { return smoothstep(Min, Max, X); }
@@ -106,7 +106,7 @@ half4 test_smoothstep_half4(half4 Min, half4 Max, half4 X) { return smoothstep(M
// SPVCHECK-LABEL: define hidden spir_func noundef nofpclass(nan inf) float @_Z21test_smoothstep_floatfff(
// SPVCHECK-SAME: float noundef nofpclass(nan inf) [[MIN:%.*]], float noundef nofpclass(nan inf) [[MAX:%.*]], float noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR0]] {
// SPVCHECK-NEXT: [[ENTRY:.*:]]
-// SPVCHECK-NEXT: [[SPV_SMOOTHSTEP_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef float @llvm.spv.smoothstep.f32(float nofpclass(nan inf) [[MIN]], float nofpclass(nan inf) [[MAX]], float nofpclass(nan inf) [[X]])
+// SPVCHECK-NEXT: [[SPV_SMOOTHSTEP_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) float @llvm.spv.smoothstep.f32(float nofpclass(nan inf) [[MIN]], float nofpclass(nan inf) [[MAX]], float nofpclass(nan inf) [[X]])
// SPVCHECK-NEXT: ret float [[SPV_SMOOTHSTEP_I]]
//
float test_smoothstep_float(float Min, float Max, float X) { return smoothstep(Min, Max, X); }
@@ -127,7 +127,7 @@ float test_smoothstep_float(float Min, float Max, float X) { return smoothstep(M
// SPVCHECK-LABEL: define hidden spir_func noundef nofpclass(nan inf) <2 x float> @_Z22test_smoothstep_float2Dv2_fS_S_(
// SPVCHECK-SAME: <2 x float> noundef nofpclass(nan inf) [[MIN:%.*]], <2 x float> noundef nofpclass(nan inf) [[MAX:%.*]], <2 x float> noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR0]] {
// SPVCHECK-NEXT: [[ENTRY:.*:]]
-// SPVCHECK-NEXT: [[SPV_SMOOTHSTEP_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef <2 x float> @llvm.spv.smoothstep.v2f32(<2 x float> nofpclass(nan inf) [[MIN]], <2 x float> nofpclass(nan inf) [[MAX]], <2 x float> nofpclass(nan inf) [[X]])
+// SPVCHECK-NEXT: [[SPV_SMOOTHSTEP_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) <2 x float> @llvm.spv.smoothstep.v2f32(<2 x float> nofpclass(nan inf) [[MIN]], <2 x float> nofpclass(nan inf) [[MAX]], <2 x float> nofpclass(nan inf) [[X]])
// SPVCHECK-NEXT: ret <2 x float> [[SPV_SMOOTHSTEP_I]]
//
float2 test_smoothstep_float2(float2 Min, float2 Max, float2 X) { return smoothstep(Min, Max, X); }
@@ -148,7 +148,7 @@ float2 test_smoothstep_float2(float2 Min, float2 Max, float2 X) { return smooths
// SPVCHECK-LABEL: define hidden spir_func noundef nofpclass(nan inf) <3 x float> @_Z22test_smoothstep_float3Dv3_fS_S_(
// SPVCHECK-SAME: <3 x float> noundef nofpclass(nan inf) [[MIN:%.*]], <3 x float> noundef nofpclass(nan inf) [[MAX:%.*]], <3 x float> noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR0]] {
// SPVCHECK-NEXT: [[ENTRY:.*:]]
-// SPVCHECK-NEXT: [[SPV_SMOOTHSTEP_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef <3 x float> @llvm.spv.smoothstep.v3f32(<3 x float> nofpclass(nan inf) [[MIN]], <3 x float> nofpclass(nan inf) [[MAX]], <3 x float> nofpclass(nan inf) [[X]])
+// SPVCHECK-NEXT: [[SPV_SMOOTHSTEP_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) <3 x float> @llvm.spv.smoothstep.v3f32(<3 x float> nofpclass(nan inf) [[MIN]], <3 x float> nofpclass(nan inf) [[MAX]], <3 x float> nofpclass(nan inf) [[X]])
// SPVCHECK-NEXT: ret <3 x float> [[SPV_SMOOTHSTEP_I]]
//
float3 test_smoothstep_float3(float3 Min, float3 Max, float3 X) { return smoothstep(Min, Max, X); }
@@ -169,7 +169,7 @@ float3 test_smoothstep_float3(float3 Min, float3 Max, float3 X) { return smooths
// SPVCHECK-LABEL: define hidden spir_func noundef nofpclass(nan inf) <4 x float> @_Z22test_smoothstep_float4Dv4_fS_S_(
// SPVCHECK-SAME: <4 x float> noundef nofpclass(nan inf) [[MIN:%.*]], <4 x float> noundef nofpclass(nan inf) [[MAX:%.*]], <4 x float> noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR0]] {
// SPVCHECK-NEXT: [[ENTRY:.*:]]
-// SPVCHECK-NEXT: [[SPV_SMOOTHSTEP_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef <4 x float> @llvm.spv.smoothstep.v4f32(<4 x float> nofpclass(nan inf) [[MIN]], <4 x float> nofpclass(nan inf) [[MAX]], <4 x float> nofpclass(nan inf) [[X]])
+// SPVCHECK-NEXT: [[SPV_SMOOTHSTEP_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef nofpclass(nan inf) <4 x float> @llvm.spv.smoothstep.v4f32(<4 x float> nofpclass(nan inf) [[MIN]], <4 x float> nofpclass(nan inf) [[MAX]], <4 x float> nofpclass(nan inf) [[X]])
// SPVCHECK-NEXT: ret <4 x float> [[SPV_SMOOTHSTEP_I]]
//
float4 test_smoothstep_float4(float4 Min, float4 Max, float4 X) { return smoothstep(Min, Max, X); }
diff --git a/clang/test/Headers/__clang_hip_cmath.hip b/clang/test/Headers/__clang_hip_cmath.hip
index 2e0b776c4bfe5..5bcd5857b3919 100644
--- a/clang/test/Headers/__clang_hip_cmath.hip
+++ b/clang/test/Headers/__clang_hip_cmath.hip
@@ -33,7 +33,7 @@
//
// FINITEONLY-LABEL: @test_fma_f16(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef half @llvm.fma.f16(half nofpclass(nan inf) [[X:%.*]], half nofpclass(nan inf) [[Y:%.*]], half nofpclass(nan inf) [[Z:%.*]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) half @llvm.fma.f16(half nofpclass(nan inf) [[X:%.*]], half nofpclass(nan inf) [[Y:%.*]], half nofpclass(nan inf) [[Z:%.*]])
// FINITEONLY-NEXT: ret half [[TMP0]]
//
extern "C" __device__ _Float16 test_fma_f16(_Float16 x, _Float16 y,
@@ -62,7 +62,7 @@ extern "C" __device__ _Float16 test_pow_f16(_Float16 x, int y) {
//
// FINITEONLY-LABEL: @test_fabs_f32(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.fabs.f32(float nofpclass(nan inf) [[X:%.*]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @llvm.fabs.f32(float nofpclass(nan inf) [[X:%.*]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
extern "C" __device__ float test_fabs_f32(float x) {
@@ -159,12 +159,12 @@ namespace user_namespace {
//
// FINITEONLY-LABEL: @test_sqrt_f32(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.sqrt.f32(float nofpclass(nan inf) [[X:%.*]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @llvm.sqrt.f32(float nofpclass(nan inf) [[X:%.*]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// NO-CORRECT-DIV-SQRT-LABEL: @test_sqrt_f32(
// NO-CORRECT-DIV-SQRT-NEXT: entry:
-// NO-CORRECT-DIV-SQRT-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.sqrt.f32(float [[X:%.*]]), !fpmath [[META4:![0-9]+]]
+// NO-CORRECT-DIV-SQRT-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.sqrt.f32(float [[X:%.*]]), !fpmath [[META8:![0-9]+]]
// NO-CORRECT-DIV-SQRT-NEXT: ret float [[TMP0]]
//
extern "C" __device__ float test_sqrt_f32(float x) {
@@ -178,7 +178,7 @@ extern "C" __device__ float test_sqrt_f32(float x) {
//
// FINITEONLY-LABEL: @test_sqrt_f64(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef double @llvm.sqrt.f64(double nofpclass(nan inf) [[X:%.*]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @llvm.sqrt.f64(double nofpclass(nan inf) [[X:%.*]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
extern "C" __device__ double test_sqrt_f64(double x) {
diff --git a/clang/test/Headers/__clang_hip_math.hip b/clang/test/Headers/__clang_hip_math.hip
index e91f5723a8460..0270fb4cf74dd 100644
--- a/clang/test/Headers/__clang_hip_math.hip
+++ b/clang/test/Headers/__clang_hip_math.hip
@@ -1000,7 +1000,7 @@ extern "C" __device__ double test_cbrt(double x) {
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) float @test_ceilf(
// FINITEONLY-SAME: float noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.ceil.f32(float nofpclass(nan inf) [[X]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @llvm.ceil.f32(float nofpclass(nan inf) [[X]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef float @test_ceilf(
@@ -1034,7 +1034,7 @@ extern "C" __device__ float test_ceilf(float x) {
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) double @test_ceil(
// FINITEONLY-SAME: double noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef double @llvm.ceil.f64(double nofpclass(nan inf) [[X]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @llvm.ceil.f64(double nofpclass(nan inf) [[X]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef double @test_ceil(
@@ -1068,7 +1068,7 @@ extern "C" __device__ double test_ceil(double x) {
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) float @test_copysignf(
// FINITEONLY-SAME: float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.copysign.f32(float nofpclass(nan inf) [[X]], float nofpclass(nan inf) [[Y]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @llvm.copysign.f32(float nofpclass(nan inf) [[X]], float nofpclass(nan inf) [[Y]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef float @test_copysignf(
@@ -1102,7 +1102,7 @@ extern "C" __device__ float test_copysignf(float x, float y) {
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) double @test_copysign(
// FINITEONLY-SAME: double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef double @llvm.copysign.f64(double nofpclass(nan inf) [[X]], double nofpclass(nan inf) [[Y]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @llvm.copysign.f64(double nofpclass(nan inf) [[X]], double nofpclass(nan inf) [[Y]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef double @test_copysign(
@@ -1612,7 +1612,7 @@ extern "C" __device__ double test_erfinv(double x) {
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) float @test_exp10f(
// FINITEONLY-SAME: float noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.exp10.f32(float nofpclass(nan inf) [[X]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @llvm.exp10.f32(float nofpclass(nan inf) [[X]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef float @test_exp10f(
@@ -1680,7 +1680,7 @@ extern "C" __device__ double test_exp10(double x) {
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) float @test_exp2f(
// FINITEONLY-SAME: float noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.exp2.f32(float nofpclass(nan inf) [[X]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @llvm.exp2.f32(float nofpclass(nan inf) [[X]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef float @test_exp2f(
@@ -1748,7 +1748,7 @@ extern "C" __device__ double test_exp2(double x) {
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) float @test_expf(
// FINITEONLY-SAME: float noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.exp.f32(float nofpclass(nan inf) [[X]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @llvm.exp.f32(float nofpclass(nan inf) [[X]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef float @test_expf(
@@ -1884,7 +1884,7 @@ extern "C" __device__ double test_expm1(double x) {
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) float @test_fabsf(
// FINITEONLY-SAME: float noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.fabs.f32(float nofpclass(nan inf) [[X]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @llvm.fabs.f32(float nofpclass(nan inf) [[X]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef float @test_fabsf(
@@ -1918,7 +1918,7 @@ extern "C" __device__ float test_fabsf(float x) {
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) double @test_fabs(
// FINITEONLY-SAME: double noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef double @llvm.fabs.f64(double nofpclass(nan inf) [[X]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @llvm.fabs.f64(double nofpclass(nan inf) [[X]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef double @test_fabs(
@@ -2054,7 +2054,7 @@ extern "C" __device__ float test_fdividef(float x, float y) {
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) float @test_floorf(
// FINITEONLY-SAME: float noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.floor.f32(float nofpclass(nan inf) [[X]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @llvm.floor.f32(float nofpclass(nan inf) [[X]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef float @test_floorf(
@@ -2088,7 +2088,7 @@ extern "C" __device__ float test_floorf(float x) {
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) double @test_floor(
// FINITEONLY-SAME: double noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef double @llvm.floor.f64(double nofpclass(nan inf) [[X]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @llvm.floor.f64(double nofpclass(nan inf) [[X]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef double @test_floor(
@@ -2122,7 +2122,7 @@ extern "C" __device__ double test_floor(double x) {
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) float @test_fmaf(
// FINITEONLY-SAME: float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]], float noundef nofpclass(nan inf) [[Z:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.fma.f32(float nofpclass(nan inf) [[X]], float nofpclass(nan inf) [[Y]], float nofpclass(nan inf) [[Z]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @llvm.fma.f32(float nofpclass(nan inf) [[X]], float nofpclass(nan inf) [[Y]], float nofpclass(nan inf) [[Z]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef float @test_fmaf(
@@ -2156,7 +2156,7 @@ extern "C" __device__ float test_fmaf(float x, float y, float z) {
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) double @test_fma(
// FINITEONLY-SAME: double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]], double noundef nofpclass(nan inf) [[Z:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef double @llvm.fma.f64(double nofpclass(nan inf) [[X]], double nofpclass(nan inf) [[Y]], double nofpclass(nan inf) [[Z]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @llvm.fma.f64(double nofpclass(nan inf) [[X]], double nofpclass(nan inf) [[Y]], double nofpclass(nan inf) [[Z]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef double @test_fma(
@@ -2190,7 +2190,7 @@ extern "C" __device__ double test_fma(double x, double y, double z) {
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) double @test_fma_rn(
// FINITEONLY-SAME: double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]], double noundef nofpclass(nan inf) [[Z:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef double @llvm.fma.f64(double nofpclass(nan inf) [[X]], double nofpclass(nan inf) [[Y]], double nofpclass(nan inf) [[Z]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @llvm.fma.f64(double nofpclass(nan inf) [[X]], double nofpclass(nan inf) [[Y]], double nofpclass(nan inf) [[Z]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef double @test_fma_rn(
@@ -2224,7 +2224,7 @@ extern "C" __device__ double test_fma_rn(double x, double y, double z) {
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) float @test_fmaxf(
// FINITEONLY-SAME: float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.maxnum.f32(float nofpclass(nan inf) [[X]], float nofpclass(nan inf) [[Y]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @llvm.maxnum.f32(float nofpclass(nan inf) [[X]], float nofpclass(nan inf) [[Y]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef float @test_fmaxf(
@@ -2258,7 +2258,7 @@ extern "C" __device__ float test_fmaxf(float x, float y) {
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) double @test_fmax(
// FINITEONLY-SAME: double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef double @llvm.maxnum.f64(double nofpclass(nan inf) [[X]], double nofpclass(nan inf) [[Y]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @llvm.maxnum.f64(double nofpclass(nan inf) [[X]], double nofpclass(nan inf) [[Y]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef double @test_fmax(
@@ -2292,7 +2292,7 @@ extern "C" __device__ double test_fmax(double x, double y) {
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) float @test_fminf(
// FINITEONLY-SAME: float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.minnum.f32(float nofpclass(nan inf) [[X]], float nofpclass(nan inf) [[Y]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @llvm.minnum.f32(float nofpclass(nan inf) [[X]], float nofpclass(nan inf) [[Y]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef float @test_fminf(
@@ -2326,7 +2326,7 @@ extern "C" __device__ float test_fminf(float x, float y) {
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) double @test_fmin(
// FINITEONLY-SAME: double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef double @llvm.minnum.f64(double nofpclass(nan inf) [[X]], double nofpclass(nan inf) [[Y]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @llvm.minnum.f64(double nofpclass(nan inf) [[X]], double nofpclass(nan inf) [[Y]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef double @test_fmin(
@@ -3384,7 +3384,7 @@ extern "C" __device__ double test_jn(int x, double y) {
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) float @test_ldexpf(
// FINITEONLY-SAME: float noundef nofpclass(nan inf) [[X:%.*]], i32 noundef [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.ldexp.f32.i32(float nofpclass(nan inf) [[X]], i32 [[Y]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @llvm.ldexp.f32.i32(float nofpclass(nan inf) [[X]], i32 [[Y]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef float @test_ldexpf(
@@ -3418,7 +3418,7 @@ extern "C" __device__ float test_ldexpf(float x, int y) {
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) double @test_ldexp(
// FINITEONLY-SAME: double noundef nofpclass(nan inf) [[X:%.*]], i32 noundef [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef double @llvm.ldexp.f64.i32(double nofpclass(nan inf) [[X]], i32 [[Y]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @llvm.ldexp.f64.i32(double nofpclass(nan inf) [[X]], i32 [[Y]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef double @test_ldexp(
@@ -3676,7 +3676,7 @@ extern "C" __device__ long long int test_llround(double x) {
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) float @test_log10f(
// FINITEONLY-SAME: float noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.log10.f32(float nofpclass(nan inf) [[X]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @llvm.log10.f32(float nofpclass(nan inf) [[X]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef float @test_log10f(
@@ -3812,7 +3812,7 @@ extern "C" __device__ double test_log1p(double x) {
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) float @test_log2f(
// FINITEONLY-SAME: float noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.log2.f32(float nofpclass(nan inf) [[X]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @llvm.log2.f32(float nofpclass(nan inf) [[X]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef float @test_log2f(
@@ -4018,7 +4018,7 @@ extern "C" __device__ double test_logb(double x) {
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) float @test_logf(
// FINITEONLY-SAME: float noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.log.f32(float nofpclass(nan inf) [[X]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @llvm.log.f32(float nofpclass(nan inf) [[X]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef float @test_logf(
@@ -5148,7 +5148,7 @@ extern "C" __device__ double test_nan_fill() {
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) float @test_nearbyintf(
// FINITEONLY-SAME: float noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.nearbyint.f32(float nofpclass(nan inf) [[X]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @llvm.nearbyint.f32(float nofpclass(nan inf) [[X]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef float @test_nearbyintf(
@@ -5182,7 +5182,7 @@ extern "C" __device__ float test_nearbyintf(float x) {
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) double @test_nearbyint(
// FINITEONLY-SAME: double noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef double @llvm.nearbyint.f64(double nofpclass(nan inf) [[X]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @llvm.nearbyint.f64(double nofpclass(nan inf) [[X]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef double @test_nearbyint(
@@ -5587,7 +5587,7 @@ extern "C" __device__ double test_normcdfinv(double x) {
// FINITEONLY-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I]], label %[[_ZL5NORMFIPKF_EXIT_LOOPEXIT:.*]], label %[[WHILE_BODY_I]], !llvm.loop [[LOOP20:![0-9]+]]
// FINITEONLY: [[_ZL5NORMFIPKF_EXIT_LOOPEXIT]]:
-// FINITEONLY-NEXT: [[TMP1:%.*]] = tail call nnan ninf contract float @llvm.sqrt.f32(float [[ADD_I]])
+// FINITEONLY-NEXT: [[TMP1:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @llvm.sqrt.f32(float [[ADD_I]])
// FINITEONLY-NEXT: br label %[[_ZL5NORMFIPKF_EXIT]]
// FINITEONLY: [[_ZL5NORMFIPKF_EXIT]]:
// FINITEONLY-NEXT: [[__R_0_I_LCSSA:%.*]] = phi float [ 0.000000e+00, %[[ENTRY]] ], [ [[TMP1]], %[[_ZL5NORMFIPKF_EXIT_LOOPEXIT]] ]
@@ -5706,7 +5706,7 @@ extern "C" __device__ float test_normf(int x, const float *y) {
// FINITEONLY-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[DEC_I]], 0
// FINITEONLY-NEXT: br i1 [[TOBOOL_NOT_I]], label %[[_ZL4NORMIPKD_EXIT_LOOPEXIT:.*]], label %[[WHILE_BODY_I]], !llvm.loop [[LOOP21:![0-9]+]]
// FINITEONLY: [[_ZL4NORMIPKD_EXIT_LOOPEXIT]]:
-// FINITEONLY-NEXT: [[TMP1:%.*]] = tail call nnan ninf contract double @llvm.sqrt.f64(double [[ADD_I]])
+// FINITEONLY-NEXT: [[TMP1:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @llvm.sqrt.f64(double [[ADD_I]])
// FINITEONLY-NEXT: br label %[[_ZL4NORMIPKD_EXIT]]
// FINITEONLY: [[_ZL4NORMIPKD_EXIT]]:
// FINITEONLY-NEXT: [[__R_0_I_LCSSA:%.*]] = phi double [ 0.000000e+00, %[[ENTRY]] ], [ [[TMP1]], %[[_ZL4NORMIPKD_EXIT_LOOPEXIT]] ]
@@ -6254,7 +6254,7 @@ extern "C" __device__ double test_rhypot(double x, double y) {
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) float @test_rintf(
// FINITEONLY-SAME: float noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.rint.f32(float nofpclass(nan inf) [[X]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @llvm.rint.f32(float nofpclass(nan inf) [[X]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef float @test_rintf(
@@ -6288,7 +6288,7 @@ extern "C" __device__ float test_rintf(float x) {
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) double @test_rint(
// FINITEONLY-SAME: double noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef double @llvm.rint.f64(double nofpclass(nan inf) [[X]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @llvm.rint.f64(double nofpclass(nan inf) [[X]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef double @test_rint(
@@ -6676,7 +6676,7 @@ extern "C" __device__ double test_rnorm4d(double x, double y, double z, double w
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) float @test_roundf(
// FINITEONLY-SAME: float noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.round.f32(float nofpclass(nan inf) [[X]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @llvm.round.f32(float nofpclass(nan inf) [[X]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef float @test_roundf(
@@ -6710,7 +6710,7 @@ extern "C" __device__ float test_roundf(float x) {
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) double @test_round(
// FINITEONLY-SAME: double noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef double @llvm.round.f64(double nofpclass(nan inf) [[X]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @llvm.round.f64(double nofpclass(nan inf) [[X]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef double @test_round(
@@ -6816,7 +6816,7 @@ extern "C" __device__ double test_rsqrt(double x) {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
// FINITEONLY-NEXT: [[SPEC_STORE_SELECT_I:%.*]] = tail call i64 @llvm.smax.i64(i64 [[Y]], i64 -2147483648)
// FINITEONLY-NEXT: [[CONV_I:%.*]] = trunc i64 [[SPEC_STORE_SELECT_I]] to i32
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.ldexp.f32.i32(float nofpclass(nan inf) [[X]], i32 [[CONV_I]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @llvm.ldexp.f32.i32(float nofpclass(nan inf) [[X]], i32 [[CONV_I]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef float @test_scalblnf(
@@ -6860,7 +6860,7 @@ extern "C" __device__ float test_scalblnf(float x, long int y) {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
// FINITEONLY-NEXT: [[SPEC_STORE_SELECT_I:%.*]] = tail call i64 @llvm.smax.i64(i64 [[Y]], i64 -2147483648)
// FINITEONLY-NEXT: [[CONV_I:%.*]] = trunc i64 [[SPEC_STORE_SELECT_I]] to i32
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef double @llvm.ldexp.f64.i32(double nofpclass(nan inf) [[X]], i32 [[CONV_I]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @llvm.ldexp.f64.i32(double nofpclass(nan inf) [[X]], i32 [[CONV_I]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef double @test_scalbln(
@@ -6900,7 +6900,7 @@ extern "C" __device__ double test_scalbln(double x, long int y) {
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) float @test_scalbnf(
// FINITEONLY-SAME: float noundef nofpclass(nan inf) [[X:%.*]], i32 noundef [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.ldexp.f32.i32(float nofpclass(nan inf) [[X]], i32 [[Y]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @llvm.ldexp.f32.i32(float nofpclass(nan inf) [[X]], i32 [[Y]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef float @test_scalbnf(
@@ -6934,7 +6934,7 @@ extern "C" __device__ float test_scalbnf(float x, int y) {
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) double @test_scalbn(
// FINITEONLY-SAME: double noundef nofpclass(nan inf) [[X:%.*]], i32 noundef [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef double @llvm.ldexp.f64.i32(double nofpclass(nan inf) [[X]], i32 [[Y]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @llvm.ldexp.f64.i32(double nofpclass(nan inf) [[X]], i32 [[Y]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef double @test_scalbn(
@@ -7447,7 +7447,7 @@ extern "C" __device__ double test_sinpi(double x) {
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) float @test_sqrtf(
// FINITEONLY-SAME: float noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.sqrt.f32(float nofpclass(nan inf) [[X]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @llvm.sqrt.f32(float nofpclass(nan inf) [[X]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef float @test_sqrtf(
@@ -7481,7 +7481,7 @@ extern "C" __device__ float test_sqrtf(float x) {
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) double @test_sqrt(
// FINITEONLY-SAME: double noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef double @llvm.sqrt.f64(double nofpclass(nan inf) [[X]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @llvm.sqrt.f64(double nofpclass(nan inf) [[X]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef double @test_sqrt(
@@ -7719,7 +7719,7 @@ extern "C" __device__ double test_tgamma(double x) {
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) float @test_truncf(
// FINITEONLY-SAME: float noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.trunc.f32(float nofpclass(nan inf) [[X]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @llvm.trunc.f32(float nofpclass(nan inf) [[X]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef float @test_truncf(
@@ -7753,7 +7753,7 @@ extern "C" __device__ float test_truncf(float x) {
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) double @test_trunc(
// FINITEONLY-SAME: double noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef double @llvm.trunc.f64(double nofpclass(nan inf) [[X]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @llvm.trunc.f64(double nofpclass(nan inf) [[X]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef double @test_trunc(
@@ -8307,7 +8307,7 @@ extern "C" __device__ float test___cosf(float x) {
// FINITEONLY-SAME: float noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
// FINITEONLY-NEXT: [[MUL_I:%.*]] = fmul nnan ninf contract float [[X]], 0x400A934F00000000
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.amdgcn.exp2.f32(float [[MUL_I]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @llvm.amdgcn.exp2.f32(float [[MUL_I]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef float @test___exp10f(
@@ -8346,7 +8346,7 @@ extern "C" __device__ float test___exp10f(float x) {
// FINITEONLY-SAME: float noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
// FINITEONLY-NEXT: [[MUL_I:%.*]] = fmul nnan ninf contract float [[X]], 0x3FF7154760000000
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.amdgcn.exp2.f32(float [[MUL_I]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @llvm.amdgcn.exp2.f32(float [[MUL_I]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef float @test___expf(
@@ -8451,7 +8451,7 @@ extern "C" __device__ float test___fdividef(float x, float y) {
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) float @test__fmaf_rn(
// FINITEONLY-SAME: float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]], float noundef nofpclass(nan inf) [[Z:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.fma.f32(float nofpclass(nan inf) [[X]], float nofpclass(nan inf) [[Y]], float nofpclass(nan inf) [[Z]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @llvm.fma.f32(float nofpclass(nan inf) [[X]], float nofpclass(nan inf) [[Y]], float nofpclass(nan inf) [[Z]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef float @test__fmaf_rn(
@@ -8553,7 +8553,7 @@ extern "C" __device__ float test___frcp_rn(float x) {
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) float @test___frsqrt_rn(
// FINITEONLY-SAME: float noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.amdgcn.rsq.f32(float nofpclass(nan inf) [[X]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @llvm.amdgcn.rsq.f32(float nofpclass(nan inf) [[X]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef float @test___frsqrt_rn(
@@ -8655,7 +8655,7 @@ extern "C" __device__ float test___fsub_rn(float x, float y) {
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) float @test___log10f(
// FINITEONLY-SAME: float noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.log10.f32(float nofpclass(nan inf) [[X]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @llvm.log10.f32(float nofpclass(nan inf) [[X]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef float @test___log10f(
@@ -8689,7 +8689,7 @@ extern "C" __device__ float test___log10f(float x) {
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) float @test___log2f(
// FINITEONLY-SAME: float noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.amdgcn.log.f32(float nofpclass(nan inf) [[X]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @llvm.amdgcn.log.f32(float nofpclass(nan inf) [[X]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef float @test___log2f(
@@ -8723,7 +8723,7 @@ extern "C" __device__ float test___log2f(float x) {
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) float @test___logf(
// FINITEONLY-SAME: float noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.log.f32(float nofpclass(nan inf) [[X]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @llvm.log.f32(float nofpclass(nan inf) [[X]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef float @test___logf(
@@ -9108,7 +9108,7 @@ extern "C" __device__ double test___drcp_rn(double x) {
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) double @test___dsqrt_rn(
// FINITEONLY-SAME: double noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef double @llvm.sqrt.f64(double nofpclass(nan inf) [[X]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @llvm.sqrt.f64(double nofpclass(nan inf) [[X]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef double @test___dsqrt_rn(
@@ -9142,7 +9142,7 @@ extern "C" __device__ double test___dsqrt_rn(double x) {
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) double @test__fma_rn(
// FINITEONLY-SAME: double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]], double noundef nofpclass(nan inf) [[Z:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef double @llvm.fma.f64(double nofpclass(nan inf) [[X]], double nofpclass(nan inf) [[Y]], double nofpclass(nan inf) [[Z]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @llvm.fma.f64(double nofpclass(nan inf) [[X]], double nofpclass(nan inf) [[Y]], double nofpclass(nan inf) [[Z]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef double @test__fma_rn(
@@ -9176,7 +9176,7 @@ extern "C" __device__ double test__fma_rn(double x, double y, double z) {
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) float @test_float_min(
// FINITEONLY-SAME: float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.minnum.f32(float nofpclass(nan inf) [[X]], float nofpclass(nan inf) [[Y]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @llvm.minnum.f32(float nofpclass(nan inf) [[X]], float nofpclass(nan inf) [[Y]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef float @test_float_min(
@@ -9210,7 +9210,7 @@ extern "C" __device__ float test_float_min(float x, float y) {
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) float @test_float_max(
// FINITEONLY-SAME: float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.maxnum.f32(float nofpclass(nan inf) [[X]], float nofpclass(nan inf) [[Y]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @llvm.maxnum.f32(float nofpclass(nan inf) [[X]], float nofpclass(nan inf) [[Y]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef float @test_float_max(
@@ -9244,7 +9244,7 @@ extern "C" __device__ float test_float_max(float x, float y) {
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) double @test_double_min(
// FINITEONLY-SAME: double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef double @llvm.minnum.f64(double nofpclass(nan inf) [[X]], double nofpclass(nan inf) [[Y]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @llvm.minnum.f64(double nofpclass(nan inf) [[X]], double nofpclass(nan inf) [[Y]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef double @test_double_min(
@@ -9278,7 +9278,7 @@ extern "C" __device__ double test_double_min(double x, double y) {
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) double @test_double_max(
// FINITEONLY-SAME: double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef double @llvm.maxnum.f64(double nofpclass(nan inf) [[X]], double nofpclass(nan inf) [[Y]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @llvm.maxnum.f64(double nofpclass(nan inf) [[X]], double nofpclass(nan inf) [[Y]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef double @test_double_max(
More information about the llvm-commits
mailing list