AMDGPU: Add llvm.amdgcn.sqrt intrinsic

[arsenm]

Summary

I spread the GlobalISel test into the regular one, which I've been
avoiding so far.

[foad]

Thanks! I assume the intention is that it should map directly to a v_sqrt_f* instruction, even if llvm.sqrt does something more sophisticated?

llvm/lib/Target/AMDGPU/AMDGPUInstructions.td

819

Typo "shoulud".

[foad]

Can we have f16 tests too, if it works for f16?

[arsenm]

In D82560#2114645, @foad wrote:

Can we have f16 tests too, if it works for f16?

I didn't realize there was an f16 version, but apparently there is

[arsenm]

Handle f16

[foad]

Typo is still there but otherwise LGTM.

[arsenm]

9e03bdebc17a223416d682f64ef2046b8bf0fc98

clang/include/clang/Basic/BuiltinsAMDGPU.def

 BUILTIN(__builtin_amdgcn_div_fmas, "ddddb", "nc")
 BUILTIN(__builtin_amdgcn_div_fmasf, "ffffb", "nc")
 BUILTIN(__builtin_amdgcn_div_fixup, "dddd", "nc")
 BUILTIN(__builtin_amdgcn_div_fixupf, "ffff", "nc")
 BUILTIN(__builtin_amdgcn_trig_preop, "ddi", "nc")
 BUILTIN(__builtin_amdgcn_trig_preopf, "ffi", "nc")
 BUILTIN(__builtin_amdgcn_rcp, "dd", "nc")
 BUILTIN(__builtin_amdgcn_rcpf, "ff", "nc")
+BUILTIN(__builtin_amdgcn_sqrt, "dd", "nc")
+BUILTIN(__builtin_amdgcn_sqrtf, "ff", "nc")
 BUILTIN(__builtin_amdgcn_rsq, "dd", "nc")
 BUILTIN(__builtin_amdgcn_rsqf, "ff", "nc")
 BUILTIN(__builtin_amdgcn_rsq_clamp, "dd", "nc")
 BUILTIN(__builtin_amdgcn_rsq_clampf, "ff", "nc")
 BUILTIN(__builtin_amdgcn_sinf, "ff", "nc")
 BUILTIN(__builtin_amdgcn_cosf, "ff", "nc")
 BUILTIN(__builtin_amdgcn_log_clampf, "ff", "nc")
 BUILTIN(__builtin_amdgcn_ldexp, "ddi", "nc")

 BUILTIN(__builtin_amdgcn_interp_mov, "fUiUiUiUi", "nc")
 
 //===----------------------------------------------------------------------===//
 // VI+ only builtins.
 //===----------------------------------------------------------------------===//
 
 TARGET_BUILTIN(__builtin_amdgcn_div_fixuph, "hhhh", "nc", "16-bit-insts")
 TARGET_BUILTIN(__builtin_amdgcn_rcph, "hh", "nc", "16-bit-insts")
+TARGET_BUILTIN(__builtin_amdgcn_sqrth, "hh", "nc", "16-bit-insts")
 TARGET_BUILTIN(__builtin_amdgcn_rsqh, "hh", "nc", "16-bit-insts")
 TARGET_BUILTIN(__builtin_amdgcn_sinh, "hh", "nc", "16-bit-insts")
 TARGET_BUILTIN(__builtin_amdgcn_cosh, "hh", "nc", "16-bit-insts")
 TARGET_BUILTIN(__builtin_amdgcn_ldexph, "hhi", "nc", "16-bit-insts")
 TARGET_BUILTIN(__builtin_amdgcn_frexp_manth, "hh", "nc", "16-bit-insts")
 TARGET_BUILTIN(__builtin_amdgcn_frexp_exph, "sh", "nc", "16-bit-insts")
 TARGET_BUILTIN(__builtin_amdgcn_fracth, "hh", "nc", "16-bit-insts")
 TARGET_BUILTIN(__builtin_amdgcn_classh, "bhi", "nc", "16-bit-insts")

clang/lib/CodeGen/CGBuiltin.cpp

case AMDGPU::BI__builtin_amdgcn_div_fixuph:
     return emitTernaryBuiltin(*this, E, Intrinsic::amdgcn_div_fixup);
   case AMDGPU::BI__builtin_amdgcn_trig_preop:
   case AMDGPU::BI__builtin_amdgcn_trig_preopf:
     return emitFPIntBuiltin(*this, E, Intrinsic::amdgcn_trig_preop);
   case AMDGPU::BI__builtin_amdgcn_rcp:
   case AMDGPU::BI__builtin_amdgcn_rcpf:
   case AMDGPU::BI__builtin_amdgcn_rcph:
     return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_rcp);
