RPCS3/llvm-mirror
1
0
Fork 0
mirror of https://github.com/RPCS3/llvm-mirror.git synced 2024-11-23 19:23:23 +01:00
Code Issues Projects Releases Wiki Activity

AMDGPU: Add fp legacy instruction intrinsics

Browse Source 
This could use some additional optimization work
to use mad/mac legacy.

llvm-svn: 276764
This commit is contained in:
Matt Arsenault 2016-07-26 16:45:45 +00:00
parent 29dba078d8
commit d60a4f902f
8 changed files with 125 additions and 2 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

8
include/llvm/IR/IntrinsicsAMDGPU.td
View File

@ -145,10 +145,18 @@ def int_amdgcn_log_clamp : Intrinsic<
[llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]
>;
def int_amdgcn_fmul_legacy : GCCBuiltin<"__builtin_amdgcn_fmul_legacy">,
Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty], [IntrNoMem]
>;
def int_amdgcn_rcp : Intrinsic<
[llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]
>;
def int_amdgcn_rcp_legacy : GCCBuiltin<"__builtin_amdgcn_rcp_legacy">,
Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]
>;
def int_amdgcn_rsq : Intrinsic<
[llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]
>;

2
lib/Target/AMDGPU/AMDGPUISelLowering.cpp
View File

