clang/OpenCL: Add inline implementations of sqrt in builtin header

[arsenm]

Summary

We want the !fpmath metadata to be attached to the sqrt intrinsic to
make it to the backend lowering. Emit an available_externally
definition which uses the builtin, which emits the !fpmath.

Fixes #64264

[arsenm]

Move to base header and fix test. update_cc_test_checks --include-generated-funcs does a terrible job here so write manually checks

[Anastasia]

clang/lib/Headers/opencl-c-base.h

832

Is this a generic implementation enough? Would some targets not need to do something different for this built-in?

Ideally this header is to be kept light so I am a bit worried about adding definitions of the functions here. Otherwise we will end up in the same situation as we one day were with opencl-c.h. So could these be left there instead? It might be good to check with @svenvh if TableGen header has already a way to do this function forwarding or can be extended to do such a thing. Then it would be implementable in the both header mechanisms. I don't know if Sven has some other ideas or opinions...

[svenvh]

If we go with this approach, please also remove sqrt from clang/lib/Sema/OpenCLBuiltins.td (and ideally add a comment pointing out that sqrt is handled in opencl-c-base.h)

clang/lib/Headers/opencl-c-base.h

832

We did already discuss this a bit on the GitHub issue: https://github.com/llvm/llvm-project/issues/64264

856–858

[arsenm]

clang/lib/Headers/opencl-c-base.h

832

As I mentioned on the ticket, it's only this one case so I'm not worried about adding a lot more to the base header. I think we can start by assuming llvm.sqrt always works correctly, I don't want to add more complexity to handle this case without a specific reason

[Anastasia]

clang/lib/Headers/opencl-c-base.h

832

As I mentioned on the ticket, it's only this one case so I'm not worried about adding a lot more to the base header.

This is how things normally start. Someone else might want to continue this approach because it is already there.

I think we can start by assuming llvm.sqrt always works correctly, I don't want to add more complexity to handle this case without a specific reason

Do you mean it would apply to all implementations? What I am missing here is why it is required to be in the headers? Is this because it needs to be inlined or is it because the compiler must see __builtin_elementwise_sqrt with the surrounding code where it is called from?

[arsenm]

clang/lib/Headers/opencl-c-base.h

832

Yes, I would expect any llvm consumer to correctly lower llvm.sqrt, and such an implementation would be correctly rounded and pass conformance with -cl-fp32-correctly-rounded-divide-sqrt

The goal is to get !fpmath attached to an llvm.sqrt call. The way this currently happens is based on the language, it gets emitted from the various __builtin_sqrt* calls

[arsenm]

ping

The alternative is to directly put the !fpmath on the sqrt call sites but I have no idea how to do that

[arsenm]

ping

[arsenm]

ping

[Anastasia]

If we think there are no better alternatives and implementation is generic enough for every vendor, LGTM!

[arsenm]

In D156743#4644285, @Anastasia wrote:

If we think there are no better alternatives and implementation is generic enough for every vendor, LGTM!

You could argue annotating the raw callsite is better but I don't know how to implement that

[arsenm]

15e0fe0b6122e32657b98daf74a1fce028d2e5bf

clang/docs/LanguageExtensions.rst

  T __builtin_elementwise_floor(T x)          return the largest integral value less than or equal to x        floating point types
  T __builtin_elementwise_log(T x)            return the natural logarithm of x                                floating point types
  T __builtin_elementwise_log2(T x)           return the base 2 logarithm of x                                 floating point types
  T __builtin_elementwise_log10(T x)          return the base 10 logarithm of x                                floating point types
  T __builtin_elementwise_pow(T x, T y)       return x raised to the power of y                                floating point types
  T __builtin_elementwise_bitreverse(T x)     return the integer represented after reversing the bits of x     integer types
  T __builtin_elementwise_exp(T x)            returns the base-e exponential, e^x, of the specified value      floating point types
  T __builtin_elementwise_exp2(T x)           returns the base-2 exponential, 2^x, of the specified value      floating point types
+
+ T __builtin_elementwise_sqrt(T x)           return the square root of a floating-point number                floating point types
  T __builtin_elementwise_roundeven(T x)      round x to the nearest integer value in floating point format,   floating point types
                                              rounding halfway cases to even (that is, to the nearest value
                                              that is an even integer), regardless of the current rounding
                                              direction.
  T __builtin_elementwise_round(T x)          round x to the nearest  integer value in floating point format,      floating point types
                                              rounding halfway cases away from zero, regardless of the
                                              current rounding direction. May raise floating-point
                                              exceptions.