+  case AMDGPU::BI__builtin_amdgcn_sqrt:
+  case AMDGPU::BI__builtin_amdgcn_sqrtf:
+  case AMDGPU::BI__builtin_amdgcn_sqrth:
+    return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_sqrt);
   case AMDGPU::BI__builtin_amdgcn_rsq:
   case AMDGPU::BI__builtin_amdgcn_rsqf:
   case AMDGPU::BI__builtin_amdgcn_rsqh:
     return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_rsq);
   case AMDGPU::BI__builtin_amdgcn_rsq_clamp:
   case AMDGPU::BI__builtin_amdgcn_rsq_clampf:
     return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_rsq_clamp);
   case AMDGPU::BI__builtin_amdgcn_sinf:

clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl

 
 // CHECK-LABEL: @test_rcp_f16
 // CHECK: call half @llvm.amdgcn.rcp.f16
 void test_rcp_f16(global half* out, half a)
 {
   *out = __builtin_amdgcn_rcph(a);
 }
 
+// CHECK-LABEL: @test_sqrt_f16
+// CHECK: call half @llvm.amdgcn.sqrt.f16
+void test_sqrt_f16(global half* out, half a)
+{
+  *out = __builtin_amdgcn_sqrth(a);
+}
+
 // CHECK-LABEL: @test_rsq_f16
 // CHECK: call half @llvm.amdgcn.rsq.f16
 void test_rsq_f16(global half* out, half a)
 {
   *out = __builtin_amdgcn_rsqh(a);
 }
 
 // CHECK-LABEL: @test_sin_f16

clang/test/CodeGenOpenCL/builtins-amdgcn.cl

 
 // CHECK-LABEL: @test_rcp_f64
 // CHECK: call double @llvm.amdgcn.rcp.f64
 void test_rcp_f64(global double* out, double a)
 {
   *out = __builtin_amdgcn_rcp(a);
 }
 
+// CHECK-LABEL: @test_sqrt_f32
+// CHECK: call float @llvm.amdgcn.sqrt.f32
+void test_sqrt_f32(global float* out, float a)
+{
+  *out = __builtin_amdgcn_sqrtf(a);
+}
+
+// CHECK-LABEL: @test_sqrt_f64
+// CHECK: call double @llvm.amdgcn.sqrt.f64
+void test_sqrt_f64(global double* out, double a)
+{
+  *out = __builtin_amdgcn_sqrt(a);
+}
+
 // CHECK-LABEL: @test_rsq_f32
 // CHECK: call float @llvm.amdgcn.rsq.f32
 void test_rsq_f32(global float* out, float a)
 {
   *out = __builtin_amdgcn_rsqf(a);
 }
 
 // CHECK-LABEL: @test_rsq_f64

clang/test/SemaOpenCL/builtins-amdgcn-error-f16.cl

 // REQUIRES: amdgpu-registered-target
 // RUN: %clang_cc1 -triple amdgcn-- -verify -S -o - %s
 
 #pragma OPENCL EXTENSION cl_khr_fp16 : enable
 
 __attribute__((target("arch=tahiti")))
 void test_f16_tahiti(global half *out, half a, half b, half c)
 {
   *out = __builtin_amdgcn_div_fixuph(a, b, c); // expected-error {{'__builtin_amdgcn_div_fixuph' needs target feature 16-bit-insts}}
   *out = __builtin_amdgcn_rcph(a); // expected-error {{'__builtin_amdgcn_rcph' needs target feature 16-bit-insts}}
+  *out = __builtin_amdgcn_sqrth(a); // expected-error {{'__builtin_amdgcn_sqrth' needs target feature 16-bit-insts}}
   *out = __builtin_amdgcn_rsqh(a); // expected-error {{'__builtin_amdgcn_rsqh' needs target feature 16-bit-insts}}
   *out = __builtin_amdgcn_sinh(a); // expected-error {{'__builtin_amdgcn_sinh' needs target feature 16-bit-insts}}
   *out = __builtin_amdgcn_cosh(a); // expected-error {{'__builtin_amdgcn_cosh' needs target feature 16-bit-insts}}
   *out = __builtin_amdgcn_ldexph(a, b); // expected-error {{'__builtin_amdgcn_ldexph' needs target feature 16-bit-insts}}
   *out = __builtin_amdgcn_frexp_manth(a); // expected-error {{'__builtin_amdgcn_frexp_manth' needs target feature 16-bit-insts}}
   *out = __builtin_amdgcn_frexp_exph(a); // expected-error {{'__builtin_amdgcn_frexp_exph' needs target feature 16-bit-insts}}
   *out = __builtin_amdgcn_fracth(a); // expected-error {{'__builtin_amdgcn_fracth' needs target feature 16-bit-insts}}
   *out = __builtin_amdgcn_classh(a, b); // expected-error {{'__builtin_amdgcn_classh' needs target feature 16-bit-insts}}
   *out = __builtin_amdgcn_fmed3h(a, b, c); // expected-error {{'__builtin_amdgcn_fmed3h' needs target feature gfx9-insts}}
 }

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