@ -2699,7 +2699,9 @@ const char* AMDGPUTargetLowering::getTargetNodeName(unsigned Opcode) const {
NODE_NAME_CASE(TRIG_PREOP)
NODE_NAME_CASE(RCP)
NODE_NAME_CASE(RSQ)
NODE_NAME_CASE(RCP_LEGACY)
NODE_NAME_CASE(RSQ_LEGACY)
NODE_NAME_CASE(FMUL_LEGACY)
NODE_NAME_CASE(RSQ_CLAMP)
NODE_NAME_CASE(LDEXP)
NODE_NAME_CASE(FP_CLASS)

2
lib/Target/AMDGPU/AMDGPUISelLowering.h
View File

@ -249,7 +249,9 @@ enum NodeType : unsigned {
// For f64, max error 2^29 ULP, handles denormals.
RCP,
RSQ,
RCP_LEGACY,
RSQ_LEGACY,
FMUL_LEGACY,
RSQ_CLAMP,
LDEXP,
FP_CLASS,

5
lib/Target/AMDGPU/AMDGPUInstrInfo.td
View File

@ -67,6 +67,7 @@ def AMDGPUrcp : SDNode<"AMDGPUISD::RCP", SDTFPUnaryOp>;
def AMDGPUrsq : SDNode<"AMDGPUISD::RSQ", SDTFPUnaryOp>;
// out = 1.0 / sqrt(a)
def AMDGPUrcp_legacy : SDNode<"AMDGPUISD::RCP_LEGACY", SDTFPUnaryOp>;
def AMDGPUrsq_legacy : SDNode<"AMDGPUISD::RSQ_LEGACY", SDTFPUnaryOp>;
// out = 1.0 / sqrt(a) result clamped to +/- max_float.
@ -84,6 +85,10 @@ def AMDGPUfmax_legacy : SDNode<"AMDGPUISD::FMAX_LEGACY", SDTFPBinOp,
[]
>;
def AMDGPUfmul_legacy : SDNode<"AMDGPUISD::FMUL_LEGACY", SDTFPBinOp,
[SDNPCommutative, SDNPAssociative]
>;
def AMDGPUclamp : SDNode<"AMDGPUISD::CLAMP", SDTFPTernaryOp, []>;
// out = max(a, b) a and b are signed ints

9
lib/Target/AMDGPU/SIISelLowering.cpp
View File

@ -1996,6 +1996,11 @@ SDValue SITargetLowering::LowerINTRINSIC_WO_CHAIN(SDValue Op,
return DAG.getNode(AMDGPUISD::RSQ_LEGACY, DL, VT, Op.getOperand(1));
}
case Intrinsic::amdgcn_rcp_legacy: {
if (Subtarget->getGeneration() >= SISubtarget::VOLCANIC_ISLANDS)
return emitRemovedIntrinsicError(DAG, DL, VT);
return DAG.getNode(AMDGPUISD::RCP_LEGACY, DL, VT, Op.getOperand(1));
}
case Intrinsic::amdgcn_rsq_clamp: {
if (Subtarget->getGeneration() < SISubtarget::VOLCANIC_ISLANDS)
return DAG.getNode(AMDGPUISD::RSQ_CLAMP, DL, VT, Op.getOperand(1));
@ -2208,6 +2213,9 @@ SDValue SITargetLowering::LowerINTRINSIC_WO_CHAIN(SDValue Op,
return DAG.getNode(AMDGPUISD::DIV_SCALE, DL, Op->getVTList(), Src0,
Denominator, Numerator);
}
case Intrinsic::amdgcn_fmul_legacy:
return DAG.getNode(AMDGPUISD::FMUL_LEGACY, DL, VT,
Op.getOperand(1), Op.getOperand(2));
case Intrinsic::amdgcn_sffbh:
case AMDGPUIntrinsic::AMDGPU_flbit_i32: // Legacy name.
return DAG.getNode(AMDGPUISD::FFBH_I32, DL, VT, Op.getOperand(1));
@ -3356,6 +3364,7 @@ SDValue SITargetLowering::PerformDAGCombine(SDNode *N,
case AMDGPUISD::FRACT:
case AMDGPUISD::RCP:
case AMDGPUISD::RSQ:
case AMDGPUISD::RCP_LEGACY:
case AMDGPUISD::RSQ_LEGACY:
case AMDGPUISD::RSQ_CLAMP:
case AMDGPUISD::LDEXP: {

5
lib/Target/AMDGPU/SIInstructions.td
View File

@ -1402,7 +1402,8 @@ defm V_MOV_FED_B32 : VOP1InstSI <vop1<0x9>, "v_mov_fed_b32", VOP_I32_I32>;
defm V_LOG_CLAMP_F32 : VOP1InstSI <vop1<0x26>, "v_log_clamp_f32",
VOP_F32_F32, int_amdgcn_log_clamp>;
defm V_RCP_CLAMP_F32 : VOP1InstSI <vop1<0x28>, "v_rcp_clamp_f32", VOP_F32_F32>;
defm V_RCP_LEGACY_F32 : VOP1InstSI <vop1<0x29>, "v_rcp_legacy_f32", VOP_F32_F32>;
defm V_RCP_LEGACY_F32 : VOP1InstSI <vop1<0x29>, "v_rcp_legacy_f32",
VOP_F32_F32, AMDGPUrcp_legacy>;
defm V_RSQ_CLAMP_F32 : VOP1InstSI <vop1<0x2c>, "v_rsq_clamp_f32",
VOP_F32_F32, AMDGPUrsq_clamp
>;
@ -1496,7 +1497,7 @@ defm V_SUBREV_F32 : VOP2Inst <vop2<0x5, 0x3>, "v_subrev_f32",
let isCommutable = 1 in {
defm V_MUL_LEGACY_F32 : VOP2Inst <vop2<0x7, 0x4>, "v_mul_legacy_f32",
VOP_F32_F32_F32
VOP_F32_F32_F32, AMDGPUfmul_legacy
>;
defm V_MUL_F32 : VOP2Inst <vop2<0x8, 0x5>, "v_mul_f32",

54
test/CodeGen/AMDGPU/llvm.amdgcn.fmul.legacy.ll Normal file
View File

@ -0,0 +1,54 @@
; RUN: llc -march=amdgcn -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
; RUN: llc -march=amdgcn -mcpu=tonga -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
; GCN-LABEL: {{^}}test_mul_legacy_f32:
; GCN: v_mul_legacy_f32_e32 v{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}
define void @test_mul_legacy_f32(float addrspace(1)* %out, float %a, float %b) #0 {
%result = call float @llvm.amdgcn.fmul.legacy(float %a, float %b)
store float %result, float addrspace(1)* %out, align 4
ret void
}
; GCN-LABEL: {{^}}test_mul_legacy_undef0_f32:
; GCN: v_mul_legacy_f32_e32
define void @test_mul_legacy_undef0_f32(float addrspace(1)* %out, float %a) #0 {
%result = call float @llvm.amdgcn.fmul.legacy(float undef, float %a)
store float %result, float addrspace(1)* %out, align 4
ret void
}
; GCN-LABEL: {{^}}test_mul_legacy_undef1_f32:
; GCN: v_mul_legacy_f32_e32
define void @test_mul_legacy_undef1_f32(float addrspace(1)* %out, float %a) #0 {
%result = call float @llvm.amdgcn.fmul.legacy(float %a, float undef)
store float %result, float addrspace(1)* %out, align 4
ret void
}
; GCN-LABEL: {{^}}test_mul_legacy_fabs_f32:
; GCN: v_mul_legacy_f32_e64 v{{[0-9]+}}, |s{{[0-9]+}}|, |v{{[0-9]+}}|
define void @test_mul_legacy_fabs_f32(float addrspace(1)* %out, float %a, float %b) #0 {
%a.fabs = call float @llvm.fabs.f32(float %a)
%b.fabs = call float @llvm.fabs.f32(float %b)
%result = call float @llvm.amdgcn.fmul.legacy(float %a.fabs, float %b.fabs)
store float %result, float addrspace(1)* %out, align 4
ret void
}
; TODO: Should match mac_legacy/mad_legacy
; GCN-LABEL: {{^}}test_mad_legacy_f32:
; GCN: v_mul_legacy_f32_e32 v{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}
; GCN: v_add_f32_e32
define void @test_mad_legacy_f32(float addrspace(1)* %out, float %a, float %b, float %c) #0 {
%mul = call float @llvm.amdgcn.fmul.legacy(float %a, float %b)
%add = fadd float %mul, %c
store float %add, float addrspace(1)* %out, align 4
ret void
}
declare float @llvm.fabs.f32(float) #1
declare float @llvm.amdgcn.fmul.legacy(float, float) #1
attributes #0 = { nounwind }
attributes #1 = { nounwind readnone }

42
test/CodeGen/AMDGPU/llvm.amdgcn.rcp.legacy.ll Normal file
View File

@ -0,0 +1,42 @@
; RUN: llc -march=amdgcn -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
; RUN: not llc -march=amdgcn -mcpu=fiji -verify-machineinstrs < %s 2>&1 | FileCheck -check-prefix=ERROR %s
; ERROR: error: <unknown>:0:0: in function rcp_legacy_f32 void (float addrspace(1)*, float): intrinsic not supported on subtarget
declare float @llvm.amdgcn.rcp.legacy(float) #0
; GCN-LABEL: {{^}}rcp_legacy_f32:
; GCN: v_rcp_legacy_f32_e32 {{v[0-9]+}}, {{s[0-9]+}}
define void @rcp_legacy_f32(float addrspace(1)* %out, float %src) #1 {
%rcp = call float @llvm.amdgcn.rcp.legacy(float %src) #0
store float %rcp, float addrspace(1)* %out, align 4
ret void
}
; TODO: Really these should be constant folded
; GCN-LABEL: {{^}}rcp_legacy_f32_constant_4.0
; GCN: v_rcp_legacy_f32_e32 {{v[0-9]+}}, 4.0
define void @rcp_legacy_f32_constant_4.0(float addrspace(1)* %out) #1 {
%rcp = call float @llvm.amdgcn.rcp.legacy(float 4.0) #0
store float %rcp, float addrspace(1)* %out, align 4
ret void
}
; GCN-LABEL: {{^}}rcp_legacy_f32_constant_100.0
; GCN: v_rcp_legacy_f32_e32 {{v[0-9]+}}, 0x42c80000
define void @rcp_legacy_f32_constant_100.0(float addrspace(1)* %out) #1 {
%rcp = call float @llvm.amdgcn.rcp.legacy(float 100.0) #0
store float %rcp, float addrspace(1)* %out, align 4
ret void
}
; GCN-LABEL: {{^}}rcp_legacy_undef_f32:
; GCN-NOT: v_rcp_legacy_f32
define void @rcp_legacy_undef_f32(float addrspace(1)* %out) #1 {
%rcp = call float @llvm.amdgcn.rcp.legacy(float undef)
store float %rcp, float addrspace(1)* %out, align 4
ret void
}
attributes #0 = { nounwind readnone }
attributes #1 = { nounwind }