Author: Tex Riddell Date: 2024-10-01T14:41:43-07:00 New Revision: b70d32789c9463d49b46b4e57c44607d1c956e4c
URL: https://github.com/llvm/llvm-project/commit/b70d32789c9463d49b46b4e57c44607d1c956e4c DIFF: https://github.com/llvm/llvm-project/commit/b70d32789c9463d49b46b4e57c44607d1c956e4c.diff LOG: [HLSL][clang] Add elementwise builtin for atan2 (p3) (#110187) This change is part of this proposal: https://discourse.llvm.org/t/rfc-all-the-math-intrinsics/78294 - Add HLSL frontend for atan2 - Add clang Builtin, map to new llvm.atan2 - SemaChecking restrict to floating point and 2 args - SemaHLSL restrict to float or half. - Add to clang ReleaseNotes.rst and LanguageExtensions.rst - Add half-float-only-errors2.hlsl for 2 arg intrinsics, and update half-float-only-errors.hlsl with scalar case for consistency - Remove fmod-errors.hlsl and pow-errors.hlsl now covered in half-float-only-errors2.hlsl Part 3 for Implement the atan2 HLSL Function #70096. Added: clang/test/CodeGenHLSL/builtins/atan2.hlsl clang/test/SemaHLSL/BuiltIns/half-float-only-errors2.hlsl Modified: clang/docs/LanguageExtensions.rst clang/docs/ReleaseNotes.rst clang/include/clang/Basic/Builtins.td clang/lib/CodeGen/CGBuiltin.cpp clang/lib/Headers/hlsl/hlsl_intrinsics.h clang/lib/Sema/SemaChecking.cpp clang/lib/Sema/SemaHLSL.cpp clang/test/CodeGen/builtins-elementwise-math.c clang/test/CodeGen/strictfp-elementwise-bulitins.cpp clang/test/Sema/aarch64-sve-vector-trig-ops.c clang/test/Sema/builtins-elementwise-math.c clang/test/Sema/riscv-rvv-vector-trig-ops.c clang/test/SemaCXX/builtins-elementwise-math.cpp clang/test/SemaHLSL/BuiltIns/half-float-only-errors.hlsl Removed: clang/test/SemaHLSL/BuiltIns/fmod-errors.hlsl clang/test/SemaHLSL/BuiltIns/pow-errors.hlsl ################################################################################ diff --git a/clang/docs/LanguageExtensions.rst b/clang/docs/LanguageExtensions.rst index ea4b4bcec55e77..c86b85d45b064c 100644 --- a/clang/docs/LanguageExtensions.rst +++ b/clang/docs/LanguageExtensions.rst @@ -660,6 +660,7 @@ Unless specified otherwise operation(±0) = ±0 and operation(±infinity) = ±in T __builtin_elementwise_asin(T x) return the arcsine of x interpreted as an angle in radians floating point types T __builtin_elementwise_acos(T x) return the arccosine of x interpreted as an angle in radians floating point types T __builtin_elementwise_atan(T x) return the arctangent of x interpreted as an angle in radians floating point types + T __builtin_elementwise_atan2(T y, T x) return the arctangent of y/x floating point types T __builtin_elementwise_sinh(T x) return the hyperbolic sine of angle x in radians floating point types T __builtin_elementwise_cosh(T x) return the hyperbolic cosine of angle x in radians floating point types T __builtin_elementwise_tanh(T x) return the hyperbolic tangent of angle x in radians floating point types diff --git a/clang/docs/ReleaseNotes.rst b/clang/docs/ReleaseNotes.rst index 34d2b584274a5f..7210ae13fbce55 100644 --- a/clang/docs/ReleaseNotes.rst +++ b/clang/docs/ReleaseNotes.rst @@ -575,6 +575,8 @@ DWARF Support in Clang Floating Point Support in Clang ------------------------------- +- Add ``__builtin_elementwise_atan2`` builtin for floating point types only. + Fixed Point Support in Clang ---------------------------- diff --git a/clang/include/clang/Basic/Builtins.td b/clang/include/clang/Basic/Builtins.td index 8090119e512fbb..b2eb747391ce07 100644 --- a/clang/include/clang/Basic/Builtins.td +++ b/clang/include/clang/Basic/Builtins.td @@ -1250,6 +1250,12 @@ def ElementwiseATan : Builtin { let Prototype = "void(...)"; } +def ElementwiseATan2 : Builtin { + let Spellings = ["__builtin_elementwise_atan2"]; + let Attributes = [NoThrow, Const, CustomTypeChecking]; + let Prototype = "void(...)"; +} + def ElementwiseBitreverse : Builtin { let Spellings = ["__builtin_elementwise_bitreverse"]; let Attributes = [NoThrow, Const, CustomTypeChecking]; diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp index da3eca73bfb575..8433b8741c3d18 100644 --- a/clang/lib/CodeGen/CGBuiltin.cpp +++ b/clang/lib/CodeGen/CGBuiltin.cpp @@ -3836,6 +3836,9 @@ RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID, case Builtin::BI__builtin_elementwise_atan: return RValue::get(emitBuiltinWithOneOverloadedType<1>( *this, E, llvm::Intrinsic::atan, "elt.atan")); + case Builtin::BI__builtin_elementwise_atan2: + return RValue::get(emitBuiltinWithOneOverloadedType<2>( + *this, E, llvm::Intrinsic::atan2, "elt.atan2")); case Builtin::BI__builtin_elementwise_ceil: return RValue::get(emitBuiltinWithOneOverloadedType<1>( *this, E, llvm::Intrinsic::ceil, "elt.ceil")); diff --git a/clang/lib/Headers/hlsl/hlsl_intrinsics.h b/clang/lib/Headers/hlsl/hlsl_intrinsics.h index 810a16d75f0228..d28f204e352de5 100644 --- a/clang/lib/Headers/hlsl/hlsl_intrinsics.h +++ b/clang/lib/Headers/hlsl/hlsl_intrinsics.h @@ -466,6 +466,36 @@ float3 atan(float3); _HLSL_BUILTIN_ALIAS(__builtin_elementwise_atan) float4 atan(float4); +//===----------------------------------------------------------------------===// +// atan2 builtins +//===----------------------------------------------------------------------===// + +/// \fn T atan2(T y, T x) +/// \brief Returns the arctangent of y/x, using the signs of the arguments to +/// determine the correct quadrant. +/// \param y The y-coordinate. +/// \param x The x-coordinate. + +#ifdef __HLSL_ENABLE_16_BIT +_HLSL_BUILTIN_ALIAS(__builtin_elementwise_atan2) +half atan2(half y, half x); +_HLSL_BUILTIN_ALIAS(__builtin_elementwise_atan2) +half2 atan2(half2 y, half2 x); +_HLSL_BUILTIN_ALIAS(__builtin_elementwise_atan2) +half3 atan2(half3 y, half3 x); +_HLSL_BUILTIN_ALIAS(__builtin_elementwise_atan2) +half4 atan2(half4 y, half4 x); +#endif + +_HLSL_BUILTIN_ALIAS(__builtin_elementwise_atan2) +float atan2(float y, float x); +_HLSL_BUILTIN_ALIAS(__builtin_elementwise_atan2) +float2 atan2(float2 y, float2 x); +_HLSL_BUILTIN_ALIAS(__builtin_elementwise_atan2) +float3 atan2(float3 y, float3 x); +_HLSL_BUILTIN_ALIAS(__builtin_elementwise_atan2) +float4 atan2(float4 y, float4 x); + //===----------------------------------------------------------------------===// // ceil builtins //===----------------------------------------------------------------------===// diff --git a/clang/lib/Sema/SemaChecking.cpp b/clang/lib/Sema/SemaChecking.cpp index 8634b54b0535d0..732024984f0d29 100644 --- a/clang/lib/Sema/SemaChecking.cpp +++ b/clang/lib/Sema/SemaChecking.cpp @@ -2755,6 +2755,7 @@ Sema::CheckBuiltinFunctionCall(FunctionDecl *FDecl, unsigned BuiltinID, // These builtins restrict the element type to floating point // types only, and take in two arguments. + case Builtin::BI__builtin_elementwise_atan2: case Builtin::BI__builtin_elementwise_fmod: case Builtin::BI__builtin_elementwise_pow: { if (BuiltinElementwiseMath(TheCall)) diff --git a/clang/lib/Sema/SemaHLSL.cpp b/clang/lib/Sema/SemaHLSL.cpp index 43cc6c81ae5cb0..5eda22f560c9d2 100644 --- a/clang/lib/Sema/SemaHLSL.cpp +++ b/clang/lib/Sema/SemaHLSL.cpp @@ -1959,6 +1959,7 @@ bool SemaHLSL::CheckBuiltinFunctionCall(unsigned BuiltinID, CallExpr *TheCall) { case Builtin::BI__builtin_elementwise_acos: case Builtin::BI__builtin_elementwise_asin: case Builtin::BI__builtin_elementwise_atan: + case Builtin::BI__builtin_elementwise_atan2: case Builtin::BI__builtin_elementwise_ceil: case Builtin::BI__builtin_elementwise_cos: case Builtin::BI__builtin_elementwise_cosh: diff --git a/clang/test/CodeGen/builtins-elementwise-math.c b/clang/test/CodeGen/builtins-elementwise-math.c index 0e53d3e141b01b..979dd74feb6029 100644 --- a/clang/test/CodeGen/builtins-elementwise-math.c +++ b/clang/test/CodeGen/builtins-elementwise-math.c @@ -441,6 +441,26 @@ void test_builtin_elementwise_atan(float f1, float f2, double d1, double d2, vf2 = __builtin_elementwise_atan(vf1); } +void test_builtin_elementwise_atan2(float f1, float f2, float f3, double d1, + double d2, double d3, float4 vf1, + float4 vf2, float4 vf3) { + // CHECK-LABEL: define void @test_builtin_elementwise_atan2( + // CHECK: [[F1:%.+]] = load float, ptr %f1.addr, align 4 + // CHECK-NEXT: [[F2:%.+]] = load float, ptr %f2.addr, align 4 + // CHECK-NEXT: call float @llvm.atan2.f32(float [[F1]], float [[F2]]) + f3 = __builtin_elementwise_atan2(f1, f2); + + // CHECK: [[D1:%.+]] = load double, ptr %d1.addr, align 8 + // CHECK-NEXT: [[D2:%.+]] = load double, ptr %d2.addr, align 8 + // CHECK-NEXT: call double @llvm.atan2.f64(double [[D1]], double [[D2]]) + d3 = __builtin_elementwise_atan2(d1, d2); + + // CHECK: [[VF1:%.+]] = load <4 x float>, ptr %vf1.addr, align 16 + // CHECK-NEXT: [[VF2:%.+]] = load <4 x float>, ptr %vf2.addr, align 16 + // CHECK-NEXT: call <4 x float> @llvm.atan2.v4f32(<4 x float> [[VF1]], <4 x float> [[VF2]]) + vf3 = __builtin_elementwise_atan2(vf1, vf2); +} + void test_builtin_elementwise_cos(float f1, float f2, double d1, double d2, float4 vf1, float4 vf2) { // CHECK-LABEL: define void @test_builtin_elementwise_cos( diff --git a/clang/test/CodeGen/strictfp-elementwise-bulitins.cpp b/clang/test/CodeGen/strictfp-elementwise-bulitins.cpp index 651f5bfc94c6c4..7396f31c4218b2 100644 --- a/clang/test/CodeGen/strictfp-elementwise-bulitins.cpp +++ b/clang/test/CodeGen/strictfp-elementwise-bulitins.cpp @@ -257,6 +257,16 @@ float4 strict_elementwise_tanh(float4 a) { return __builtin_elementwise_tanh(a); } +// CHECK-LABEL: define dso_local noundef <4 x float> @_Z24strict_elementwise_atan2Dv4_fS_ +// CHECK-SAME: (<4 x float> noundef [[A:%.*]], <4 x float> noundef [[B:%.*]]) local_unnamed_addr #[[ATTR2]] { +// CHECK-NEXT: entry: +// CHECK-NEXT: [[ELT_ATAN2:%.*]] = tail call <4 x float> @llvm.atan2.v4f32(<4 x float> [[A]], <4 x float> [[B]]) #[[ATTR4]] +// CHECK-NEXT: ret <4 x float> [[ELT_ATAN2]] +// +float4 strict_elementwise_atan2(float4 a, float4 b) { + return __builtin_elementwise_atan2(a, b); +} + // CHECK-LABEL: define dso_local noundef <4 x float> @_Z24strict_elementwise_truncDv4_f // CHECK-SAME: (<4 x float> noundef [[A:%.*]]) local_unnamed_addr #[[ATTR2]] { // CHECK-NEXT: entry: diff --git a/clang/test/CodeGenHLSL/builtins/atan2.hlsl b/clang/test/CodeGenHLSL/builtins/atan2.hlsl new file mode 100644 index 00000000000000..40796052e608fe --- /dev/null +++ b/clang/test/CodeGenHLSL/builtins/atan2.hlsl @@ -0,0 +1,59 @@ +// RUN: %clang_cc1 -finclude-default-header -x hlsl -triple \ +// RUN: dxil-pc-shadermodel6.3-library %s -fnative-half-type \ +// RUN: -emit-llvm -disable-llvm-passes -o - | FileCheck %s \ +// RUN: --check-prefixes=CHECK,NATIVE_HALF +// RUN: %clang_cc1 -finclude-default-header -x hlsl -triple \ +// RUN: spirv-unknown-vulkan-compute %s -emit-llvm -disable-llvm-passes \ +// RUN: -o - | FileCheck %s --check-prefixes=CHECK,NO_HALF + +// CHECK-LABEL: test_atan2_half +// NATIVE_HALF: call half @llvm.atan2.f16 +// NO_HALF: call float @llvm.atan2.f32 +half test_atan2_half (half p0, half p1) { + return atan2(p0, p1); +} + +// CHECK-LABEL: test_atan2_half2 +// NATIVE_HALF: call <2 x half> @llvm.atan2.v2f16 +// NO_HALF: call <2 x float> @llvm.atan2.v2f32 +half2 test_atan2_half2 (half2 p0, half2 p1) { + return atan2(p0, p1); +} + +// CHECK-LABEL: test_atan2_half3 +// NATIVE_HALF: call <3 x half> @llvm.atan2.v3f16 +// NO_HALF: call <3 x float> @llvm.atan2.v3f32 +half3 test_atan2_half3 (half3 p0, half3 p1) { + return atan2(p0, p1); +} + +// CHECK-LABEL: test_atan2_half4 +// NATIVE_HALF: call <4 x half> @llvm.atan2.v4f16 +// NO_HALF: call <4 x float> @llvm.atan2.v4f32 +half4 test_atan2_half4 (half4 p0, half4 p1) { + return atan2(p0, p1); +} + +// CHECK-LABEL: test_atan2_float +// CHECK: call float @llvm.atan2.f32 +float test_atan2_float (float p0, float p1) { + return atan2(p0, p1); +} + +// CHECK-LABEL: test_atan2_float2 +// CHECK: call <2 x float> @llvm.atan2.v2f32 +float2 test_atan2_float2 (float2 p0, float2 p1) { + return atan2(p0, p1); +} + +// CHECK-LABEL: test_atan2_float3 +// CHECK: call <3 x float> @llvm.atan2.v3f32 +float3 test_atan2_float3 (float3 p0, float3 p1) { + return atan2(p0, p1); +} + +// CHECK-LABEL: test_atan2_float4 +// CHECK: call <4 x float> @llvm.atan2.v4f32 +float4 test_atan2_float4 (float4 p0, float4 p1) { + return atan2(p0, p1); +} diff --git a/clang/test/Sema/aarch64-sve-vector-trig-ops.c b/clang/test/Sema/aarch64-sve-vector-trig-ops.c index dfa77d20e949f9..31f608bf151099 100644 --- a/clang/test/Sema/aarch64-sve-vector-trig-ops.c +++ b/clang/test/Sema/aarch64-sve-vector-trig-ops.c @@ -22,6 +22,12 @@ svfloat32_t test_atan_vv_i8mf8(svfloat32_t v) { // expected-error@-1 {{1st argument must be a vector, integer or floating point type}} } +svfloat32_t test_atan2_vv_i8mf8(svfloat32_t v) { + + return __builtin_elementwise_atan2(v, v); + // expected-error@-1 {{1st argument must be a vector, integer or floating point type}} +} + svfloat32_t test_sin_vv_i8mf8(svfloat32_t v) { return __builtin_elementwise_sin(v); diff --git a/clang/test/Sema/builtins-elementwise-math.c b/clang/test/Sema/builtins-elementwise-math.c index 26b153dd5b210b..85082e9f2ad097 100644 --- a/clang/test/Sema/builtins-elementwise-math.c +++ b/clang/test/Sema/builtins-elementwise-math.c @@ -789,6 +789,30 @@ void test_builtin_elementwise_atan(int i, float f, double d, float4 v, int3 iv, // expected-error@-1 {{1st argument must be a floating point type (was 'unsigned4' (vector of 4 'unsigned int' values))}} } +void test_builtin_elementwise_atan2(int i, float f, double d, float4 v, int3 iv, unsigned u, unsigned4 uv) { + + struct Foo s = __builtin_elementwise_atan2(f, f); + // expected-error@-1 {{initializing 'struct Foo' with an expression of incompatible type 'float'}} + + i = __builtin_elementwise_atan2(); + // expected-error@-1 {{too few arguments to function call, expected 2, have 0}} + + i = __builtin_elementwise_atan2(f); + // expected-error@-1 {{too few arguments to function call, expected 2, have 1}} + + i = __builtin_elementwise_atan2(i, i); + // expected-error@-1 {{1st argument must be a floating point type (was 'int')}} + + i = __builtin_elementwise_atan2(f, f, f); + // expected-error@-1 {{too many arguments to function call, expected 2, have 3}} + + u = __builtin_elementwise_atan2(u, u); + // expected-error@-1 {{1st argument must be a floating point type (was 'unsigned int')}} + + uv = __builtin_elementwise_atan2(uv, uv); + // expected-error@-1 {{1st argument must be a floating point type (was 'unsigned4' (vector of 4 'unsigned int' values))}} +} + void test_builtin_elementwise_tan(int i, float f, double d, float4 v, int3 iv, unsigned u, unsigned4 uv) { struct Foo s = __builtin_elementwise_tan(f); diff --git a/clang/test/Sema/riscv-rvv-vector-trig-ops.c b/clang/test/Sema/riscv-rvv-vector-trig-ops.c index f0cd5ca4a1de1f..7b27f10f2afa97 100644 --- a/clang/test/Sema/riscv-rvv-vector-trig-ops.c +++ b/clang/test/Sema/riscv-rvv-vector-trig-ops.c @@ -23,6 +23,12 @@ vfloat32mf2_t test_asin_vv_i8mf8(vfloat32mf2_t v) { // expected-error@-1 {{1st argument must be a vector, integer or floating point type}} } +vfloat32mf2_t test_atan2_vv_i8mf8(vfloat32mf2_t v) { + + return __builtin_elementwise_atan2(v, v); + // expected-error@-1 {{1st argument must be a vector, integer or floating point type}} +} + vfloat32mf2_t test_sin_vv_i8mf8(vfloat32mf2_t v) { return __builtin_elementwise_sin(v); diff --git a/clang/test/SemaCXX/builtins-elementwise-math.cpp b/clang/test/SemaCXX/builtins-elementwise-math.cpp index 5910796c5d2983..a857d4fb18322f 100644 --- a/clang/test/SemaCXX/builtins-elementwise-math.cpp +++ b/clang/test/SemaCXX/builtins-elementwise-math.cpp @@ -146,6 +146,13 @@ void test_builtin_elementwise_atan() { static_assert(!is_const<decltype(__builtin_elementwise_atan(b))>::value); } +void test_builtin_elementwise_atan2() { + const float a = 42.0; + float b = 42.3; + static_assert(!is_const<decltype(__builtin_elementwise_atan2(a, a))>::value); + static_assert(!is_const<decltype(__builtin_elementwise_atan2(b, b))>::value); +} + void test_builtin_elementwise_tan() { const float a = 42.0; float b = 42.3; diff --git a/clang/test/SemaHLSL/BuiltIns/fmod-errors.hlsl b/clang/test/SemaHLSL/BuiltIns/fmod-errors.hlsl deleted file mode 100644 index e4fa609dd6a05d..00000000000000 --- a/clang/test/SemaHLSL/BuiltIns/fmod-errors.hlsl +++ /dev/null @@ -1,22 +0,0 @@ - -// RUN: %clang_cc1 -finclude-default-header -triple dxil-pc-shadermodel6.6-library %s -fnative-half-type -emit-llvm-only -disable-llvm-passes -verify -verify-ignore-unexpected - -float builtin_bool_to_float_type_promotion(bool p1, bool p2) { - return __builtin_elementwise_fmod(p1, p2); - // expected-error@-1 {{1st argument must be a vector, integer or floating point type (was 'bool')}} -} - -float2 builtin_fmod_int2_to_float2_promotion(int2 p1, int2 p2) { - return __builtin_elementwise_fmod(p1, p2); - // expected-error@-1 {{1st argument must be a floating point type (was 'int2' (aka 'vector<int, 2>'))}} -} - -half builtin_fmod_double_type (double p0, double p1) { - return __builtin_elementwise_fmod(p0, p1); - // expected-error@-1 {{passing 'double' to parameter of incompatible type 'float'}} -} - -half builtin_fmod_double2_type (double2 p0, double2 p1) { - return __builtin_elementwise_fmod(p0, p1); - // expected-error@-1 {{passing 'double2' (aka 'vector<double, 2>') to parameter of incompatible type '__attribute__((__vector_size__(2 * sizeof(float)))) float' (vector of 2 'float' values)}} -} diff --git a/clang/test/SemaHLSL/BuiltIns/half-float-only-errors.hlsl b/clang/test/SemaHLSL/BuiltIns/half-float-only-errors.hlsl index 1841f60568b568..0c1f7022bad609 100644 --- a/clang/test/SemaHLSL/BuiltIns/half-float-only-errors.hlsl +++ b/clang/test/SemaHLSL/BuiltIns/half-float-only-errors.hlsl @@ -18,7 +18,12 @@ // RUN: %clang_cc1 -finclude-default-header -triple dxil-pc-shadermodel6.6-library %s -fnative-half-type -emit-llvm-only -disable-llvm-passes -verify -DTEST_FUNC=__builtin_elementwise_tanh // RUN: %clang_cc1 -finclude-default-header -triple dxil-pc-shadermodel6.6-library %s -fnative-half-type -emit-llvm-only -disable-llvm-passes -verify -DTEST_FUNC=__builtin_elementwise_trunc -double2 test_double_builtin(double2 p0) { +double test_double_builtin(double p0) { + return TEST_FUNC(p0); + // expected-error@-1 {{passing 'double' to parameter of incompatible type 'float'}} +} + +double2 test_vec_double_builtin(double2 p0) { return TEST_FUNC(p0); // expected-error@-1 {{passing 'double2' (aka 'vector<double, 2>') to parameter of incompatible type '__attribute__((__vector_size__(2 * sizeof(float)))) float' (vector of 2 'float' values)}} } diff --git a/clang/test/SemaHLSL/BuiltIns/half-float-only-errors2.hlsl b/clang/test/SemaHLSL/BuiltIns/half-float-only-errors2.hlsl new file mode 100644 index 00000000000000..bfbd8b28257a3b --- /dev/null +++ b/clang/test/SemaHLSL/BuiltIns/half-float-only-errors2.hlsl @@ -0,0 +1,13 @@ +// RUN: %clang_cc1 -finclude-default-header -triple dxil-pc-shadermodel6.6-library %s -fnative-half-type -emit-llvm-only -disable-llvm-passes -verify -DTEST_FUNC=__builtin_elementwise_atan2 +// RUN: %clang_cc1 -finclude-default-header -triple dxil-pc-shadermodel6.6-library %s -fnative-half-type -emit-llvm-only -disable-llvm-passes -verify -DTEST_FUNC=__builtin_elementwise_fmod +// RUN: %clang_cc1 -finclude-default-header -triple dxil-pc-shadermodel6.6-library %s -fnative-half-type -emit-llvm-only -disable-llvm-passes -verify -DTEST_FUNC=__builtin_elementwise_pow + +double test_double_builtin(double p0, double p1) { + return TEST_FUNC(p0, p1); + // expected-error@-1 {{passing 'double' to parameter of incompatible type 'float'}} +} + +double2 test_vec_double_builtin(double2 p0, double2 p1) { + return TEST_FUNC(p0, p1); + // expected-error@-1 {{passing 'double2' (aka 'vector<double, 2>') to parameter of incompatible type '__attribute__((__vector_size__(2 * sizeof(float)))) float' (vector of 2 'float' values)}} +} diff --git a/clang/test/SemaHLSL/BuiltIns/pow-errors.hlsl b/clang/test/SemaHLSL/BuiltIns/pow-errors.hlsl deleted file mode 100644 index 5a2c9a3ef62ff5..00000000000000 --- a/clang/test/SemaHLSL/BuiltIns/pow-errors.hlsl +++ /dev/null @@ -1,6 +0,0 @@ -// RUN: %clang_cc1 -finclude-default-header -triple dxil-pc-shadermodel6.6-library %s -fnative-half-type -emit-llvm-only -disable-llvm-passes -verify - -double2 test_double_builtin(double2 p0, double2 p1) { - return __builtin_elementwise_pow(p0,p1); - // expected-error@-1 {{passing 'double2' (aka 'vector<double, 2>') to parameter of incompatible type '__attribute__((__vector_size__(2 * sizeof(float)))) float' (vector of 2 'float' values)}} -} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
