[clang] 9e03bde - AMDGPU: Add llvm.amdgcn.sqrt intrinsic
Matt Arsenault via cfe-commits
cfe-commits at lists.llvm.org
Fri Jun 26 12:07:16 PDT 2020
Author: Matt Arsenault
Date: 2020-06-26T15:07:07-04:00
New Revision: 9e03bdebc17a223416d682f64ef2046b8bf0fc98
URL: https://github.com/llvm/llvm-project/commit/9e03bdebc17a223416d682f64ef2046b8bf0fc98
DIFF: https://github.com/llvm/llvm-project/commit/9e03bdebc17a223416d682f64ef2046b8bf0fc98.diff
LOG: AMDGPU: Add llvm.amdgcn.sqrt intrinsic
I spread the GlobalISel test into the regular one, which I've been
avoiding so far.
Added:
llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sqrt.f16.ll
llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sqrt.ll
Modified:
clang/include/clang/Basic/BuiltinsAMDGPU.def
clang/lib/CodeGen/CGBuiltin.cpp
clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
clang/test/CodeGenOpenCL/builtins-amdgcn.cl
clang/test/SemaOpenCL/builtins-amdgcn-error-f16.cl
llvm/include/llvm/IR/IntrinsicsAMDGPU.td
llvm/lib/Target/AMDGPU/AMDGPUInstructions.td
llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
llvm/lib/Target/AMDGPU/VOP1Instructions.td
Removed:
################################################################################
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def
index 9add10c64962..60be0525fabc 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def
@@ -77,6 +77,8 @@ 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")
@@ -162,6 +164,7 @@ BUILTIN(__builtin_amdgcn_interp_mov, "fUiUiUiUi", "nc")
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")
diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp
index 2eef4f284271..b5c4841578c4 100644
--- a/clang/lib/CodeGen/CGBuiltin.cpp
+++ b/clang/lib/CodeGen/CGBuiltin.cpp
@@ -14702,6 +14702,10 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
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:
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
index e3e6b81271d1..5884f84ab081 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
@@ -22,6 +22,13 @@ 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)
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
index 3563ad464c66..56c83df6b6b4 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -116,6 +116,20 @@ 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)
diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-error-f16.cl b/clang/test/SemaOpenCL/builtins-amdgcn-error-f16.cl
index 3487b1a5a803..fdb2f3f3c981 100644
--- a/clang/test/SemaOpenCL/builtins-amdgcn-error-f16.cl
+++ b/clang/test/SemaOpenCL/builtins-amdgcn-error-f16.cl
@@ -8,6 +8,7 @@ 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}}
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index 734128d68218..01380afae006 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -267,6 +267,10 @@ def int_amdgcn_rcp_legacy : GCCBuiltin<"__builtin_amdgcn_rcp_legacy">,
[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]
>;
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUInstructions.td b/llvm/lib/Target/AMDGPU/AMDGPUInstructions.td
index 706053a4e315..16b061d9251f 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUInstructions.td
+++ b/llvm/lib/Target/AMDGPU/AMDGPUInstructions.td
@@ -815,3 +815,8 @@ 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 should not select directly
+def any_amdgcn_sqrt : PatFrags<(ops node:$src0),
+ [(fsqrt node:$src0), (int_amdgcn_sqrt node:$src0)]
+>;
diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp b/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
index b8f338df9213..b84f4c2cb63e 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
@@ -3952,6 +3952,7 @@ AMDGPURegisterBankInfo::getInstrMapping(const MachineInstr &MI) const {
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:
diff --git a/llvm/lib/Target/AMDGPU/VOP1Instructions.td b/llvm/lib/Target/AMDGPU/VOP1Instructions.td
index c83442acf5e1..17f334f62a30 100644
--- a/llvm/lib/Target/AMDGPU/VOP1Instructions.td
+++ b/llvm/lib/Target/AMDGPU/VOP1Instructions.td
@@ -248,13 +248,13 @@ 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 {
@@ -388,7 +388,7 @@ 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>;
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sqrt.f16.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sqrt.f16.ll
new file mode 100644
index 000000000000..bbfb88a4b22a
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sqrt.f16.ll
@@ -0,0 +1,41 @@
+; 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 }
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sqrt.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sqrt.ll
new file mode 100644
index 000000000000..0257a3d11142
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sqrt.ll
@@ -0,0 +1,78 @@
+; 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 }
More information about the cfe-commits
mailing list