clang: Add __builtin_elementwise_sqrt

[arsenm]

Summary

This will be used in the opencl builtin headers to provide direct
intrinsic access with proper !fpmath metadata.

[bob80905]

I think you need to mention this new builtin in clang/docs/ReleaseNotes.rst.

clang/lib/CodeGen/CGBuiltin.cpp

2534

Nit: I think despite this code working, it should be moved to be grouped with the other elementwise builtins at around line 3240.
Consider BI__builtin_log at around 2436 and the other log builtins, the bultin_elementwise_log builtin is not in that group since it's with the other elementwise builtins. Writing this case here would make an inconsistency with the placement of the other elementwise builtins.

[arsenm]

clang/lib/CodeGen/CGBuiltin.cpp

2534

I think it's more important to group with the emission of the metadata

Also don't understand why there is a broken constrained path and a redundant set of constrained handling in IRBuilder

[arsenm]

ping

[bob80905]

Shouldn't this new builtin be mentioned in Release Notes?

[arsenm]

Release note

[arsenm]

9e3d9c9eae03910d93e2312e1e0845433c779998

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/docs/ReleaseNotes.rst

 - Add ``__builtin_elementwise_log`` builtin for floating point types only.
 - Add ``__builtin_elementwise_log10`` builtin for floating point types only.
 - Add ``__builtin_elementwise_log2`` builtin for floating point types only.
 - Add ``__builtin_elementwise_exp`` builtin for floating point types only.
 - Add ``__builtin_elementwise_exp2`` builtin for floating point types only.
 - Add ``__builtin_set_flt_rounds`` builtin for X86, x86_64, Arm and AArch64 only.
 - Add ``__builtin_elementwise_pow`` builtin for floating point types only.
 - Add ``__builtin_elementwise_bitreverse`` builtin for integer types only.
+- Add ``__builtin_elementwise_sqrt`` builtin for floating point types only.
 
 AST Matchers
 ------------
 - Add ``convertVectorExpr``.
 - Add ``dependentSizedExtVectorType``.
 
 clang-format
 ------------

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/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/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() {