[clang] 3f38cd0 - Revert "Inline: Propagate callsite nofpclass attribute"
Matt Arsenault via cfe-commits
cfe-commits at lists.llvm.org
Tue Apr 8 09:15:25 PDT 2025
Author: Matt Arsenault
Date: 2025-04-08T23:15:00+07:00
New Revision: 3f38cd07d820248fd2043efb1341fabaac2d84a6
URL: https://github.com/llvm/llvm-project/commit/3f38cd07d820248fd2043efb1341fabaac2d84a6
DIFF: https://github.com/llvm/llvm-project/commit/3f38cd07d820248fd2043efb1341fabaac2d84a6.diff
LOG: Revert "Inline: Propagate callsite nofpclass attribute"
This reverts commit b0cb672b9968eeee6eb022e98476957dbdf8e6e2.
Breaks bot
Added:
Modified:
clang/test/CodeGenHLSL/builtins/distance.hlsl
clang/test/CodeGenHLSL/builtins/length.hlsl
clang/test/CodeGenHLSL/builtins/reflect.hlsl
clang/test/CodeGenHLSL/builtins/smoothstep.hlsl
clang/test/Headers/__clang_hip_cmath.hip
clang/test/Headers/__clang_hip_math.hip
llvm/lib/Transforms/Utils/InlineFunction.cpp
llvm/test/Transforms/Inline/access-attributes-prop.ll
Removed:
################################################################################
diff --git a/clang/test/CodeGenHLSL/builtins/distance.hlsl b/clang/test/CodeGenHLSL/builtins/distance.hlsl
index ac38cf1853799..e830903261c8c 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 half @llvm.fabs.f16(half [[SUB_I]])
// CHECK-NEXT: ret half [[ELT_ABS_I]]
//
// SPVCHECK-LABEL: define 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 half @llvm.fabs.f16(half [[SUB_I]])
// SPVCHECK-NEXT: ret half [[ELT_ABS_I]]
//
half test_distance_half(half X, half Y) { return distance(X, Y); }
@@ -26,7 +26,7 @@ half test_distance_half(half X, half Y) { return distance(X, Y); }
// CHECK-SAME: <2 x half> noundef nofpclass(nan inf) [[X:%.*]], <2 x half> noundef nofpclass(nan inf) [[Y:%.*]]) local_unnamed_addr #[[ATTR0]] {
// 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: [[HLSL_DOT_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn half @llvm.dx.fdot.v2f16(<2 x half> [[SUB_I]], <2 x half> [[SUB_I]])
// CHECK-NEXT: [[TMP0:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef half @llvm.sqrt.f16(half [[HLSL_DOT_I]])
// CHECK-NEXT: ret half [[TMP0]]
//
@@ -34,7 +34,7 @@ half test_distance_half(half X, half Y) { return distance(X, Y); }
// 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 half @llvm.spv.length.v2f16(<2 x half> [[SUB_I]])
// SPVCHECK-NEXT: ret half [[SPV_LENGTH_I]]
//
half test_distance_half2(half2 X, half2 Y) { return distance(X, Y); }
@@ -43,7 +43,7 @@ half test_distance_half2(half2 X, half2 Y) { return distance(X, Y); }
// CHECK-SAME: <3 x half> noundef nofpclass(nan inf) [[X:%.*]], <3 x half> noundef nofpclass(nan inf) [[Y:%.*]]) local_unnamed_addr #[[ATTR0]] {
// 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: [[HLSL_DOT_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn half @llvm.dx.fdot.v3f16(<3 x half> [[SUB_I]], <3 x half> [[SUB_I]])
// CHECK-NEXT: [[TMP0:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef half @llvm.sqrt.f16(half [[HLSL_DOT_I]])
// CHECK-NEXT: ret half [[TMP0]]
//
@@ -51,7 +51,7 @@ half test_distance_half2(half2 X, half2 Y) { return distance(X, Y); }
// 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 half @llvm.spv.length.v3f16(<3 x half> [[SUB_I]])
// SPVCHECK-NEXT: ret half [[SPV_LENGTH_I]]
//
half test_distance_half3(half3 X, half3 Y) { return distance(X, Y); }
@@ -60,7 +60,7 @@ half test_distance_half3(half3 X, half3 Y) { return distance(X, Y); }
// CHECK-SAME: <4 x half> noundef nofpclass(nan inf) [[X:%.*]], <4 x half> noundef nofpclass(nan inf) [[Y:%.*]]) local_unnamed_addr #[[ATTR0]] {
// 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: [[HLSL_DOT_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn half @llvm.dx.fdot.v4f16(<4 x half> [[SUB_I]], <4 x half> [[SUB_I]])
// CHECK-NEXT: [[TMP0:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef half @llvm.sqrt.f16(half [[HLSL_DOT_I]])
// CHECK-NEXT: ret half [[TMP0]]
//
@@ -68,7 +68,7 @@ half test_distance_half3(half3 X, half3 Y) { return distance(X, Y); }
// 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 half @llvm.spv.length.v4f16(<4 x half> [[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 float @llvm.fabs.f32(float [[SUB_I]])
// CHECK-NEXT: ret float [[ELT_ABS_I]]
//
// SPVCHECK-LABEL: define 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 float @llvm.fabs.f32(float [[SUB_I]])
// SPVCHECK-NEXT: ret float [[ELT_ABS_I]]
//
float test_distance_float(float X, float Y) { return distance(X, Y); }
@@ -93,7 +93,7 @@ float test_distance_float(float X, float Y) { return distance(X, Y); }
// CHECK-SAME: <2 x float> noundef nofpclass(nan inf) [[X:%.*]], <2 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 <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: [[HLSL_DOT_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn float @llvm.dx.fdot.v2f32(<2 x float> [[SUB_I]], <2 x float> [[SUB_I]])
// CHECK-NEXT: [[TMP0:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef float @llvm.sqrt.f32(float [[HLSL_DOT_I]])
// CHECK-NEXT: ret float [[TMP0]]
//
@@ -101,7 +101,7 @@ float test_distance_float(float X, float Y) { return distance(X, Y); }
// 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 float @llvm.spv.length.v2f32(<2 x float> [[SUB_I]])
// SPVCHECK-NEXT: ret float [[SPV_LENGTH_I]]
//
float test_distance_float2(float2 X, float2 Y) { return distance(X, Y); }
@@ -110,7 +110,7 @@ float test_distance_float2(float2 X, float2 Y) { return distance(X, Y); }
// CHECK-SAME: <3 x float> noundef nofpclass(nan inf) [[X:%.*]], <3 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 <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: [[HLSL_DOT_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn float @llvm.dx.fdot.v3f32(<3 x float> [[SUB_I]], <3 x float> [[SUB_I]])
// CHECK-NEXT: [[TMP0:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef float @llvm.sqrt.f32(float [[HLSL_DOT_I]])
// CHECK-NEXT: ret float [[TMP0]]
//
@@ -118,7 +118,7 @@ float test_distance_float2(float2 X, float2 Y) { return distance(X, Y); }
// 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 float @llvm.spv.length.v3f32(<3 x float> [[SUB_I]])
// SPVCHECK-NEXT: ret float [[SPV_LENGTH_I]]
//
float test_distance_float3(float3 X, float3 Y) { return distance(X, Y); }
@@ -127,7 +127,7 @@ float test_distance_float3(float3 X, float3 Y) { return distance(X, Y); }
// CHECK-SAME: <4 x float> noundef nofpclass(nan inf) [[X:%.*]], <4 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 <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: [[HLSL_DOT_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn float @llvm.dx.fdot.v4f32(<4 x float> [[SUB_I]], <4 x float> [[SUB_I]])
// CHECK-NEXT: [[TMP0:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef float @llvm.sqrt.f32(float [[HLSL_DOT_I]])
// CHECK-NEXT: ret float [[TMP0]]
//
@@ -135,7 +135,7 @@ float test_distance_float3(float3 X, float3 Y) { return distance(X, Y); }
// 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 float @llvm.spv.length.v4f32(<4 x float> [[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 0b17d03d7097d..2d4bbd995298f 100644
--- a/clang/test/CodeGenHLSL/builtins/length.hlsl
+++ b/clang/test/CodeGenHLSL/builtins/length.hlsl
@@ -14,13 +14,13 @@
// CHECK-LABEL: define 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 half @llvm.fabs.f16(half [[P0]])
// CHECK-NEXT: ret half [[ELT_ABS_I]]
//
// SPVCHECK-LABEL: define 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 half @llvm.fabs.f16(half [[P0]])
// SPVCHECK-NEXT: ret half [[ELT_ABS_I]]
//
half test_length_half(half p0)
@@ -35,14 +35,14 @@ half test_length_half(half p0)
// CHECK-LABEL: define noundef nofpclass(nan inf) half @_Z17test_length_half2Dv2_Dh(
// 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: [[HLSL_DOT_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn half @llvm.dx.fdot.v2f16(<2 x half> [[P0]], <2 x half> [[P0]])
// CHECK-NEXT: [[TMP0:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef half @llvm.sqrt.f16(half [[HLSL_DOT_I]])
// CHECK-NEXT: ret half [[TMP0]]
//
// SPVCHECK-LABEL: define 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 half @llvm.spv.length.v2f16(<2 x half> [[P0]])
// SPVCHECK-NEXT: ret half [[SPV_LENGTH_I]]
//
half test_length_half2(half2 p0)
@@ -54,14 +54,14 @@ half test_length_half2(half2 p0)
// CHECK-LABEL: define noundef nofpclass(nan inf) half @_Z17test_length_half3Dv3_Dh(
// 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: [[HLSL_DOT_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn half @llvm.dx.fdot.v3f16(<3 x half> [[P0]], <3 x half> [[P0]])
// CHECK-NEXT: [[TMP0:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef half @llvm.sqrt.f16(half [[HLSL_DOT_I]])
// CHECK-NEXT: ret half [[TMP0]]
//
// SPVCHECK-LABEL: define 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 half @llvm.spv.length.v3f16(<3 x half> [[P0]])
// SPVCHECK-NEXT: ret half [[SPV_LENGTH_I]]
//
half test_length_half3(half3 p0)
@@ -73,14 +73,14 @@ half test_length_half3(half3 p0)
// CHECK-LABEL: define noundef nofpclass(nan inf) half @_Z17test_length_half4Dv4_Dh(
// 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: [[HLSL_DOT_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn half @llvm.dx.fdot.v4f16(<4 x half> [[P0]], <4 x half> [[P0]])
// CHECK-NEXT: [[TMP0:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef half @llvm.sqrt.f16(half [[HLSL_DOT_I]])
// CHECK-NEXT: ret half [[TMP0]]
//
// SPVCHECK-LABEL: define 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 half @llvm.spv.length.v4f16(<4 x half> [[P0]])
// SPVCHECK-NEXT: ret half [[SPV_LENGTH_I]]
//
half test_length_half4(half4 p0)
@@ -92,13 +92,13 @@ half test_length_half4(half4 p0)
// CHECK-LABEL: define 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 float @llvm.fabs.f32(float [[P0]])
// CHECK-NEXT: ret float [[ELT_ABS_I]]
//
// SPVCHECK-LABEL: define 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 float @llvm.fabs.f32(float [[P0]])
// SPVCHECK-NEXT: ret float [[ELT_ABS_I]]
//
float test_length_float(float p0)
@@ -110,14 +110,14 @@ float test_length_float(float p0)
// CHECK-LABEL: define noundef nofpclass(nan inf) float @_Z18test_length_float2Dv2_f(
// 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: [[HLSL_DOT_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn float @llvm.dx.fdot.v2f32(<2 x float> [[P0]], <2 x float> [[P0]])
// CHECK-NEXT: [[TMP0:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef float @llvm.sqrt.f32(float [[HLSL_DOT_I]])
// CHECK-NEXT: ret float [[TMP0]]
//
// SPVCHECK-LABEL: define 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 float @llvm.spv.length.v2f32(<2 x float> [[P0]])
// SPVCHECK-NEXT: ret float [[SPV_LENGTH_I]]
//
float test_length_float2(float2 p0)
@@ -129,14 +129,14 @@ float test_length_float2(float2 p0)
// CHECK-LABEL: define noundef nofpclass(nan inf) float @_Z18test_length_float3Dv3_f(
// 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: [[HLSL_DOT_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn float @llvm.dx.fdot.v3f32(<3 x float> [[P0]], <3 x float> [[P0]])
// CHECK-NEXT: [[TMP0:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef float @llvm.sqrt.f32(float [[HLSL_DOT_I]])
// CHECK-NEXT: ret float [[TMP0]]
//
// SPVCHECK-LABEL: define 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 float @llvm.spv.length.v3f32(<3 x float> [[P0]])
// SPVCHECK-NEXT: ret float [[SPV_LENGTH_I]]
//
float test_length_float3(float3 p0)
@@ -148,14 +148,14 @@ float test_length_float3(float3 p0)
// CHECK-LABEL: define noundef nofpclass(nan inf) float @_Z18test_length_float4Dv4_f(
// 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: [[HLSL_DOT_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn float @llvm.dx.fdot.v4f32(<4 x float> [[P0]], <4 x float> [[P0]])
// CHECK-NEXT: [[TMP0:%.*]] = tail call reassoc nnan ninf nsz arcp afn noundef float @llvm.sqrt.f32(float [[HLSL_DOT_I]])
// CHECK-NEXT: ret float [[TMP0]]
//
// SPVCHECK-LABEL: define 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 float @llvm.spv.length.v4f32(<4 x float> [[P0]])
// SPVCHECK-NEXT: ret float [[SPV_LENGTH_I]]
//
float test_length_float4(float4 p0)
diff --git a/clang/test/CodeGenHLSL/builtins/reflect.hlsl b/clang/test/CodeGenHLSL/builtins/reflect.hlsl
index c082e63ac1da6..35ee059697c4b 100644
--- a/clang/test/CodeGenHLSL/builtins/reflect.hlsl
+++ b/clang/test/CodeGenHLSL/builtins/reflect.hlsl
@@ -31,7 +31,7 @@ half test_reflect_half(half I, half N) {
// CHECK-LABEL: define noundef nofpclass(nan inf) <2 x half> @_Z18test_reflect_half2Dv2_DhS_(
// CHECK-SAME: <2 x half> noundef nofpclass(nan inf) [[I:%.*]], <2 x half> noundef nofpclass(nan inf) [[N:%.*]]) 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) [[I]], <2 x half> nofpclass(nan inf) [[N]])
+// CHECK-NEXT: [[HLSL_DOT_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn half @llvm.dx.fdot.v2f16(<2 x half> [[I]], <2 x half> [[N]])
// CHECK-NEXT: [[DOTSCALAR:%.*]] = fmul reassoc nnan ninf nsz arcp afn half [[HLSL_DOT_I]], 0xH4000
// CHECK-NEXT: [[TMP0:%.*]] = insertelement <2 x half> poison, half [[DOTSCALAR]], i64 0
// CHECK-NEXT: [[TMP1:%.*]] = shufflevector <2 x half> [[TMP0]], <2 x half> poison, <2 x i32> zeroinitializer
@@ -42,7 +42,7 @@ half test_reflect_half(half I, half N) {
// SPVCHECK-LABEL: define 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 <2 x half> @llvm.spv.reflect.v2f16(<2 x half> [[I]], <2 x half> [[N]])
// SPVCHECK-NEXT: ret <2 x half> [[SPV_REFLECT_I]]
//
half2 test_reflect_half2(half2 I, half2 N) {
@@ -52,7 +52,7 @@ half2 test_reflect_half2(half2 I, half2 N) {
// CHECK-LABEL: define noundef nofpclass(nan inf) <3 x half> @_Z18test_reflect_half3Dv3_DhS_(
// CHECK-SAME: <3 x half> noundef nofpclass(nan inf) [[I:%.*]], <3 x half> noundef nofpclass(nan inf) [[N:%.*]]) 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) [[I]], <3 x half> nofpclass(nan inf) [[N]])
+// CHECK-NEXT: [[HLSL_DOT_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn half @llvm.dx.fdot.v3f16(<3 x half> [[I]], <3 x half> [[N]])
// CHECK-NEXT: [[DOTSCALAR:%.*]] = fmul reassoc nnan ninf nsz arcp afn half [[HLSL_DOT_I]], 0xH4000
// CHECK-NEXT: [[TMP0:%.*]] = insertelement <3 x half> poison, half [[DOTSCALAR]], i64 0
// CHECK-NEXT: [[TMP1:%.*]] = shufflevector <3 x half> [[TMP0]], <3 x half> poison, <3 x i32> zeroinitializer
@@ -63,7 +63,7 @@ half2 test_reflect_half2(half2 I, half2 N) {
// SPVCHECK-LABEL: define 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 <3 x half> @llvm.spv.reflect.v3f16(<3 x half> [[I]], <3 x half> [[N]])
// SPVCHECK-NEXT: ret <3 x half> [[SPV_REFLECT_I]]
//
half3 test_reflect_half3(half3 I, half3 N) {
@@ -73,7 +73,7 @@ half3 test_reflect_half3(half3 I, half3 N) {
// CHECK-LABEL: define noundef nofpclass(nan inf) <4 x half> @_Z18test_reflect_half4Dv4_DhS_(
// CHECK-SAME: <4 x half> noundef nofpclass(nan inf) [[I:%.*]], <4 x half> noundef nofpclass(nan inf) [[N:%.*]]) 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) [[I]], <4 x half> nofpclass(nan inf) [[N]])
+// CHECK-NEXT: [[HLSL_DOT_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn half @llvm.dx.fdot.v4f16(<4 x half> [[I]], <4 x half> [[N]])
// CHECK-NEXT: [[DOTSCALAR:%.*]] = fmul reassoc nnan ninf nsz arcp afn half [[HLSL_DOT_I]], 0xH4000
// CHECK-NEXT: [[TMP0:%.*]] = insertelement <4 x half> poison, half [[DOTSCALAR]], i64 0
// CHECK-NEXT: [[TMP1:%.*]] = shufflevector <4 x half> [[TMP0]], <4 x half> poison, <4 x i32> zeroinitializer
@@ -84,7 +84,7 @@ half3 test_reflect_half3(half3 I, half3 N) {
// SPVCHECK-LABEL: define 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 <4 x half> @llvm.spv.reflect.v4f16(<4 x half> [[I]], <4 x half> [[N]])
// SPVCHECK-NEXT: ret <4 x half> [[SPV_REFLECT_I]]
//
half4 test_reflect_half4(half4 I, half4 N) {
@@ -116,7 +116,7 @@ float test_reflect_float(float I, float N) {
// CHECK-LABEL: define noundef nofpclass(nan inf) <2 x float> @_Z19test_reflect_float2Dv2_fS_(
// CHECK-SAME: <2 x float> noundef nofpclass(nan inf) [[I:%.*]], <2 x float> noundef nofpclass(nan inf) [[N:%.*]]) 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) [[I]], <2 x float> nofpclass(nan inf) [[N]])
+// CHECK-NEXT: [[HLSL_DOT_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn float @llvm.dx.fdot.v2f32(<2 x float> [[I]], <2 x float> [[N]])
// CHECK-NEXT: [[DOTSCALAR:%.*]] = fmul reassoc nnan ninf nsz arcp afn float [[HLSL_DOT_I]], 2.000000e+00
// CHECK-NEXT: [[TMP0:%.*]] = insertelement <2 x float> poison, float [[DOTSCALAR]], i64 0
// CHECK-NEXT: [[TMP1:%.*]] = shufflevector <2 x float> [[TMP0]], <2 x float> poison, <2 x i32> zeroinitializer
@@ -127,7 +127,7 @@ float test_reflect_float(float I, float N) {
// SPVCHECK-LABEL: define 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 <2 x float> @llvm.spv.reflect.v2f32(<2 x float> [[I]], <2 x float> [[N]])
// SPVCHECK-NEXT: ret <2 x float> [[SPV_REFLECT_I]]
//
float2 test_reflect_float2(float2 I, float2 N) {
@@ -137,7 +137,7 @@ float2 test_reflect_float2(float2 I, float2 N) {
// CHECK-LABEL: define noundef nofpclass(nan inf) <3 x float> @_Z19test_reflect_float3Dv3_fS_(
// CHECK-SAME: <3 x float> noundef nofpclass(nan inf) [[I:%.*]], <3 x float> noundef nofpclass(nan inf) [[N:%.*]]) 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) [[I]], <3 x float> nofpclass(nan inf) [[N]])
+// CHECK-NEXT: [[HLSL_DOT_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn float @llvm.dx.fdot.v3f32(<3 x float> [[I]], <3 x float> [[N]])
// CHECK-NEXT: [[DOTSCALAR:%.*]] = fmul reassoc nnan ninf nsz arcp afn float [[HLSL_DOT_I]], 2.000000e+00
// CHECK-NEXT: [[TMP0:%.*]] = insertelement <3 x float> poison, float [[DOTSCALAR]], i64 0
// CHECK-NEXT: [[TMP1:%.*]] = shufflevector <3 x float> [[TMP0]], <3 x float> poison, <3 x i32> zeroinitializer
@@ -148,7 +148,7 @@ float2 test_reflect_float2(float2 I, float2 N) {
// SPVCHECK-LABEL: define 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 <3 x float> @llvm.spv.reflect.v3f32(<3 x float> [[I]], <3 x float> [[N]])
// SPVCHECK-NEXT: ret <3 x float> [[SPV_REFLECT_I]]
//
float3 test_reflect_float3(float3 I, float3 N) {
@@ -158,7 +158,7 @@ float3 test_reflect_float3(float3 I, float3 N) {
// CHECK-LABEL: define noundef nofpclass(nan inf) <4 x float> @_Z19test_reflect_float4Dv4_fS_(
// CHECK-SAME: <4 x float> noundef nofpclass(nan inf) [[I:%.*]], <4 x float> noundef nofpclass(nan inf) [[N:%.*]]) 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) [[I]], <4 x float> nofpclass(nan inf) [[N]])
+// CHECK-NEXT: [[HLSL_DOT_I:%.*]] = tail call reassoc nnan ninf nsz arcp afn float @llvm.dx.fdot.v4f32(<4 x float> [[I]], <4 x float> [[N]])
// CHECK-NEXT: [[DOTSCALAR:%.*]] = fmul reassoc nnan ninf nsz arcp afn float [[HLSL_DOT_I]], 2.000000e+00
// CHECK-NEXT: [[TMP0:%.*]] = insertelement <4 x float> poison, float [[DOTSCALAR]], i64 0
// CHECK-NEXT: [[TMP1:%.*]] = shufflevector <4 x float> [[TMP0]], <4 x float> poison, <4 x i32> zeroinitializer
@@ -169,7 +169,7 @@ float3 test_reflect_float3(float3 I, float3 N) {
// SPVCHECK-LABEL: define 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 <4 x float> @llvm.spv.reflect.v4f32(<4 x float> [[I]], <4 x float> [[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/smoothstep.hlsl b/clang/test/CodeGenHLSL/builtins/smoothstep.hlsl
index d3e5c1059029c..f2328c7330e6c 100644
--- a/clang/test/CodeGenHLSL/builtins/smoothstep.hlsl
+++ b/clang/test/CodeGenHLSL/builtins/smoothstep.hlsl
@@ -22,7 +22,7 @@
// SPVCHECK-LABEL: define 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 half @llvm.spv.smoothstep.f16(half [[MIN]], half [[MAX]], half [[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 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 <2 x half> @llvm.spv.smoothstep.v2f16(<2 x half> [[MIN]], <2 x half> [[MAX]], <2 x half> [[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 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 <3 x half> @llvm.spv.smoothstep.v3f16(<3 x half> [[MIN]], <3 x half> [[MAX]], <3 x half> [[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 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 <4 x half> @llvm.spv.smoothstep.v4f16(<4 x half> [[MIN]], <4 x half> [[MAX]], <4 x half> [[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 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 float @llvm.spv.smoothstep.f32(float [[MIN]], float [[MAX]], float [[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 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 <2 x float> @llvm.spv.smoothstep.v2f32(<2 x float> [[MIN]], <2 x float> [[MAX]], <2 x float> [[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 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 <3 x float> @llvm.spv.smoothstep.v3f32(<3 x float> [[MIN]], <3 x float> [[MAX]], <3 x float> [[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 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 <4 x float> @llvm.spv.smoothstep.v4f32(<4 x float> [[MIN]], <4 x float> [[MAX]], <4 x float> [[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 7d812fd0265a6..0c9ff4cdd7808 100644
--- a/clang/test/Headers/__clang_hip_cmath.hip
+++ b/clang/test/Headers/__clang_hip_cmath.hip
@@ -24,7 +24,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 half @llvm.fma.f16(half [[X:%.*]], half [[Y:%.*]], half [[Z:%.*]])
// FINITEONLY-NEXT: ret half [[TMP0]]
//
extern "C" __device__ _Float16 test_fma_f16(_Float16 x, _Float16 y,
@@ -34,12 +34,12 @@ extern "C" __device__ _Float16 test_fma_f16(_Float16 x, _Float16 y,
// DEFAULT-LABEL: @test_pow_f16(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract noundef half @__ocml_pown_f16(half noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR9:[0-9]+]]
+// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract noundef half @__ocml_pown_f16(half noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR7:[0-9]+]]
// DEFAULT-NEXT: ret half [[CALL_I]]
//
// FINITEONLY-LABEL: @test_pow_f16(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) half @__ocml_pown_f16(half noundef nofpclass(nan inf) [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR9:[0-9]+]]
+// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) half @__ocml_pown_f16(half noundef nofpclass(nan inf) [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR7:[0-9]+]]
// FINITEONLY-NEXT: ret half [[CALL_I]]
//
extern "C" __device__ _Float16 test_pow_f16(_Float16 x, int y) {
@@ -53,7 +53,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 float @llvm.fabs.f32(float [[X:%.*]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
extern "C" __device__ float test_fabs_f32(float x) {
@@ -62,12 +62,12 @@ extern "C" __device__ float test_fabs_f32(float x) {
// DEFAULT-LABEL: @test_sin_f32(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I1:%.*]] = tail call contract noundef float @__ocml_sin_f32(float noundef [[X:%.*]]) #[[ATTR10:[0-9]+]]
+// DEFAULT-NEXT: [[CALL_I1:%.*]] = tail call contract noundef float @__ocml_sin_f32(float noundef [[X:%.*]]) #[[ATTR8:[0-9]+]]
// DEFAULT-NEXT: ret float [[CALL_I1]]
//
// FINITEONLY-LABEL: @test_sin_f32(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I1:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_sin_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR10:[0-9]+]]
+// FINITEONLY-NEXT: [[CALL_I1:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_sin_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR8:[0-9]+]]
// FINITEONLY-NEXT: ret float [[CALL_I1]]
//
extern "C" __device__ float test_sin_f32(float x) {
@@ -76,12 +76,12 @@ extern "C" __device__ float test_sin_f32(float x) {
// DEFAULT-LABEL: @test_cos_f32(
// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[CALL_I1:%.*]] = tail call contract noundef float @__ocml_cos_f32(float noundef [[X:%.*]]) #[[ATTR10]]
+// DEFAULT-NEXT: [[CALL_I1:%.*]] = tail call contract noundef float @__ocml_cos_f32(float noundef [[X:%.*]]) #[[ATTR8]]
// DEFAULT-NEXT: ret float [[CALL_I1]]
//
// FINITEONLY-LABEL: @test_cos_f32(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[CALL_I1:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_cos_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR10]]
+// FINITEONLY-NEXT: [[CALL_I1:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_cos_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR8]]
// FINITEONLY-NEXT: ret float [[CALL_I1]]
//
extern "C" __device__ float test_cos_f32(float x) {
@@ -97,46 +97,10 @@ struct user_bfloat16 {
};
namespace user_namespace {
-// DEFAULT-LABEL: @_ZN14user_namespace3fmaE13user_bfloat16S0_S0_(
-// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: ret void
-//
-// FINITEONLY-LABEL: @_ZN14user_namespace3fmaE13user_bfloat16S0_S0_(
-// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: ret void
-//
__device__ user_bfloat16 fma(const user_bfloat16 a, const user_bfloat16 b, const user_bfloat16 c) {
return a;
}
-// DEFAULT-LABEL: @_ZN14user_namespace8test_fmaEv(
-// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[A:%.*]] = alloca [[STRUCT_USER_BFLOAT16:%.*]], align 1, addrspace(5)
-// DEFAULT-NEXT: [[B:%.*]] = alloca [[STRUCT_USER_BFLOAT16]], align 1, addrspace(5)
-// DEFAULT-NEXT: [[A_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A]] to ptr
-// DEFAULT-NEXT: [[B_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[B]] to ptr
-// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 1, ptr addrspace(5) [[A]]) #[[ATTR11:[0-9]+]]
-// DEFAULT-NEXT: call void @_ZN13user_bfloat16C1Ef(ptr noundef nonnull align 1 dereferenceable(1) [[A_ASCAST]], float noundef 1.000000e+00) #[[ATTR10]]
-// DEFAULT-NEXT: call void @llvm.lifetime.start.p5(i64 1, ptr addrspace(5) [[B]]) #[[ATTR11]]
-// DEFAULT-NEXT: call void @_ZN13user_bfloat16C1Ef(ptr noundef nonnull align 1 dereferenceable(1) [[B_ASCAST]], float noundef 2.000000e+00) #[[ATTR10]]
-// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 1, ptr addrspace(5) [[B]]) #[[ATTR11]]
-// DEFAULT-NEXT: call void @llvm.lifetime.end.p5(i64 1, ptr addrspace(5) [[A]]) #[[ATTR11]]
-// DEFAULT-NEXT: ret void
-//
-// FINITEONLY-LABEL: @_ZN14user_namespace8test_fmaEv(
-// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[A:%.*]] = alloca [[STRUCT_USER_BFLOAT16:%.*]], align 1, addrspace(5)
-// FINITEONLY-NEXT: [[B:%.*]] = alloca [[STRUCT_USER_BFLOAT16]], align 1, addrspace(5)
-// FINITEONLY-NEXT: [[A_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A]] to ptr
-// FINITEONLY-NEXT: [[B_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[B]] to ptr
-// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 1, ptr addrspace(5) [[A]]) #[[ATTR11:[0-9]+]]
-// FINITEONLY-NEXT: call void @_ZN13user_bfloat16C1Ef(ptr noundef nonnull align 1 dereferenceable(1) [[A_ASCAST]], float noundef nofpclass(nan inf) 1.000000e+00) #[[ATTR10]]
-// FINITEONLY-NEXT: call void @llvm.lifetime.start.p5(i64 1, ptr addrspace(5) [[B]]) #[[ATTR11]]
-// FINITEONLY-NEXT: call void @_ZN13user_bfloat16C1Ef(ptr noundef nonnull align 1 dereferenceable(1) [[B_ASCAST]], float noundef nofpclass(nan inf) 2.000000e+00) #[[ATTR10]]
-// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 1, ptr addrspace(5) [[B]]) #[[ATTR11]]
-// FINITEONLY-NEXT: call void @llvm.lifetime.end.p5(i64 1, ptr addrspace(5) [[A]]) #[[ATTR11]]
-// FINITEONLY-NEXT: ret void
-//
__global__ void test_fma() {
user_bfloat16 a = 1.0f, b = 2.0f;
fma(a, b, b);
diff --git a/clang/test/Headers/__clang_hip_math.hip b/clang/test/Headers/__clang_hip_math.hip
index df1cd716342a5..e879fec0ebe5a 100644
--- a/clang/test/Headers/__clang_hip_math.hip
+++ b/clang/test/Headers/__clang_hip_math.hip
@@ -842,7 +842,7 @@ extern "C" __device__ double test_cbrt(double x) {
//
// FINITEONLY-LABEL: @test_ceilf(
// 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 float @llvm.ceil.f32(float [[X:%.*]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: @test_ceilf(
@@ -866,7 +866,7 @@ extern "C" __device__ float test_ceilf(float x) {
//
// FINITEONLY-LABEL: @test_ceil(
// 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 double @llvm.ceil.f64(double [[X:%.*]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: @test_ceil(
@@ -890,7 +890,7 @@ extern "C" __device__ double test_ceil(double x) {
//
// FINITEONLY-LABEL: @test_copysignf(
// 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 float @llvm.copysign.f32(float [[X:%.*]], float [[Y:%.*]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: @test_copysignf(
@@ -914,7 +914,7 @@ extern "C" __device__ float test_copysignf(float x, float y) {
//
// FINITEONLY-LABEL: @test_copysign(
// 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 double @llvm.copysign.f64(double [[X:%.*]], double [[Y:%.*]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: @test_copysign(
@@ -1274,7 +1274,7 @@ extern "C" __device__ double test_erfinv(double x) {
//
// FINITEONLY-LABEL: @test_exp10f(
// 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 float @llvm.exp10.f32(float [[X:%.*]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: @test_exp10f(
@@ -1322,7 +1322,7 @@ extern "C" __device__ double test_exp10(double x) {
//
// FINITEONLY-LABEL: @test_exp2f(
// 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 float @llvm.exp2.f32(float [[X:%.*]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: @test_exp2f(
@@ -1370,7 +1370,7 @@ extern "C" __device__ double test_exp2(double x) {
//
// FINITEONLY-LABEL: @test_expf(
// 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 float @llvm.exp.f32(float [[X:%.*]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: @test_expf(
@@ -1466,7 +1466,7 @@ extern "C" __device__ double test_expm1(double x) {
//
// FINITEONLY-LABEL: @test_fabsf(
// 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 float @llvm.fabs.f32(float [[X:%.*]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: @test_fabsf(
@@ -1490,7 +1490,7 @@ extern "C" __device__ float test_fabsf(float x) {
//
// FINITEONLY-LABEL: @test_fabs(
// 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 double @llvm.fabs.f64(double [[X:%.*]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: @test_fabs(
@@ -1586,7 +1586,7 @@ extern "C" __device__ float test_fdividef(float x, float y) {
//
// FINITEONLY-LABEL: @test_floorf(
// 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 float @llvm.floor.f32(float [[X:%.*]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: @test_floorf(
@@ -1610,7 +1610,7 @@ extern "C" __device__ float test_floorf(float x) {
//
// FINITEONLY-LABEL: @test_floor(
// 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 double @llvm.floor.f64(double [[X:%.*]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: @test_floor(
@@ -1634,7 +1634,7 @@ extern "C" __device__ double test_floor(double x) {
//
// FINITEONLY-LABEL: @test_fmaf(
// 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 float @llvm.fma.f32(float [[X:%.*]], float [[Y:%.*]], float [[Z:%.*]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: @test_fmaf(
@@ -1658,7 +1658,7 @@ extern "C" __device__ float test_fmaf(float x, float y, float z) {
//
// FINITEONLY-LABEL: @test_fma(
// 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 double @llvm.fma.f64(double [[X:%.*]], double [[Y:%.*]], double [[Z:%.*]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: @test_fma(
@@ -1682,7 +1682,7 @@ extern "C" __device__ double test_fma(double x, double y, double z) {
//
// FINITEONLY-LABEL: @test_fma_rn(
// 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 double @llvm.fma.f64(double [[X:%.*]], double [[Y:%.*]], double [[Z:%.*]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: @test_fma_rn(
@@ -1706,7 +1706,7 @@ extern "C" __device__ double test_fma_rn(double x, double y, double z) {
//
// FINITEONLY-LABEL: @test_fmaxf(
// 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 float @llvm.maxnum.f32(float [[X:%.*]], float [[Y:%.*]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: @test_fmaxf(
@@ -1730,7 +1730,7 @@ extern "C" __device__ float test_fmaxf(float x, float y) {
//
// FINITEONLY-LABEL: @test_fmax(
// 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 double @llvm.maxnum.f64(double [[X:%.*]], double [[Y:%.*]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: @test_fmax(
@@ -1754,7 +1754,7 @@ extern "C" __device__ double test_fmax(double x, double y) {
//
// FINITEONLY-LABEL: @test_fminf(
// 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 float @llvm.minnum.f32(float [[X:%.*]], float [[Y:%.*]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: @test_fminf(
@@ -1778,7 +1778,7 @@ extern "C" __device__ float test_fminf(float x, float y) {
//
// FINITEONLY-LABEL: @test_fmin(
// 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 double @llvm.minnum.f64(double [[X:%.*]], double [[Y:%.*]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: @test_fmin(
@@ -1843,29 +1843,13 @@ extern "C" __device__ double test_fmod(double x, double y) {
return fmod(x, y);
}
-// DEFAULT-LABEL: @test_frexpf(
-// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[TMP0:%.*]] = tail call { float, i32 } @llvm.frexp.f32.i32(float [[X:%.*]])
-// DEFAULT-NEXT: [[TMP1:%.*]] = extractvalue { float, i32 } [[TMP0]], 1
-// DEFAULT-NEXT: store i32 [[TMP1]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA12:![0-9]+]]
-// DEFAULT-NEXT: [[TMP2:%.*]] = extractvalue { float, i32 } [[TMP0]], 0
-// DEFAULT-NEXT: ret float [[TMP2]]
-//
-// FINITEONLY-LABEL: @test_frexpf(
-// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call { float, i32 } @llvm.frexp.f32.i32(float nofpclass(nan inf) [[X:%.*]])
-// FINITEONLY-NEXT: [[TMP1:%.*]] = extractvalue { float, i32 } [[TMP0]], 1
-// FINITEONLY-NEXT: store i32 [[TMP1]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA12:![0-9]+]]
-// FINITEONLY-NEXT: [[TMP2:%.*]] = extractvalue { float, i32 } [[TMP0]], 0
-// FINITEONLY-NEXT: ret float [[TMP2]]
-//
-// APPROX-LABEL: @test_frexpf(
-// APPROX-NEXT: entry:
-// APPROX-NEXT: [[TMP0:%.*]] = tail call { float, i32 } @llvm.frexp.f32.i32(float [[X:%.*]])
-// APPROX-NEXT: [[TMP1:%.*]] = extractvalue { float, i32 } [[TMP0]], 1
-// APPROX-NEXT: store i32 [[TMP1]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA12:![0-9]+]]
-// APPROX-NEXT: [[TMP2:%.*]] = extractvalue { float, i32 } [[TMP0]], 0
-// APPROX-NEXT: ret float [[TMP2]]
+// CHECK-LABEL: @test_frexpf(
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[TMP0:%.*]] = tail call { float, i32 } @llvm.frexp.f32.i32(float [[X:%.*]])
+// CHECK-NEXT: [[TMP1:%.*]] = extractvalue { float, i32 } [[TMP0]], 1
+// CHECK-NEXT: store i32 [[TMP1]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA12:![0-9]+]]
+// CHECK-NEXT: [[TMP2:%.*]] = extractvalue { float, i32 } [[TMP0]], 0
+// CHECK-NEXT: ret float [[TMP2]]
//
// AMDGCNSPIRV-LABEL: @test_frexpf(
// AMDGCNSPIRV-NEXT: entry:
@@ -1879,29 +1863,13 @@ extern "C" __device__ float test_frexpf(float x, int* y) {
return frexpf(x, y);
}
-// DEFAULT-LABEL: @test_frexp(
-// DEFAULT-NEXT: entry:
-// DEFAULT-NEXT: [[TMP0:%.*]] = tail call { double, i32 } @llvm.frexp.f64.i32(double [[X:%.*]])
-// DEFAULT-NEXT: [[TMP1:%.*]] = extractvalue { double, i32 } [[TMP0]], 1
-// DEFAULT-NEXT: store i32 [[TMP1]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA12]]
-// DEFAULT-NEXT: [[TMP2:%.*]] = extractvalue { double, i32 } [[TMP0]], 0
-// DEFAULT-NEXT: ret double [[TMP2]]
-//
-// FINITEONLY-LABEL: @test_frexp(
-// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call { double, i32 } @llvm.frexp.f64.i32(double nofpclass(nan inf) [[X:%.*]])
-// FINITEONLY-NEXT: [[TMP1:%.*]] = extractvalue { double, i32 } [[TMP0]], 1
-// FINITEONLY-NEXT: store i32 [[TMP1]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA12]]
-// FINITEONLY-NEXT: [[TMP2:%.*]] = extractvalue { double, i32 } [[TMP0]], 0
-// FINITEONLY-NEXT: ret double [[TMP2]]
-//
-// APPROX-LABEL: @test_frexp(
-// APPROX-NEXT: entry:
-// APPROX-NEXT: [[TMP0:%.*]] = tail call { double, i32 } @llvm.frexp.f64.i32(double [[X:%.*]])
-// APPROX-NEXT: [[TMP1:%.*]] = extractvalue { double, i32 } [[TMP0]], 1
-// APPROX-NEXT: store i32 [[TMP1]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA12]]
-// APPROX-NEXT: [[TMP2:%.*]] = extractvalue { double, i32 } [[TMP0]], 0
-// APPROX-NEXT: ret double [[TMP2]]
+// CHECK-LABEL: @test_frexp(
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[TMP0:%.*]] = tail call { double, i32 } @llvm.frexp.f64.i32(double [[X:%.*]])
+// CHECK-NEXT: [[TMP1:%.*]] = extractvalue { double, i32 } [[TMP0]], 1
+// CHECK-NEXT: store i32 [[TMP1]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA12]]
+// CHECK-NEXT: [[TMP2:%.*]] = extractvalue { double, i32 } [[TMP0]], 0
+// CHECK-NEXT: ret double [[TMP2]]
//
// AMDGCNSPIRV-LABEL: @test_frexp(
// AMDGCNSPIRV-NEXT: entry:
@@ -2554,7 +2522,7 @@ extern "C" __device__ double test_jn(int x, double y) {
//
// FINITEONLY-LABEL: @test_ldexpf(
// 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 float @llvm.ldexp.f32.i32(float [[X:%.*]], i32 [[Y:%.*]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: @test_ldexpf(
@@ -2578,7 +2546,7 @@ extern "C" __device__ float test_ldexpf(float x, int y) {
//
// FINITEONLY-LABEL: @test_ldexp(
// 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 double @llvm.ldexp.f64.i32(double [[X:%.*]], i32 [[Y:%.*]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: @test_ldexp(
@@ -2651,7 +2619,7 @@ extern "C" __device__ double test_lgamma(double x) {
//
// FINITEONLY-LABEL: @test_llrintf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.rint.f32(float nofpclass(nan inf) [[X:%.*]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.rint.f32(float [[X:%.*]])
// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64
// FINITEONLY-NEXT: ret i64 [[CONV_I]]
//
@@ -2679,7 +2647,7 @@ extern "C" __device__ long long int test_llrintf(float x) {
//
// FINITEONLY-LABEL: @test_llrint(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.rint.f64(double nofpclass(nan inf) [[X:%.*]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.rint.f64(double [[X:%.*]])
// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64
// FINITEONLY-NEXT: ret i64 [[CONV_I]]
//
@@ -2707,7 +2675,7 @@ extern "C" __device__ long long int test_llrint(double x) {
//
// FINITEONLY-LABEL: @test_llroundf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.round.f32(float nofpclass(nan inf) [[X:%.*]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.round.f32(float [[X:%.*]])
// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64
// FINITEONLY-NEXT: ret i64 [[CONV_I]]
//
@@ -2735,7 +2703,7 @@ extern "C" __device__ long long int test_llroundf(float x) {
//
// FINITEONLY-LABEL: @test_llround(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.round.f64(double nofpclass(nan inf) [[X:%.*]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.round.f64(double [[X:%.*]])
// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64
// FINITEONLY-NEXT: ret i64 [[CONV_I]]
//
@@ -2762,7 +2730,7 @@ extern "C" __device__ long long int test_llround(double x) {
//
// FINITEONLY-LABEL: @test_log10f(
// 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 float @llvm.log10.f32(float [[X:%.*]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: @test_log10f(
@@ -2858,7 +2826,7 @@ extern "C" __device__ double test_log1p(double x) {
//
// FINITEONLY-LABEL: @test_log2f(
// 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 float @llvm.log2.f32(float [[X:%.*]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: @test_log2f(
@@ -2954,7 +2922,7 @@ extern "C" __device__ double test_logb(double x) {
//
// FINITEONLY-LABEL: @test_logf(
// 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 float @llvm.log.f32(float [[X:%.*]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: @test_logf(
@@ -2979,7 +2947,7 @@ extern "C" __device__ float test_logf(float x) {
//
// FINITEONLY-LABEL: @test_lrintf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.rint.f32(float nofpclass(nan inf) [[X:%.*]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.rint.f32(float [[X:%.*]])
// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64
// FINITEONLY-NEXT: ret i64 [[CONV_I]]
//
@@ -3007,7 +2975,7 @@ extern "C" __device__ long int test_lrintf(float x) {
//
// FINITEONLY-LABEL: @test_lrint(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.rint.f64(double nofpclass(nan inf) [[X:%.*]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.rint.f64(double [[X:%.*]])
// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64
// FINITEONLY-NEXT: ret i64 [[CONV_I]]
//
@@ -3035,7 +3003,7 @@ extern "C" __device__ long int test_lrint(double x) {
//
// FINITEONLY-LABEL: @test_lroundf(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.round.f32(float nofpclass(nan inf) [[X:%.*]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.round.f32(float [[X:%.*]])
// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64
// FINITEONLY-NEXT: ret i64 [[CONV_I]]
//
@@ -3063,7 +3031,7 @@ extern "C" __device__ long int test_lroundf(float x) {
//
// FINITEONLY-LABEL: @test_lround(
// FINITEONLY-NEXT: entry:
-// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.round.f64(double nofpclass(nan inf) [[X:%.*]])
+// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.round.f64(double [[X:%.*]])
// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64
// FINITEONLY-NEXT: ret i64 [[CONV_I]]
//
@@ -3827,7 +3795,7 @@ extern "C" __device__ double test_nan_fill() {
//
// FINITEONLY-LABEL: @test_nearbyintf(
// 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 float @llvm.nearbyint.f32(float [[X:%.*]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: @test_nearbyintf(
@@ -3851,7 +3819,7 @@ extern "C" __device__ float test_nearbyintf(float x) {
//
// FINITEONLY-LABEL: @test_nearbyint(
// 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 double @llvm.nearbyint.f64(double [[X:%.*]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: @test_nearbyint(
@@ -4613,7 +4581,7 @@ extern "C" __device__ double test_rhypot(double x, double y) {
//
// FINITEONLY-LABEL: @test_rintf(
// 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 float @llvm.rint.f32(float [[X:%.*]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: @test_rintf(
@@ -4637,7 +4605,7 @@ extern "C" __device__ float test_rintf(float x) {
//
// FINITEONLY-LABEL: @test_rint(
// 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 double @llvm.rint.f64(double [[X:%.*]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: @test_rint(
@@ -4925,7 +4893,7 @@ extern "C" __device__ double test_rnorm4d(double x, double y, double z, double w
//
// FINITEONLY-LABEL: @test_roundf(
// 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 float @llvm.round.f32(float [[X:%.*]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: @test_roundf(
@@ -4949,7 +4917,7 @@ extern "C" __device__ float test_roundf(float x) {
//
// FINITEONLY-LABEL: @test_round(
// 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 double @llvm.round.f64(double [[X:%.*]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: @test_round(
@@ -5025,7 +4993,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 float @llvm.ldexp.f32.i32(float [[X:%.*]], i32 [[CONV_I]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: @test_scalblnf(
@@ -5057,7 +5025,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 double @llvm.ldexp.f64.i32(double [[X:%.*]], i32 [[CONV_I]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: @test_scalbln(
@@ -5085,7 +5053,7 @@ extern "C" __device__ double test_scalbln(double x, long int y) {
//
// FINITEONLY-LABEL: @test_scalbnf(
// 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 float @llvm.ldexp.f32.i32(float [[X:%.*]], i32 [[Y:%.*]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: @test_scalbnf(
@@ -5109,7 +5077,7 @@ extern "C" __device__ float test_scalbnf(float x, int y) {
//
// FINITEONLY-LABEL: @test_scalbn(
// 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 double @llvm.ldexp.f64.i32(double [[X:%.*]], i32 [[Y:%.*]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: @test_scalbn(
@@ -5459,7 +5427,7 @@ extern "C" __device__ double test_sinpi(double x) {
//
// FINITEONLY-LABEL: @test_sqrtf(
// 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 float @llvm.sqrt.f32(float [[X:%.*]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: @test_sqrtf(
@@ -5483,7 +5451,7 @@ extern "C" __device__ float test_sqrtf(float x) {
//
// FINITEONLY-LABEL: @test_sqrt(
// 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 double @llvm.sqrt.f64(double [[X:%.*]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: @test_sqrt(
@@ -5651,7 +5619,7 @@ extern "C" __device__ double test_tgamma(double x) {
//
// FINITEONLY-LABEL: @test_truncf(
// 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 float @llvm.trunc.f32(float [[X:%.*]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: @test_truncf(
@@ -5675,7 +5643,7 @@ extern "C" __device__ float test_truncf(float x) {
//
// FINITEONLY-LABEL: @test_trunc(
// 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 double @llvm.trunc.f64(double [[X:%.*]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: @test_trunc(
@@ -6195,7 +6163,7 @@ extern "C" __device__ float test___fdividef(float x, float y) {
//
// FINITEONLY-LABEL: @test__fmaf_rn(
// 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 float @llvm.fma.f32(float [[X:%.*]], float [[Y:%.*]], float [[Z:%.*]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: @test__fmaf_rn(
@@ -6267,7 +6235,7 @@ extern "C" __device__ float test___frcp_rn(float x) {
//
// FINITEONLY-LABEL: @test___frsqrt_rn(
// 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 float @llvm.amdgcn.rsq.f32(float [[X:%.*]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: @test___frsqrt_rn(
@@ -6339,7 +6307,7 @@ extern "C" __device__ float test___fsub_rn(float x, float y) {
//
// FINITEONLY-LABEL: @test___log10f(
// 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 float @llvm.log10.f32(float [[X:%.*]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: @test___log10f(
@@ -6363,7 +6331,7 @@ extern "C" __device__ float test___log10f(float x) {
//
// FINITEONLY-LABEL: @test___log2f(
// 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 float @llvm.amdgcn.log.f32(float [[X:%.*]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: @test___log2f(
@@ -6387,7 +6355,7 @@ extern "C" __device__ float test___log2f(float x) {
//
// FINITEONLY-LABEL: @test___logf(
// 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 float @llvm.log.f32(float [[X:%.*]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: @test___logf(
@@ -6663,7 +6631,7 @@ extern "C" __device__ double test___drcp_rn(double x) {
//
// FINITEONLY-LABEL: @test___dsqrt_rn(
// 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 double @llvm.sqrt.f64(double [[X:%.*]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: @test___dsqrt_rn(
@@ -6687,7 +6655,7 @@ extern "C" __device__ double test___dsqrt_rn(double x) {
//
// FINITEONLY-LABEL: @test__fma_rn(
// 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 double @llvm.fma.f64(double [[X:%.*]], double [[Y:%.*]], double [[Z:%.*]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: @test__fma_rn(
@@ -6711,7 +6679,7 @@ extern "C" __device__ double test__fma_rn(double x, double y, double z) {
//
// FINITEONLY-LABEL: @test_float_min(
// 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 float @llvm.minnum.f32(float [[X:%.*]], float [[Y:%.*]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: @test_float_min(
@@ -6735,7 +6703,7 @@ extern "C" __device__ float test_float_min(float x, float y) {
//
// FINITEONLY-LABEL: @test_float_max(
// 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 float @llvm.maxnum.f32(float [[X:%.*]], float [[Y:%.*]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: @test_float_max(
@@ -6759,7 +6727,7 @@ extern "C" __device__ float test_float_max(float x, float y) {
//
// FINITEONLY-LABEL: @test_double_min(
// 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 double @llvm.minnum.f64(double [[X:%.*]], double [[Y:%.*]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: @test_double_min(
@@ -6783,7 +6751,7 @@ extern "C" __device__ double test_double_min(double x, double y) {
//
// FINITEONLY-LABEL: @test_double_max(
// 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 double @llvm.maxnum.f64(double [[X:%.*]], double [[Y:%.*]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: @test_double_max(
diff --git a/llvm/lib/Transforms/Utils/InlineFunction.cpp b/llvm/lib/Transforms/Utils/InlineFunction.cpp
index c65bf16b6a937..5beee1f681b81 100644
--- a/llvm/lib/Transforms/Utils/InlineFunction.cpp
+++ b/llvm/lib/Transforms/Utils/InlineFunction.cpp
@@ -1381,8 +1381,7 @@ static void AddParamAndFnBasicAttributes(const CallBase &CB,
// behavior was just using a poison value.
static const Attribute::AttrKind ExactAttrsToPropagate[] = {
Attribute::Dereferenceable, Attribute::DereferenceableOrNull,
- Attribute::NonNull, Attribute::NoFPClass,
- Attribute::Alignment, Attribute::Range};
+ Attribute::NonNull, Attribute::Alignment, Attribute::Range};
for (unsigned I = 0, E = CB.arg_size(); I < E; ++I) {
ValidObjParamAttrs.emplace_back(AttrBuilder{CB.getContext()});
@@ -1464,13 +1463,6 @@ static void AddParamAndFnBasicAttributes(const CallBase &CB,
NewAB.addRangeAttr(CombinedRange);
}
}
-
- if (FPClassTest ExistingNoFP = AL.getParamNoFPClass(I)) {
- FPClassTest NewNoFP =
- NewAB.getAttribute(Attribute::NoFPClass).getNoFPClass();
- NewAB.addNoFPClassAttr(ExistingNoFP | NewNoFP);
- }
-
AL = AL.addParamAttributes(Context, I, NewAB);
} else if (NewInnerCB->getArgOperand(I)->getType()->isPointerTy()) {
// Check if the underlying value for the parameter is an argument.
diff --git a/llvm/test/Transforms/Inline/access-attributes-prop.ll b/llvm/test/Transforms/Inline/access-attributes-prop.ll
index 5a102d14b5c90..5bf845d5ba94b 100644
--- a/llvm/test/Transforms/Inline/access-attributes-prop.ll
+++ b/llvm/test/Transforms/Inline/access-attributes-prop.ll
@@ -750,47 +750,3 @@ define void @prop_range_direct(i32 %v) {
call void @foo4(i32 range(i32 1, 11) %v)
ret void
}
-
-declare void @bar_fp(float %x)
-
-define void @foo_fp(float %x) {
-; CHECK-LABEL: define {{[^@]+}}@foo_fp
-; CHECK-SAME: (float [[X:%.*]]) {
-; CHECK-NEXT: call void @bar_fp(float [[X]])
-; CHECK-NEXT: ret void
-;
- call void @bar_fp(float %x)
- ret void
-}
-
-define void @prop_param_nofpclass(float %x) {
-; CHECK-LABEL: define {{[^@]+}}@prop_param_nofpclass
-; CHECK-SAME: (float [[X:%.*]]) {
-; CHECK-NEXT: call void @bar_fp(float nofpclass(nan inf) [[X]])
-; CHECK-NEXT: ret void
-;
- call void @foo_fp(float nofpclass(nan inf) %x)
- ret void
-}
-
-declare void @func_fp(float)
-
-define void @union_nofpclass(float %v) {
-; CHECK-LABEL: define {{[^@]+}}@union_nofpclass
-; CHECK-SAME: (float [[V:%.*]]) {
-; CHECK-NEXT: call void @func_fp(float nofpclass(inf) [[V]])
-; CHECK-NEXT: ret void
-;
- call void @func_fp(float nofpclass(inf) %v)
- ret void
-}
-
-define void @prop_nofpclass_union(float %v) {
-; CHECK-LABEL: define {{[^@]+}}@prop_nofpclass_union
-; CHECK-SAME: (float [[V:%.*]]) {
-; CHECK-NEXT: call void @func_fp(float nofpclass(nan inf) [[V]])
-; CHECK-NEXT: ret void
-;
- call void @union_nofpclass(float nofpclass(nan) %v)
- ret void
-}
More information about the cfe-commits
mailing list