clang/include/clang/Basic/Builtins.def

 BUILTIN(__builtin_elementwise_log2, "v.", "nct")
 BUILTIN(__builtin_elementwise_log10, "v.", "nct")
 BUILTIN(__builtin_elementwise_pow, "v.", "nct")
 BUILTIN(__builtin_elementwise_roundeven, "v.", "nct")
 BUILTIN(__builtin_elementwise_round, "v.", "nct")
 BUILTIN(__builtin_elementwise_rint, "v.", "nct")
 BUILTIN(__builtin_elementwise_nearbyint, "v.", "nct")
 BUILTIN(__builtin_elementwise_sin, "v.", "nct")
+BUILTIN(__builtin_elementwise_sqrt, "v.", "nct")
 BUILTIN(__builtin_elementwise_trunc, "v.", "nct")
 BUILTIN(__builtin_elementwise_canonicalize, "v.", "nct")
 BUILTIN(__builtin_elementwise_copysign, "v.", "nct")
 BUILTIN(__builtin_elementwise_fma, "v.", "nct")
 BUILTIN(__builtin_elementwise_add_sat, "v.", "nct")
 BUILTIN(__builtin_elementwise_sub_sat, "v.", "nct")
 BUILTIN(__builtin_reduce_max, "v.", "nct")
 BUILTIN(__builtin_reduce_min, "v.", "nct")

clang/lib/CodeGen/CGBuiltin.cpp