def int_amdgcn_rcp : Intrinsic<
   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable, IntrWillReturn]
 >;
 
 def int_amdgcn_rcp_legacy : GCCBuiltin<"__builtin_amdgcn_rcp_legacy">,
   Intrinsic<[llvm_float_ty], [llvm_float_ty],
   [IntrNoMem, IntrSpeculatable, IntrWillReturn]
 >;
 
+def int_amdgcn_sqrt :  Intrinsic<
+  [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable, IntrWillReturn]
+>;
+
 def int_amdgcn_rsq :  Intrinsic<
   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable, IntrWillReturn]
 >;
 
 def int_amdgcn_rsq_legacy :  GCCBuiltin<"__builtin_amdgcn_rsq_legacy">,
   Intrinsic<
   [llvm_float_ty], [llvm_float_ty], [IntrNoMem, IntrSpeculatable, IntrWillReturn]
 >;

llvm/lib/Target/AMDGPU/AMDGPUInstructions.td

def fmaxnum_like_oneuse : PatFrags<(ops node:$src0, node:$src1),
   [(fmaxnum_ieee_oneuse node:$src0, node:$src1),
    (fmaxnum_oneuse node:$src0, node:$src1)]
 >;
 
 def any_fmad : PatFrags<(ops node:$src0, node:$src1, node:$src2),
   [(fmad node:$src0, node:$src1, node:$src2),
    (AMDGPUfmad_ftz node:$src0, node:$src1, node:$src2)]
 >;
+
+// FIXME: fsqrt shoulud not select directly
+def any_amdgcn_sqrt : PatFrags<(ops node:$src0),
+  [(fsqrt node:$src0), (int_amdgcn_sqrt node:$src0)]
+>;

llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp

