Author: Matt Arsenault Date: 2025-04-08T22:53:17+07:00 New Revision: b0cb672b9968eeee6eb022e98476957dbdf8e6e2
URL: https://github.com/llvm/llvm-project/commit/b0cb672b9968eeee6eb022e98476957dbdf8e6e2 DIFF: https://github.com/llvm/llvm-project/commit/b0cb672b9968eeee6eb022e98476957dbdf8e6e2.diff LOG: Inline: Propagate callsite nofpclass attribute (#134800) Fixes #134070 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 e830903261c8c..ac38cf1853799 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 [[SUB_I]]) +// 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: 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 [[SUB_I]]) +// 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: 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> [[SUB_I]], <2 x half> [[SUB_I]]) +// 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: 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> [[SUB_I]]) +// 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: 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> [[SUB_I]], <3 x half> [[SUB_I]]) +// 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: 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> [[SUB_I]]) +// 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: 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> [[SUB_I]], <4 x half> [[SUB_I]]) +// 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: 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> [[SUB_I]]) +// 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: 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 [[SUB_I]]) +// 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: 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 [[SUB_I]]) +// 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: 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> [[SUB_I]], <2 x float> [[SUB_I]]) +// 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: 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> [[SUB_I]]) +// 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: 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> [[SUB_I]], <3 x float> [[SUB_I]]) +// 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: 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> [[SUB_I]]) +// 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: 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> [[SUB_I]], <4 x float> [[SUB_I]]) +// 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: 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> [[SUB_I]]) +// 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: 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 2d4bbd995298f..0b17d03d7097d 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 [[P0]]) +// 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: 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 [[P0]]) +// 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: 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> [[P0]], <2 x half> [[P0]]) +// 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: 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> [[P0]]) +// 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: 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> [[P0]], <3 x half> [[P0]]) +// 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: 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> [[P0]]) +// 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: 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> [[P0]], <4 x half> [[P0]]) +// 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: 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> [[P0]]) +// 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: 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 [[P0]]) +// 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: 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 [[P0]]) +// 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: 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> [[P0]], <2 x float> [[P0]]) +// 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: 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> [[P0]]) +// 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: 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> [[P0]], <3 x float> [[P0]]) +// 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: 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> [[P0]]) +// 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: 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> [[P0]], <4 x float> [[P0]]) +// 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: 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> [[P0]]) +// 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: 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 35ee059697c4b..c082e63ac1da6 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> [[I]], <2 x half> [[N]]) +// 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: [[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> [[I]], <2 x half> [[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> 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) { @@ -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> [[I]], <3 x half> [[N]]) +// 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: [[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> [[I]], <3 x half> [[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> 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) { @@ -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> [[I]], <4 x half> [[N]]) +// 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: [[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> [[I]], <4 x half> [[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> 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) { @@ -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> [[I]], <2 x float> [[N]]) +// 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: [[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> [[I]], <2 x float> [[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> 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) { @@ -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> [[I]], <3 x float> [[N]]) +// 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: [[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> [[I]], <3 x float> [[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> 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) { @@ -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> [[I]], <4 x float> [[N]]) +// 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: [[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> [[I]], <4 x float> [[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> 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/smoothstep.hlsl b/clang/test/CodeGenHLSL/builtins/smoothstep.hlsl index f2328c7330e6c..d3e5c1059029c 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 [[MIN]], half [[MAX]], half [[X]]) +// 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: 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> [[MIN]], <2 x half> [[MAX]], <2 x half> [[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> 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 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> [[MIN]], <3 x half> [[MAX]], <3 x half> [[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> 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 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> [[MIN]], <4 x half> [[MAX]], <4 x half> [[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> 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 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 [[MIN]], float [[MAX]], float [[X]]) +// 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: 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> [[MIN]], <2 x float> [[MAX]], <2 x float> [[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> 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 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> [[MIN]], <3 x float> [[MAX]], <3 x float> [[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> 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 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> [[MIN]], <4 x float> [[MAX]], <4 x float> [[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> 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 0c9ff4cdd7808..7d812fd0265a6 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 [[X:%.*]], half [[Y:%.*]], half [[Z:%.*]]) +// 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: 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:%.*]]) #[[ATTR7:[0-9]+]] +// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract noundef half @__ocml_pown_f16(half noundef [[X:%.*]], i32 noundef [[Y:%.*]]) #[[ATTR9:[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:%.*]]) #[[ATTR7:[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:%.*]]) #[[ATTR9:[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 [[X:%.*]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.fabs.f32(float nofpclass(nan inf) [[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:%.*]]) #[[ATTR8:[0-9]+]] +// DEFAULT-NEXT: [[CALL_I1:%.*]] = tail call contract noundef float @__ocml_sin_f32(float noundef [[X:%.*]]) #[[ATTR10:[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:%.*]]) #[[ATTR8:[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:%.*]]) #[[ATTR10:[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:%.*]]) #[[ATTR8]] +// DEFAULT-NEXT: [[CALL_I1:%.*]] = tail call contract noundef float @__ocml_cos_f32(float noundef [[X:%.*]]) #[[ATTR10]] // 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:%.*]]) #[[ATTR8]] +// 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: ret float [[CALL_I1]] // extern "C" __device__ float test_cos_f32(float x) { @@ -97,10 +97,46 @@ 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 e879fec0ebe5a..df1cd716342a5 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 [[X:%.*]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.ceil.f32(float nofpclass(nan inf) [[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 [[X:%.*]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef double @llvm.ceil.f64(double nofpclass(nan inf) [[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 [[X:%.*]], float [[Y:%.*]]) +// 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: 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 [[X:%.*]], double [[Y:%.*]]) +// 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: 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 [[X:%.*]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.exp10.f32(float nofpclass(nan inf) [[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 [[X:%.*]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.exp2.f32(float nofpclass(nan inf) [[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 [[X:%.*]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.exp.f32(float nofpclass(nan inf) [[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 [[X:%.*]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.fabs.f32(float nofpclass(nan inf) [[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 [[X:%.*]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef double @llvm.fabs.f64(double nofpclass(nan inf) [[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 [[X:%.*]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.floor.f32(float nofpclass(nan inf) [[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 [[X:%.*]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef double @llvm.floor.f64(double nofpclass(nan inf) [[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 [[X:%.*]], float [[Y:%.*]], float [[Z:%.*]]) +// 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: 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 [[X:%.*]], double [[Y:%.*]], double [[Z:%.*]]) +// 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: 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 [[X:%.*]], double [[Y:%.*]], double [[Z:%.*]]) +// 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: 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 [[X:%.*]], float [[Y:%.*]]) +// 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: 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 [[X:%.*]], double [[Y:%.*]]) +// 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: 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 [[X:%.*]], float [[Y:%.*]]) +// 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: 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 [[X:%.*]], double [[Y:%.*]]) +// 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: ret double [[TMP0]] // // APPROX-LABEL: @test_fmin( @@ -1843,13 +1843,29 @@ extern "C" __device__ double test_fmod(double x, double y) { return fmod(x, y); } -// 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]] +// 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]] // // AMDGCNSPIRV-LABEL: @test_frexpf( // AMDGCNSPIRV-NEXT: entry: @@ -1863,13 +1879,29 @@ extern "C" __device__ float test_frexpf(float x, int* y) { return frexpf(x, y); } -// 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]] +// 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]] // // AMDGCNSPIRV-LABEL: @test_frexp( // AMDGCNSPIRV-NEXT: entry: @@ -2522,7 +2554,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 [[X:%.*]], i32 [[Y:%.*]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.ldexp.f32.i32(float nofpclass(nan inf) [[X:%.*]], i32 [[Y:%.*]]) // FINITEONLY-NEXT: ret float [[TMP0]] // // APPROX-LABEL: @test_ldexpf( @@ -2546,7 +2578,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 [[X:%.*]], i32 [[Y:%.*]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef double @llvm.ldexp.f64.i32(double nofpclass(nan inf) [[X:%.*]], i32 [[Y:%.*]]) // FINITEONLY-NEXT: ret double [[TMP0]] // // APPROX-LABEL: @test_ldexp( @@ -2619,7 +2651,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 [[X:%.*]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.rint.f32(float nofpclass(nan inf) [[X:%.*]]) // FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64 // FINITEONLY-NEXT: ret i64 [[CONV_I]] // @@ -2647,7 +2679,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 [[X:%.*]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.rint.f64(double nofpclass(nan inf) [[X:%.*]]) // FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64 // FINITEONLY-NEXT: ret i64 [[CONV_I]] // @@ -2675,7 +2707,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 [[X:%.*]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.round.f32(float nofpclass(nan inf) [[X:%.*]]) // FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64 // FINITEONLY-NEXT: ret i64 [[CONV_I]] // @@ -2703,7 +2735,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 [[X:%.*]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.round.f64(double nofpclass(nan inf) [[X:%.*]]) // FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64 // FINITEONLY-NEXT: ret i64 [[CONV_I]] // @@ -2730,7 +2762,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 [[X:%.*]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.log10.f32(float nofpclass(nan inf) [[X:%.*]]) // FINITEONLY-NEXT: ret float [[TMP0]] // // APPROX-LABEL: @test_log10f( @@ -2826,7 +2858,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 [[X:%.*]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.log2.f32(float nofpclass(nan inf) [[X:%.*]]) // FINITEONLY-NEXT: ret float [[TMP0]] // // APPROX-LABEL: @test_log2f( @@ -2922,7 +2954,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 [[X:%.*]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.log.f32(float nofpclass(nan inf) [[X:%.*]]) // FINITEONLY-NEXT: ret float [[TMP0]] // // APPROX-LABEL: @test_logf( @@ -2947,7 +2979,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 [[X:%.*]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.rint.f32(float nofpclass(nan inf) [[X:%.*]]) // FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64 // FINITEONLY-NEXT: ret i64 [[CONV_I]] // @@ -2975,7 +3007,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 [[X:%.*]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.rint.f64(double nofpclass(nan inf) [[X:%.*]]) // FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64 // FINITEONLY-NEXT: ret i64 [[CONV_I]] // @@ -3003,7 +3035,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 [[X:%.*]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.round.f32(float nofpclass(nan inf) [[X:%.*]]) // FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64 // FINITEONLY-NEXT: ret i64 [[CONV_I]] // @@ -3031,7 +3063,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 [[X:%.*]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.round.f64(double nofpclass(nan inf) [[X:%.*]]) // FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64 // FINITEONLY-NEXT: ret i64 [[CONV_I]] // @@ -3795,7 +3827,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 [[X:%.*]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.nearbyint.f32(float nofpclass(nan inf) [[X:%.*]]) // FINITEONLY-NEXT: ret float [[TMP0]] // // APPROX-LABEL: @test_nearbyintf( @@ -3819,7 +3851,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 [[X:%.*]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef double @llvm.nearbyint.f64(double nofpclass(nan inf) [[X:%.*]]) // FINITEONLY-NEXT: ret double [[TMP0]] // // APPROX-LABEL: @test_nearbyint( @@ -4581,7 +4613,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 [[X:%.*]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.rint.f32(float nofpclass(nan inf) [[X:%.*]]) // FINITEONLY-NEXT: ret float [[TMP0]] // // APPROX-LABEL: @test_rintf( @@ -4605,7 +4637,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 [[X:%.*]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef double @llvm.rint.f64(double nofpclass(nan inf) [[X:%.*]]) // FINITEONLY-NEXT: ret double [[TMP0]] // // APPROX-LABEL: @test_rint( @@ -4893,7 +4925,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 [[X:%.*]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.round.f32(float nofpclass(nan inf) [[X:%.*]]) // FINITEONLY-NEXT: ret float [[TMP0]] // // APPROX-LABEL: @test_roundf( @@ -4917,7 +4949,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 [[X:%.*]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef double @llvm.round.f64(double nofpclass(nan inf) [[X:%.*]]) // FINITEONLY-NEXT: ret double [[TMP0]] // // APPROX-LABEL: @test_round( @@ -4993,7 +5025,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 [[X:%.*]], i32 [[CONV_I]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.ldexp.f32.i32(float nofpclass(nan inf) [[X:%.*]], i32 [[CONV_I]]) // FINITEONLY-NEXT: ret float [[TMP0]] // // APPROX-LABEL: @test_scalblnf( @@ -5025,7 +5057,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 [[X:%.*]], i32 [[CONV_I]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef double @llvm.ldexp.f64.i32(double nofpclass(nan inf) [[X:%.*]], i32 [[CONV_I]]) // FINITEONLY-NEXT: ret double [[TMP0]] // // APPROX-LABEL: @test_scalbln( @@ -5053,7 +5085,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 [[X:%.*]], i32 [[Y:%.*]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.ldexp.f32.i32(float nofpclass(nan inf) [[X:%.*]], i32 [[Y:%.*]]) // FINITEONLY-NEXT: ret float [[TMP0]] // // APPROX-LABEL: @test_scalbnf( @@ -5077,7 +5109,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 [[X:%.*]], i32 [[Y:%.*]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef double @llvm.ldexp.f64.i32(double nofpclass(nan inf) [[X:%.*]], i32 [[Y:%.*]]) // FINITEONLY-NEXT: ret double [[TMP0]] // // APPROX-LABEL: @test_scalbn( @@ -5427,7 +5459,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 [[X:%.*]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.sqrt.f32(float nofpclass(nan inf) [[X:%.*]]) // FINITEONLY-NEXT: ret float [[TMP0]] // // APPROX-LABEL: @test_sqrtf( @@ -5451,7 +5483,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 [[X:%.*]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef double @llvm.sqrt.f64(double nofpclass(nan inf) [[X:%.*]]) // FINITEONLY-NEXT: ret double [[TMP0]] // // APPROX-LABEL: @test_sqrt( @@ -5619,7 +5651,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 [[X:%.*]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.trunc.f32(float nofpclass(nan inf) [[X:%.*]]) // FINITEONLY-NEXT: ret float [[TMP0]] // // APPROX-LABEL: @test_truncf( @@ -5643,7 +5675,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 [[X:%.*]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef double @llvm.trunc.f64(double nofpclass(nan inf) [[X:%.*]]) // FINITEONLY-NEXT: ret double [[TMP0]] // // APPROX-LABEL: @test_trunc( @@ -6163,7 +6195,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 [[X:%.*]], float [[Y:%.*]], float [[Z:%.*]]) +// 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: ret float [[TMP0]] // // APPROX-LABEL: @test__fmaf_rn( @@ -6235,7 +6267,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 [[X:%.*]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.amdgcn.rsq.f32(float nofpclass(nan inf) [[X:%.*]]) // FINITEONLY-NEXT: ret float [[TMP0]] // // APPROX-LABEL: @test___frsqrt_rn( @@ -6307,7 +6339,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 [[X:%.*]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.log10.f32(float nofpclass(nan inf) [[X:%.*]]) // FINITEONLY-NEXT: ret float [[TMP0]] // // APPROX-LABEL: @test___log10f( @@ -6331,7 +6363,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 [[X:%.*]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.amdgcn.log.f32(float nofpclass(nan inf) [[X:%.*]]) // FINITEONLY-NEXT: ret float [[TMP0]] // // APPROX-LABEL: @test___log2f( @@ -6355,7 +6387,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 [[X:%.*]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.log.f32(float nofpclass(nan inf) [[X:%.*]]) // FINITEONLY-NEXT: ret float [[TMP0]] // // APPROX-LABEL: @test___logf( @@ -6631,7 +6663,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 [[X:%.*]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef double @llvm.sqrt.f64(double nofpclass(nan inf) [[X:%.*]]) // FINITEONLY-NEXT: ret double [[TMP0]] // // APPROX-LABEL: @test___dsqrt_rn( @@ -6655,7 +6687,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 [[X:%.*]], double [[Y:%.*]], double [[Z:%.*]]) +// 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: ret double [[TMP0]] // // APPROX-LABEL: @test__fma_rn( @@ -6679,7 +6711,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 [[X:%.*]], float [[Y:%.*]]) +// 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: ret float [[TMP0]] // // APPROX-LABEL: @test_float_min( @@ -6703,7 +6735,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 [[X:%.*]], float [[Y:%.*]]) +// 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: ret float [[TMP0]] // // APPROX-LABEL: @test_float_max( @@ -6727,7 +6759,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 [[X:%.*]], double [[Y:%.*]]) +// 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: ret double [[TMP0]] // // APPROX-LABEL: @test_double_min( @@ -6751,7 +6783,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 [[X:%.*]], double [[Y:%.*]]) +// 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: 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 5beee1f681b81..c65bf16b6a937 100644 --- a/llvm/lib/Transforms/Utils/InlineFunction.cpp +++ b/llvm/lib/Transforms/Utils/InlineFunction.cpp @@ -1381,7 +1381,8 @@ 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::Alignment, Attribute::Range}; + Attribute::NonNull, Attribute::NoFPClass, + Attribute::Alignment, Attribute::Range}; for (unsigned I = 0, E = CB.arg_size(); I < E; ++I) { ValidObjParamAttrs.emplace_back(AttrBuilder{CB.getContext()}); @@ -1463,6 +1464,13 @@ 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 5bf845d5ba94b..5a102d14b5c90 100644 --- a/llvm/test/Transforms/Inline/access-attributes-prop.ll +++ b/llvm/test/Transforms/Inline/access-attributes-prop.ll @@ -750,3 +750,47 @@ 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 +} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