if (FD->hasAttr<ConstAttr>() ||
 
     case Builtin::BIsqrt:
     case Builtin::BIsqrtf:
     case Builtin::BIsqrtl:
     case Builtin::BI__builtin_sqrt:
     case Builtin::BI__builtin_sqrtf:
     case Builtin::BI__builtin_sqrtf16:
     case Builtin::BI__builtin_sqrtl:
-    case Builtin::BI__builtin_sqrtf128: {
+    case Builtin::BI__builtin_sqrtf128:
+    case Builtin::BI__builtin_elementwise_sqrt: {
       llvm::Value *Call = emitUnaryMaybeConstrainedFPBuiltin(
           *this, E, Intrinsic::sqrt, Intrinsic::experimental_constrained_sqrt);
       SetSqrtFPAccuracy(Call);
       return RValue::get(Call);
     }
     case Builtin::BItrunc:
     case Builtin::BItruncf:
     case Builtin::BItruncl:

clang/lib/Headers/opencl-c-base.h

 
 #define CLK_AVC_IME_RESULT_SINGLE_REFERENCE_STREAMOUT_INITIALIZE_INTEL 0x0
 #define CLK_AVC_IME_RESULT_SINGLE_REFERENCE_STREAMIN_INITIALIZE_INTEL 0x0
 #define CLK_AVC_IME_RESULT_DUAL_REFERENCE_STREAMOUT_INITIALIZE_INTEL 0x0
 #define CLK_AVC_IME_RESULT_DUAL_REFERENCE_STREAMIN_INITIALIZE_INTEL 0x0
 
 #endif // cl_intel_device_side_avc_motion_estimation
 
+/**
+ * Compute square root.
+ *
+ * Provide inline implementations using the builtin so that we get appropriate
+ * !fpmath based on -cl-fp32-correctly-rounded-divide-sqrt, attached to
+ * llvm.sqrt. The implementation should still provide an external definition.
+ */
+#define __ovld __attribute__((overloadable))
+#define __cnfn __attribute__((const))
+
+inline float __ovld __cnfn sqrt(float __x) {
+  return __builtin_elementwise_sqrt(__x);
+}
+
+inline float2 __ovld __cnfn sqrt(float2 __x) {
+  return __builtin_elementwise_sqrt(__x);
+}
+
+inline float3 __ovld __cnfn sqrt(float3 __x) {
+  return __builtin_elementwise_sqrt(__x);
+}
+
+inline float4 __ovld __cnfn sqrt(float4 __x) {
+  return __builtin_elementwise_sqrt(__x);
+}
+
+inline float8 __ovld __cnfn sqrt(float8 __x) {
+  return __builtin_elementwise_sqrt(__x);
+}
+
+inline float16 __ovld __cnfn sqrt(float16 __x) {
+  return __builtin_elementwise_sqrt(__x);
+}
+
+// We only really want to define the float variants here. However bad things
+// seem to happen with -fdeclare-opencl-builtins and splitting the handling of
+// different overloads.
+
+#ifdef cl_khr_fp64
+double __ovld __cnfn sqrt(double);
+double2 __ovld __cnfn sqrt(double2);
+double3 __ovld __cnfn sqrt(double3);
+double4 __ovld __cnfn sqrt(double4);
+double8 __ovld __cnfn sqrt(double8);
+double16 __ovld __cnfn sqrt(double16);
+#endif //cl_khr_fp64
+#ifdef cl_khr_fp16
+half __ovld __cnfn sqrt(half);
+half2 __ovld __cnfn sqrt(half2);
+half3 __ovld __cnfn sqrt(half3);
+half4 __ovld __cnfn sqrt(half4);
+half8 __ovld __cnfn sqrt(half8);
+half16 __ovld __cnfn sqrt(half16);
+#endif //cl_khr_fp16
+
+#undef __cnfn
+#undef __ovld
+
 // Disable any extensions we may have enabled previously.
 #pragma OPENCL EXTENSION all : disable
 
 #endif //_OPENCL_BASE_H_

clang/lib/Headers/opencl-c.h

 half2 __ovld __cnfn sinpi(half2);
 half3 __ovld __cnfn sinpi(half3);
 half4 __ovld __cnfn sinpi(half4);
 half8 __ovld __cnfn sinpi(half8);
 half16 __ovld __cnfn sinpi(half16);
 #endif //cl_khr_fp16
 
 /**
- * Compute square root.
- */
-float __ovld __cnfn sqrt(float);
-float2 __ovld __cnfn sqrt(float2);
-float3 __ovld __cnfn sqrt(float3);
-float4 __ovld __cnfn sqrt(float4);
-float8 __ovld __cnfn sqrt(float8);
-float16 __ovld __cnfn sqrt(float16);
-#ifdef cl_khr_fp64
-double __ovld __cnfn sqrt(double);
-double2 __ovld __cnfn sqrt(double2);
-double3 __ovld __cnfn sqrt(double3);
-double4 __ovld __cnfn sqrt(double4);
-double8 __ovld __cnfn sqrt(double8);
-double16 __ovld __cnfn sqrt(double16);
-#endif //cl_khr_fp64
-#ifdef cl_khr_fp16
-half __ovld __cnfn sqrt(half);
-half2 __ovld __cnfn sqrt(half2);
-half3 __ovld __cnfn sqrt(half3);
-half4 __ovld __cnfn sqrt(half4);
-half8 __ovld __cnfn sqrt(half8);
-half16 __ovld __cnfn sqrt(half16);
-#endif //cl_khr_fp16
-
-/**
  * Compute tangent.
  */
 float __ovld __cnfn tan(float);
 float2 __ovld __cnfn tan(float2);
 float3 __ovld __cnfn tan(float3);
 float4 __ovld __cnfn tan(float4);
 float8 __ovld __cnfn tan(float8);
 float16 __ovld __cnfn tan(float16);

clang/lib/Sema/SemaChecking.cpp

#include "clang/Basic/Builtins.def"
   case Builtin::BI__builtin_elementwise_log:
   case Builtin::BI__builtin_elementwise_log2:
   case Builtin::BI__builtin_elementwise_log10:
   case Builtin::BI__builtin_elementwise_roundeven:
   case Builtin::BI__builtin_elementwise_round:
   case Builtin::BI__builtin_elementwise_rint:
   case Builtin::BI__builtin_elementwise_nearbyint:
   case Builtin::BI__builtin_elementwise_sin:
+  case Builtin::BI__builtin_elementwise_sqrt:
   case Builtin::BI__builtin_elementwise_trunc:
   case Builtin::BI__builtin_elementwise_canonicalize: {
     if (PrepareBuiltinElementwiseMathOneArgCall(TheCall))
       return ExprError();
 
     QualType ArgTy = TheCall->getArg(0)->getType();
     if (checkFPMathBuiltinElementType(*this, TheCall->getArg(0)->getBeginLoc(),
                                       ArgTy, 1))

clang/test/CodeGen/builtins-elementwise-math.c

void test_builtin_elementwise_sin(float f1, float f2, double d1, double d2,
   // CHECK-NEXT: call double @llvm.sin.f64(double [[D1]])
   d2 = __builtin_elementwise_sin(d1);
 
   // CHECK:      [[VF1:%.+]] = load <4 x float>, ptr %vf1.addr, align 16
   // CHECK-NEXT: call <4 x float> @llvm.sin.v4f32(<4 x float> [[VF1]])
   vf2 = __builtin_elementwise_sin(vf1);
 }
 
+void test_builtin_elementwise_sqrt(float f1, float f2, double d1, double d2,
+                                  float4 vf1, float4 vf2) {
+  // CHECK-LABEL: define void @test_builtin_elementwise_sqrt(
+  // CHECK:      [[F1:%.+]] = load float, ptr %f1.addr, align 4
+  // CHECK-NEXT:  call float @llvm.sqrt.f32(float [[F1]])
+  f2 = __builtin_elementwise_sqrt(f1);
+
+  // CHECK:      [[D1:%.+]] = load double, ptr %d1.addr, align 8
+  // CHECK-NEXT: call double @llvm.sqrt.f64(double [[D1]])
+  d2 = __builtin_elementwise_sqrt(d1);
+
+  // CHECK:      [[VF1:%.+]] = load <4 x float>, ptr %vf1.addr, align 16
+  // CHECK-NEXT: call <4 x float> @llvm.sqrt.v4f32(<4 x float> [[VF1]])
+  vf2 = __builtin_elementwise_sqrt(vf1);
+}
+
 void test_builtin_elementwise_trunc(float f1, float f2, double d1, double d2,
                                     float4 vf1, float4 vf2) {
   // CHECK-LABEL: define void @test_builtin_elementwise_trunc(
   // CHECK:      [[F1:%.+]] = load float, ptr %f1.addr, align 4
   // CHECK-NEXT:  call float @llvm.trunc.f32(float [[F1]])
   f2 = __builtin_elementwise_trunc(f1);
 
   // CHECK:      [[D1:%.+]] = load double, ptr %d1.addr, align 8

clang/test/CodeGen/strictfp-elementwise-bulitins.cpp

 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    [[ELT_SIN:%.*]] = tail call <4 x float> @llvm.sin.v4f32(<4 x float> [[A]]) #[[ATTR4]]
 // CHECK-NEXT:    ret <4 x float> [[ELT_SIN]]
 //
 float4 strict_elementwise_sin(float4 a) {
   return __builtin_elementwise_sin(a);
 }
 
+// CHECK-LABEL: define dso_local noundef <4 x float> @_Z23strict_elementwise_sqrtDv4_f
+// CHECK-SAME: (<4 x float> noundef [[A:%.*]]) local_unnamed_addr #[[ATTR0]] {
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[TMP0:%.*]] = tail call <4 x float> @llvm.experimental.constrained.sqrt.v4f32(<4 x float> [[A]], metadata !"round.dynamic", metadata !"fpexcept.strict") #[[ATTR4]]
+// CHECK-NEXT:    ret <4 x float> [[TMP0]]
+//
+float4 strict_elementwise_sqrt(float4 a) {
+  return __builtin_elementwise_sqrt(a);
+}
+
 // 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:
 // CHECK-NEXT:    [[ELT_TRUNC:%.*]] = tail call <4 x float> @llvm.trunc.v4f32(<4 x float> [[A]]) #[[ATTR4]]
 // CHECK-NEXT:    ret <4 x float> [[ELT_TRUNC]]
 //
 float4 strict_elementwise_trunc(float4 a) {
   return __builtin_elementwise_trunc(a);

clang/test/CodeGenCUDA/correctly-rounded-div.cu

 
 // COMMON-LABEL: @_Z12dpscalarsqrt
 // COMMON: call contract double @llvm.sqrt.f64(double %{{.+}}){{$}}
 // COMMON-NOT: !fpmath
 __device__ double dpscalarsqrt(double a) {
   return __builtin_sqrt(a);
 }
 
+// COMMON-LABEL: @_Z28test_builtin_elementwise_f32f
+// NCRDIV: call contract float @llvm.sqrt.f32(float %{{.+}}), !fpmath ![[MD:[0-9]+]]
+// CRDIV: call contract float @llvm.sqrt.f32(float %{{.+}}){{$}}
+__device__ float test_builtin_elementwise_f32(float a) {
+  return __builtin_elementwise_sqrt(a);
+}
+
+// COMMON-LABEL: @_Z28test_builtin_elementwise_f64d
+// COMMON: call contract double @llvm.sqrt.f64(double %{{.+}}){{$}}
+// COMMON-NOT: !fpmath
+__device__ double test_builtin_elementwise_f64(double a) {
+  return __builtin_elementwise_sqrt(a);
+}
+
 // NCRSQRT: ![[MD]] = !{float 2.500000e+00}

clang/test/CodeGenOpenCL/fpmath.cl

 
 float spscalarsqrt(float a) {
   // CHECK-LABEL: @spscalarsqrt
   // NODIVOPT: call float @llvm.sqrt.f32(float %{{.+}}), !fpmath ![[MD_SQRT:[0-9]+]]
   // DIVOPT: call float @llvm.sqrt.f32(float %{{.+}}){{$}}
   return __builtin_sqrtf(a);
 }
 
+float elementwise_sqrt_f32(float a) {
+  // CHECK-LABEL: @elementwise_sqrt_f32
+  // NODIVOPT: call float @llvm.sqrt.f32(float %{{.+}}), !fpmath ![[MD_SQRT:[0-9]+]]
+  // DIVOPT: call float @llvm.sqrt.f32(float %{{.+}}){{$}}
+  return __builtin_elementwise_sqrt(a);
+}
+
+float4 elementwise_sqrt_v4f32(float4 a) {
+  // CHECK-LABEL: @elementwise_sqrt_v4f32
+  // NODIVOPT: call <4 x float> @llvm.sqrt.v4f32(<4 x float> %{{.+}}), !fpmath ![[MD_SQRT:[0-9]+]]
+  // DIVOPT: call <4 x float> @llvm.sqrt.v4f32(<4 x float> %{{.+}}){{$}}
+  return __builtin_elementwise_sqrt(a);
+}
+
+
 #if __OPENCL_C_VERSION__ >=120
 void printf(constant char* fmt, ...);
 
 void testdbllit(long *val) {
   // CHECK-FLT: float noundef 2.000000e+01
   // CHECK-DBL: double noundef 2.000000e+01
   printf("%f", 20.0);
 }

 }
 
 double dpscalarsqrt(double a) {
   // CHECK-LABEL: @dpscalarsqrt
   // CHECK: call double @llvm.sqrt.f64(double %{{.+}}){{$}}
   return __builtin_sqrt(a);
 }
 
+double elementwise_sqrt_f64(double a) {
+  // CHECK-LABEL: @elementwise_sqrt_f64
+  // CHECK: call double @llvm.sqrt.f64(double %{{.+}}){{$}}
+  return __builtin_elementwise_sqrt(a);
+}
+
+double4 elementwise_sqrt_v4f64(double4 a) {
+  // CHECK-LABEL: @elementwise_sqrt_v4f64
+  // CHECK: call <4 x double> @llvm.sqrt.v4f64(<4 x double> %{{.+}}){{$}}
+  return __builtin_elementwise_sqrt(a);
+}
+
 #endif
 
 // NODIVOPT: ![[MD_FDIV]] = !{float 2.500000e+00}
 // NODIVOPT: ![[MD_SQRT]] = !{float 3.000000e+00}

clang/test/CodeGenOpenCL/sqrt-fpmath.cl

+// Test that float variants of sqrt are emitted as available_externally inline
+// definitions that call the sqrt intrinsic with appropriate !fpmath metadata
+// depending on -cl-fp32-correctly-rounded-divide-sqrt
+
+// Test with -fdeclare-opencl-builtins
+// RUN: %clang_cc1 -disable-llvm-passes -triple amdgcn-unknown-unknown -fdeclare-opencl-builtins -finclude-default-header -S -emit-llvm -o - %s | FileCheck -check-prefixes=CHECK,DEFAULT %s
+// RUN: %clang_cc1 -disable-llvm-passes -triple amdgcn-unknown-unknown -fdeclare-opencl-builtins -finclude-default-header -cl-fp32-correctly-rounded-divide-sqrt -S -emit-llvm -o - %s | FileCheck -check-prefixes=CHECK,CORRECTLYROUNDED %s
+
+// RUN: %clang_cc1 -disable-llvm-passes -triple amdgcn-unknown-unknown -fdeclare-opencl-builtins -finclude-default-header -cl-unsafe-math-optimizations -S -emit-llvm -o - %s | FileCheck -check-prefixes=CHECK,DEFAULT-UNSAFE %s
+// RUN: %clang_cc1 -disable-llvm-passes -triple amdgcn-unknown-unknown -fdeclare-opencl-builtins -finclude-default-header -cl-fp32-correctly-rounded-divide-sqrt -cl-unsafe-math-optimizations -S -emit-llvm -o - %s | FileCheck -check-prefixes=CHECK,CORRECTLYROUNDED-UNSAFE %s
+
+// Test without -fdeclare-opencl-builtins
+// RUN: %clang_cc1 -disable-llvm-passes -triple amdgcn-unknown-unknown -finclude-default-header -S -emit-llvm -o - %s | FileCheck -check-prefixes=CHECK,DEFAULT %s
+// RUN: %clang_cc1 -disable-llvm-passes -triple amdgcn-unknown-unknown -finclude-default-header -cl-fp32-correctly-rounded-divide-sqrt -S -emit-llvm -o - %s | FileCheck -check-prefixes=CHECK,CORRECTLYROUNDED %s
+
+// RUN: %clang_cc1 -disable-llvm-passes -triple amdgcn-unknown-unknown -finclude-default-header -cl-unsafe-math-optimizations -S -emit-llvm -o - %s | FileCheck -check-prefixes=CHECK,DEFAULT-UNSAFE %s
+// RUN: %clang_cc1 -disable-llvm-passes -triple amdgcn-unknown-unknown -finclude-default-header -cl-fp32-correctly-rounded-divide-sqrt -cl-unsafe-math-optimizations -S -emit-llvm -o - %s | FileCheck -check-prefixes=CHECK,CORRECTLYROUNDED-UNSAFE %s
+
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+// CHECK-LABEL: define {{.*}} float @call_sqrt_f32(
+// CHECK: call {{.*}} float @_Z4sqrtf(float noundef %{{.+}}) #{{[0-9]+$}}
+float call_sqrt_f32(float x) {
+  return sqrt(x);
+}
+
+// CHECK-LABEL: define available_externally float @_Z4sqrtf(float noundef %__x)
+// DEFAULT: call float @llvm.sqrt.f32(float %{{.+}}), !fpmath [[$FPMATH:![0-9]+]]{{$}}
+// CORRECTLYROUNDED: call float @llvm.sqrt.f32(float %{{.+}}){{$}}
+
+// DEFAULT-UNSAFE: call reassoc nsz arcp contract afn float @llvm.sqrt.f32(float %{{.+}}), !fpmath [[$FPMATH:![0-9]+]]{{$}}
+// CORRECTLYROUNDED-UNSAFE: call reassoc nsz arcp contract afn float @llvm.sqrt.f32(float %{{.+}}){{$}}
+
+// CHECK-LABEL: define {{.*}} <2 x float> @call_sqrt_v2f32(
+// CHECK: call {{.*}} <2 x float> @_Z4sqrtDv2_f(<2 x float> noundef %{{.*}}) #{{[0-9]+$}}
+float2 call_sqrt_v2f32(float2 x) {
+  return sqrt(x);
+}
+
+// CHECK-LABEL: define available_externally <2 x float> @_Z4sqrtDv2_f(<2 x float> noundef %__x)
+// DEFAULT: call <2 x float> @llvm.sqrt.v2f32(<2 x float> %{{.+}}), !fpmath [[$FPMATH:![0-9]+]]{{$}}
+// CORRECTLYROUNDED: call <2 x float> @llvm.sqrt.v2f32(<2 x float> %{{.+}}){{$}}
+
+// DEFAULT-UNSAFE: call reassoc nsz arcp contract afn <2 x float> @llvm.sqrt.v2f32(<2 x float> %{{.+}}), !fpmath [[$FPMATH:![0-9]+]]{{$}}
+// CORRECTLYROUNDED-UNSAFE: call reassoc nsz arcp contract afn <2 x float> @llvm.sqrt.v2f32(<2 x float> %{{.+}}){{$}}
+
+// CHECK-LABEL: define {{.*}} <3 x float> @call_sqrt_v3f32(
+// CHECK: call {{.*}} <3 x float> @_Z4sqrtDv3_f(<3 x float> noundef %{{.*}}) #{{[0-9]+$}}
+float3 call_sqrt_v3f32(float3 x) {
+  return sqrt(x);
+}
+
+// CHECK-LABEL: define available_externally <3 x float> @_Z4sqrtDv3_f(<3 x float> noundef %__x)
+// DEFAULT: call <3 x float> @llvm.sqrt.v3f32(<3 x float> %{{.+}}), !fpmath [[$FPMATH:![0-9]+]]{{$}}
+// CORRECTLYROUNDED: call <3 x float> @llvm.sqrt.v3f32(<3 x float> %{{.+}}){{$}}
+
+// DEFAULT-UNSAFE: call reassoc nsz arcp contract afn <3 x float> @llvm.sqrt.v3f32(<3 x float> %{{.+}}), !fpmath [[$FPMATH:![0-9]+]]{{$}}
+// CORRECTLYROUNDED-UNSAFE: call reassoc nsz arcp contract afn <3 x float> @llvm.sqrt.v3f32(<3 x float> %{{.+}}){{$}}
+
+
+// CHECK-LABEL: define {{.*}} <4 x float> @call_sqrt_v4f32(
+// CHECK: call {{.*}} <4 x float> @_Z4sqrtDv4_f(<4 x float> noundef %{{.*}}) #{{[0-9]+$}}
+float4 call_sqrt_v4f32(float4 x) {
+  return sqrt(x);
+}
+
+// CHECK-LABEL: define available_externally <4 x float> @_Z4sqrtDv4_f(<4 x float> noundef %__x)
+// DEFAULT: call <4 x float> @llvm.sqrt.v4f32(<4 x float> %{{.+}}), !fpmath [[$FPMATH:![0-9]+]]{{$}}
+// CORRECTLYROUNDED: call <4 x float> @llvm.sqrt.v4f32(<4 x float> %{{.+}}){{$}}
+
+// DEFAULT-UNSAFE: call reassoc nsz arcp contract afn <4 x float> @llvm.sqrt.v4f32(<4 x float> %{{.+}}), !fpmath [[$FPMATH:![0-9]+]]{{$}}
+// CORRECTLYROUNDED-UNSAFE: call reassoc nsz arcp contract afn <4 x float> @llvm.sqrt.v4f32(<4 x float> %{{.+}}){{$}}
+
+// CHECK-LABEL: define {{.*}} <8 x float> @call_sqrt_v8f32(
+// CHECK: call {{.*}} <8 x float> @_Z4sqrtDv8_f(<8 x float> noundef %{{.*}}) #{{[0-9]+$}}
+float8 call_sqrt_v8f32(float8 x) {
+  return sqrt(x);
+}
+
+// CHECK-LABEL: define available_externally <8 x float> @_Z4sqrtDv8_f(<8 x float> noundef %__x)
+// DEFAULT: call <8 x float> @llvm.sqrt.v8f32(<8 x float> %{{.+}}), !fpmath [[$FPMATH:![0-9]+]]{{$}}
+// CORRECTLYROUNDED: call <8 x float> @llvm.sqrt.v8f32(<8 x float> %{{.+}}){{$}}
+
+// DEFAULT-UNSAFE: call reassoc nsz arcp contract afn <8 x float> @llvm.sqrt.v8f32(<8 x float> %{{.+}}), !fpmath [[$FPMATH:![0-9]+]]{{$}}
+// CORRECTLYROUNDED-UNSAFE: call reassoc nsz arcp contract afn <8 x float> @llvm.sqrt.v8f32(<8 x float> %{{.+}}){{$}}
+
+
+// CHECK-LABEL: define {{.*}} <16 x float> @call_sqrt_v16f32(
+// CHECK: call {{.*}} <16 x float> @_Z4sqrtDv16_f(<16 x float> noundef %{{.*}}) #{{[0-9]+$}}
+float16 call_sqrt_v16f32(float16 x) {
+  return sqrt(x);
+}
+
+// CHECK-LABEL: define available_externally <16 x float> @_Z4sqrtDv16_f(<16 x float> noundef %__x)
+// DEFAULT: call <16 x float> @llvm.sqrt.v16f32(<16 x float> %{{.+}}), !fpmath [[$FPMATH:![0-9]+]]{{$}}
+// CORRECTLYROUNDED: call <16 x float> @llvm.sqrt.v16f32(<16 x float> %{{.+}}){{$}}
+
+// DEFAULT-UNSAFE: call reassoc nsz arcp contract afn <16 x float> @llvm.sqrt.v16f32(<16 x float> %{{.+}}), !fpmath [[$FPMATH:![0-9]+]]{{$}}
+// CORRECTLYROUNDED-UNSAFE: call reassoc nsz arcp contract afn <16 x float> @llvm.sqrt.v16f32(<16 x float> %{{.+}}){{$}}
+
+
+// Not for f64
+// CHECK-LABEL: define {{.*}} double @call_sqrt_f64(
+// CHECK: call {{.*}} double @_Z4sqrtd(double noundef %{{.+}}) #{{[0-9]+$}}
+double call_sqrt_f64(double x) {
+  return sqrt(x);
+}
+
+// CHECK-NOT: define
+
+// Not for f64
+// CHECK-LABEL: define {{.*}} <2 x double> @call_sqrt_v2f64(
+// CHECK: call {{.*}} <2 x double> @_Z4sqrtDv2_d(<2 x double> noundef %{{.+}}) #{{[0-9]+$}}
+double2 call_sqrt_v2f64(double2 x) {
+  return sqrt(x);
+}
+
+// CHECK-NOT: define
+
+// CHECK-LABEL: define {{.*}} <3 x double> @call_sqrt_v3f64(
+// CHECK: call {{.*}} <3 x double> @_Z4sqrtDv3_d(<3 x double> noundef %{{.+}}) #{{[0-9]+$}}
+double3 call_sqrt_v3f64(double3 x) {
+  return sqrt(x);
+}
+
+// CHECK-NOT: define
+
+// CHECK-LABEL: define {{.*}} <4 x double> @call_sqrt_v4f64(
+// CHECK: call {{.*}} <4 x double> @_Z4sqrtDv4_d(<4 x double> noundef %{{.+}}) #{{[0-9]+$}}
+double4 call_sqrt_v4f64(double4 x) {
+  return sqrt(x);
+}
+
+// CHECK-NOT: define
+
+// CHECK-LABEL: define {{.*}} <8 x double> @call_sqrt_v8f64(
+// CHECK: call {{.*}} <8 x double> @_Z4sqrtDv8_d(<8 x double> noundef %{{.+}}) #{{[0-9]+$}}
+double8 call_sqrt_v8f64(double8 x) {
+  return sqrt(x);
+}
+
+// CHECK-NOT: define
+
+// CHECK-LABEL: define {{.*}} <16 x double> @call_sqrt_v16f64(
+// CHECK: call {{.*}} <16 x double> @_Z4sqrtDv16_d(<16 x double> noundef %{{.+}}) #{{[0-9]+$}}
+double16 call_sqrt_v16f64(double16 x) {
+  return sqrt(x);
+}
+
+// CHECK-NOT: define
+
+// Not for f16
+// CHECK-LABEL: define {{.*}} half @call_sqrt_f16(
+// CHECK: call {{.*}} half @_Z4sqrtDh(half noundef %{{.+}}) #{{[0-9]+$}}
+half call_sqrt_f16(half x) {
+  return sqrt(x);
+}
+
+// CHECK-NOT: define
+
+// CHECK-LABEL: define {{.*}} <2 x half> @call_sqrt_v2f16(
+// CHECK: call {{.*}} <2 x half> @_Z4sqrtDv2_Dh(<2 x half> noundef %{{.+}}) #{{[0-9]+$}}
+half2 call_sqrt_v2f16(half2 x) {
+  return sqrt(x);
+}
+
+// CHECK-NOT: define
+
+// CHECK-LABEL: define {{.*}} <3 x half> @call_sqrt_v3f16(
+// CHECK: call {{.*}} <3 x half> @_Z4sqrtDv3_Dh(<3 x half> noundef %{{.+}}) #{{[0-9]+$}}
+half3 call_sqrt_v3f16(half3 x) {
+  return sqrt(x);
+}
+
+// CHECK-NOT: define
+
+// CHECK-LABEL: define {{.*}} <4 x half> @call_sqrt_v4f16(
+// CHECK: call {{.*}} <4 x half> @_Z4sqrtDv4_Dh(<4 x half> noundef %{{.+}}) #{{[0-9]+$}}
+half4 call_sqrt_v4f16(half4 x) {
+  return sqrt(x);
+}
+
+// CHECK-NOT: define
+
+// CHECK-LABEL: define {{.*}} <8 x half> @call_sqrt_v8f16(
+// CHECK: call {{.*}} <8 x half> @_Z4sqrtDv8_Dh(<8 x half> noundef %{{.+}}) #{{[0-9]+$}}
+half8 call_sqrt_v8f16(half8 x) {
+  return sqrt(x);
+}
+
+// CHECK-NOT: define
+
+// CHECK-LABEL: define {{.*}} <16 x half> @call_sqrt_v16f16(
+// CHECK: call {{.*}} <16 x half> @_Z4sqrtDv16_Dh(<16 x half> noundef %{{.+}}) #{{[0-9]+$}}
+half16 call_sqrt_v16f16(half16 x) {
+  return sqrt(x);
+}
+
+// CHECK-NOT: define
+
+// DEFAULT: [[$FPMATH]] = !{float 3.000000e+00}

clang/test/Sema/builtins-elementwise-math.c

void test_builtin_elementwise_sin(int i, float f, double d, float4 v, int3 iv, unsigned u, unsigned4 uv) {
 
   u = __builtin_elementwise_sin(u);
   // expected-error@-1 {{1st argument must be a floating point type (was 'unsigned int')}}
 
   uv = __builtin_elementwise_sin(uv);
   // expected-error@-1 {{1st argument must be a floating point type (was 'unsigned4' (vector of 4 'unsigned int' values))}}
 }
 
+void test_builtin_elementwise_sqrt(int i, float f, double d, float4 v, int3 iv, unsigned u, unsigned4 uv) {
+
+  struct Foo s = __builtin_elementwise_sqrt(f);
+  // expected-error@-1 {{initializing 'struct Foo' with an expression of incompatible type 'float'}}
+
+  i = __builtin_elementwise_sqrt();
+  // expected-error@-1 {{too few arguments to function call, expected 1, have 0}}
+
+  i = __builtin_elementwise_sqrt(i);
+  // expected-error@-1 {{1st argument must be a floating point type (was 'int')}}
+
+  i = __builtin_elementwise_sqrt(f, f);
+  // expected-error@-1 {{too many arguments to function call, expected 1, have 2}}
+
+  u = __builtin_elementwise_sqrt(u);
+  // expected-error@-1 {{1st argument must be a floating point type (was 'unsigned int')}}
+
+  uv = __builtin_elementwise_sqrt(uv);
+  // expected-error@-1 {{1st argument must be a floating point type (was 'unsigned4' (vector of 4 'unsigned int' values))}}
+}
+
 void test_builtin_elementwise_trunc(int i, float f, double d, float4 v, int3 iv, unsigned u, unsigned4 uv) {
 
   struct Foo s = __builtin_elementwise_trunc(f);
   // expected-error@-1 {{initializing 'struct Foo' with an expression of incompatible type 'float'}}
 
   i = __builtin_elementwise_trunc();
   // expected-error@-1 {{too few arguments to function call, expected 1, have 0}}
 

clang/test/SemaCXX/builtins-elementwise-math.cpp

 
 void test_builtin_elementwise_sin() {
   const float a = 42.0;
   float b = 42.3;
   static_assert(!is_const<decltype(__builtin_elementwise_sin(a))>::value);
   static_assert(!is_const<decltype(__builtin_elementwise_sin(b))>::value);
 }
 
+void test_builtin_elementwise_sqrt() {
+  const float a = 42.0;
+  float b = 42.3;
+  static_assert(!is_const<decltype(__builtin_elementwise_sqrt(a))>::value);
+  static_assert(!is_const<decltype(__builtin_elementwise_sqrt(b))>::value);
+}
+
 void test_builtin_elementwise_log() {
   const float a = 42.0;
   float b = 42.3;
   static_assert(!is_const<decltype(__builtin_elementwise_log(a))>::value);
   static_assert(!is_const<decltype(__builtin_elementwise_log(b))>::value);
 }
 
 void test_builtin_elementwise_log10() {