case AMDGPU::G_INTRINSIC: {
     case Intrinsic::amdgcn_div_fmas:
     case Intrinsic::amdgcn_div_fixup:
     case Intrinsic::amdgcn_trig_preop:
     case Intrinsic::amdgcn_sin:
     case Intrinsic::amdgcn_cos:
     case Intrinsic::amdgcn_log_clamp:
     case Intrinsic::amdgcn_rcp:
     case Intrinsic::amdgcn_rcp_legacy:
+    case Intrinsic::amdgcn_sqrt:
     case Intrinsic::amdgcn_rsq:
     case Intrinsic::amdgcn_rsq_legacy:
     case Intrinsic::amdgcn_rsq_clamp:
     case Intrinsic::amdgcn_fmul_legacy:
     case Intrinsic::amdgcn_ldexp:
     case Intrinsic::amdgcn_frexp_mant:
     case Intrinsic::amdgcn_frexp_exp:
     case Intrinsic::amdgcn_fract:

llvm/lib/Target/AMDGPU/VOP1Instructions.td

 defm V_FLOOR_F32 : VOP1Inst <"v_floor_f32", VOP_F32_F32, ffloor>;
 
 let SchedRW = [WriteTrans32] in {
 defm V_EXP_F32 : VOP1Inst <"v_exp_f32", VOP_F32_F32, fexp2>;
 defm V_LOG_F32 : VOP1Inst <"v_log_f32", VOP_F32_F32, flog2>;
 defm V_RCP_F32 : VOP1Inst <"v_rcp_f32", VOP_F32_F32, AMDGPUrcp>;
 defm V_RCP_IFLAG_F32 : VOP1Inst <"v_rcp_iflag_f32", VOP_F32_F32, AMDGPUrcp_iflag>;
 defm V_RSQ_F32 : VOP1Inst <"v_rsq_f32", VOP_F32_F32, AMDGPUrsq>;
-defm V_SQRT_F32 : VOP1Inst <"v_sqrt_f32", VOP_F32_F32, fsqrt>;
+defm V_SQRT_F32 : VOP1Inst <"v_sqrt_f32", VOP_F32_F32, any_amdgcn_sqrt>;
 } // End SchedRW = [WriteTrans32]
 
 let SchedRW = [WriteTrans64] in {
 defm V_RCP_F64 : VOP1Inst <"v_rcp_f64", VOP_F64_F64, AMDGPUrcp>;
 defm V_RSQ_F64 : VOP1Inst <"v_rsq_f64", VOP_F64_F64, AMDGPUrsq>;
-defm V_SQRT_F64 : VOP1Inst <"v_sqrt_f64", VOP_F64_F64, fsqrt>;
+defm V_SQRT_F64 : VOP1Inst <"v_sqrt_f64", VOP_F64_F64, any_amdgcn_sqrt>;
 } // End SchedRW = [WriteTrans64]
 
 let SchedRW = [WriteTrans32] in {
 defm V_SIN_F32 : VOP1Inst <"v_sin_f32", VOP_F32_F32, AMDGPUsin>;
 defm V_COS_F32 : VOP1Inst <"v_cos_f32", VOP_F32_F32, AMDGPUcos>;
 } // End SchedRW = [WriteTrans32]
 
 defm V_NOT_B32 : VOP1Inst <"v_not_b32", VOP_I32_I32>;

 let FPDPRounding = 1 in {
 defm V_CVT_F16_U16 : VOP1Inst <"v_cvt_f16_u16", VOP1_F16_I16, uint_to_fp>;
 defm V_CVT_F16_I16 : VOP1Inst <"v_cvt_f16_i16", VOP1_F16_I16, sint_to_fp>;
 } // End FPDPRounding = 1
 defm V_CVT_U16_F16 : VOP1Inst <"v_cvt_u16_f16", VOP_I16_F16, fp_to_uint>;
 defm V_CVT_I16_F16 : VOP1Inst <"v_cvt_i16_f16", VOP_I16_F16, fp_to_sint>;
 let SchedRW = [WriteTrans32] in {
 defm V_RCP_F16 : VOP1Inst <"v_rcp_f16", VOP_F16_F16, AMDGPUrcp>;
-defm V_SQRT_F16 : VOP1Inst <"v_sqrt_f16", VOP_F16_F16, fsqrt>;
+defm V_SQRT_F16 : VOP1Inst <"v_sqrt_f16", VOP_F16_F16, any_amdgcn_sqrt>;
 defm V_RSQ_F16 : VOP1Inst <"v_rsq_f16", VOP_F16_F16, AMDGPUrsq>;
 defm V_LOG_F16 : VOP1Inst <"v_log_f16", VOP_F16_F16, flog2>;
 defm V_EXP_F16 : VOP1Inst <"v_exp_f16", VOP_F16_F16, fexp2>;
 defm V_SIN_F16 : VOP1Inst <"v_sin_f16", VOP_F16_F16, AMDGPUsin>;
 defm V_COS_F16 : VOP1Inst <"v_cos_f16", VOP_F16_F16, AMDGPUcos>;
 } // End SchedRW = [WriteTrans32]
 defm V_FREXP_MANT_F16 : VOP1Inst <"v_frexp_mant_f16", VOP_F16_F16, int_amdgcn_frexp_mant>;
 defm V_FREXP_EXP_I16_F16 : VOP1Inst <"v_frexp_exp_i16_f16", VOP_I16_F16, int_amdgcn_frexp_exp>;

llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sqrt.f16.ll

+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=fiji < %s | FileCheck -check-prefix=GCN %s
+; RUN: llc -global-isel -mtriple=amdgcn-amd-amdhsa -mcpu=fiji < %s | FileCheck -check-prefix=GCN %s
+
+define half @v_sqrt_f16(half %src)  {
+; GCN-LABEL: v_sqrt_f16:
+; GCN:       ; %bb.0:
+; GCN-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GCN-NEXT:    v_sqrt_f16_e32 v0, v0
+; GCN-NEXT:    s_setpc_b64 s[30:31]
+  %sqrt = call half @llvm.amdgcn.sqrt.f16(half %src)
+  ret half %sqrt
+}
+
+define half @v_fabs_sqrt_f16(half %src)  {
+; GCN-LABEL: v_fabs_sqrt_f16:
+; GCN:       ; %bb.0:
+; GCN-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GCN-NEXT:    v_sqrt_f16_e64 v0, |v0|
+; GCN-NEXT:    s_setpc_b64 s[30:31]
+  %fabs.src = call half @llvm.fabs.f16(half %src)
+  %sqrt = call half @llvm.amdgcn.sqrt.f16(half %fabs.src)
+  ret half %sqrt
+}
+
+define half @v_fneg_fabs_sqrt_f16(half %src)  {
+; GCN-LABEL: v_fneg_fabs_sqrt_f16:
+; GCN:       ; %bb.0:
+; GCN-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GCN-NEXT:    v_sqrt_f16_e64 v0, -|v0|
+; GCN-NEXT:    s_setpc_b64 s[30:31]
+  %fabs.src = call half @llvm.fabs.f16(half %src)
+  %neg.fabs.src = fneg half %fabs.src
+  %sqrt = call half @llvm.amdgcn.sqrt.f16(half %neg.fabs.src)
+  ret half %sqrt
+}
+
+declare half @llvm.amdgcn.sqrt.f16(half) #0
+declare half @llvm.fabs.f16(half) #0
+
+attributes #0 = { nounwind readnone speculatable willreturn }

llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sqrt.ll

+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=hawaii < %s | FileCheck -check-prefix=GCN %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=fiji < %s | FileCheck -check-prefix=GCN %s
+; RUN: llc -global-isel -mtriple=amdgcn-amd-amdhsa -mcpu=hawaii < %s | FileCheck -check-prefix=GCN %s
+; RUN: llc -global-isel -mtriple=amdgcn-amd-amdhsa -mcpu=fiji < %s | FileCheck -check-prefix=GCN %s
+
+define float @v_sqrt_f32(float %src)  {
+; GCN-LABEL: v_sqrt_f32:
+; GCN:       ; %bb.0:
+; GCN-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GCN-NEXT:    v_sqrt_f32_e32 v0, v0
+; GCN-NEXT:    s_setpc_b64 s[30:31]
+  %sqrt = call float @llvm.amdgcn.sqrt.f32(float %src)
+  ret float %sqrt
+}
+
+define float @v_fabs_sqrt_f32(float %src)  {
+; GCN-LABEL: v_fabs_sqrt_f32:
+; GCN:       ; %bb.0:
+; GCN-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GCN-NEXT:    v_sqrt_f32_e64 v0, |v0|
+; GCN-NEXT:    s_setpc_b64 s[30:31]
+  %fabs.src = call float @llvm.fabs.f32(float %src)
+  %sqrt = call float @llvm.amdgcn.sqrt.f32(float %fabs.src)
+  ret float %sqrt
+}
+
+define float @v_fneg_fabs_sqrt_f32(float %src)  {
+; GCN-LABEL: v_fneg_fabs_sqrt_f32:
+; GCN:       ; %bb.0:
+; GCN-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GCN-NEXT:    v_sqrt_f32_e64 v0, -|v0|
+; GCN-NEXT:    s_setpc_b64 s[30:31]
+  %fabs.src = call float @llvm.fabs.f32(float %src)
+  %neg.fabs.src = fneg float %fabs.src
+  %sqrt = call float @llvm.amdgcn.sqrt.f32(float %neg.fabs.src)
+  ret float %sqrt
+}
+
+define double @v_sqrt_f64(double %src)  {
+; GCN-LABEL: v_sqrt_f64:
+; GCN:       ; %bb.0:
+; GCN-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GCN-NEXT:    v_sqrt_f64_e32 v[0:1], v[0:1]
+; GCN-NEXT:    s_setpc_b64 s[30:31]
+  %sqrt = call double @llvm.amdgcn.sqrt.f64(double %src)
+  ret double %sqrt
+}
+
+define double @v_fabs_sqrt_f64(double %src)  {
+; GCN-LABEL: v_fabs_sqrt_f64:
+; GCN:       ; %bb.0:
+; GCN-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GCN-NEXT:    v_sqrt_f64_e64 v[0:1], |v[0:1]|
+; GCN-NEXT:    s_setpc_b64 s[30:31]
+  %fabs.src = call double @llvm.fabs.f64(double %src)
+  %sqrt = call double @llvm.amdgcn.sqrt.f64(double %fabs.src)
+  ret double %sqrt
+}
+
+define double @v_fneg_fabs_sqrt_f64(double %src)  {
+; GCN-LABEL: v_fneg_fabs_sqrt_f64:
+; GCN:       ; %bb.0:
+; GCN-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GCN-NEXT:    v_sqrt_f64_e64 v[0:1], -|v[0:1]|
+; GCN-NEXT:    s_setpc_b64 s[30:31]
+  %fabs.src = call double @llvm.fabs.f64(double %src)
+  %neg.fabs.src = fneg double %fabs.src
+  %sqrt = call double @llvm.amdgcn.sqrt.f64(double %neg.fabs.src)
+  ret double %sqrt
+}
+
+declare float @llvm.amdgcn.sqrt.f32(float) #0
+declare double @llvm.amdgcn.sqrt.f64(double) #0
+declare float @llvm.fabs.f32(float) #0
+declare double @llvm.fabs.f64(double) #0
+
+attributes #0 = { nounwind readnone speculatable willreturn }