This is an archive of the discontinued LLVM Phabricator instance.

Differential D112394

[AMDGPU] Implement llvm.amdgcn.mulhi.[i,u]24 intrinsics.
ClosedPublic

Authored by abinavpp on Oct 24 2021, 8:18 PM.

Details

Reviewers
arsenm
foad
rampitec
Commits
rG9bd5cfeb1fa4: [AMDGPU] Implement llvm.amdgcn.mulhi.[i,u]24 intrinsics.
Summary

These intrinsics maps to the 24-bit v_mul_hi instructions.

Diff Detail

Event Timeline

abinavpp created this revision.Oct 24 2021, 8:18 PM
Herald added subscribers: kerbowa, hiraditya, t-tye and 6 others. · View Herald TranscriptOct 24 2021, 8:18 PM
abinavpp requested review of this revision.Oct 24 2021, 8:18 PM
Herald added a project: Restricted Project. · View Herald TranscriptOct 24 2021, 8:18 PM
Herald added subscribers: llvm-commits, wdng. · View Herald Transcript
abinavpp mentioned this in D111523: [AMDGPU] Fix 24 bit mul intrinsic generation for > 32 bit result..Oct 24 2021, 8:26 PM
abinavpp added a child revision: D112395: [AMDGPU] Enable 48-bit mul in AMDGPUCodeGenPrepare..Oct 24 2021, 10:20 PM
Harbormaster completed remote builds in B130342: Diff 381824.Oct 24 2021, 10:48 PM
foad added inline comments.Oct 25 2021, 7:40 AM
llvm/lib/Target/AMDGPU/AMDGPUInstrInfo.td
283

Not your fault, but are these mulhi nodes really associative??

llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mulhi.i24.ll
2

Can you add a -global-isel run line? Does it Just Work?

abinavpp updated this revision to Diff 382271.Oct 26 2021, 4:55 AM

Addressed review comments.

abinavpp marked an inline comment as done.Oct 26 2021, 4:57 AM
abinavpp added inline comments.
llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mulhi.i24.ll
2

It's not working at the moment. I've added the mir version for now.

foad added inline comments.Oct 26 2021, 5:24 AM
llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mulhi.i24.ll
2

It should be trivial to implement, just handle the mulhi intrinsics the same as the mul intrinsics in AMDGPURegisterBankInfo.

Harbormaster completed remote builds in B130665: Diff 382271.Oct 26 2021, 5:48 AM
abinavpp updated this revision to Diff 382283.Oct 26 2021, 6:16 AM

Rebased; Addressed review comment.

abinavpp marked an inline comment as done.Oct 26 2021, 6:17 AM
foad accepted this revision.Oct 26 2021, 6:20 AM

LGTM, thanks.

This revision is now accepted and ready to land.Oct 26 2021, 6:20 AM
This revision was landed with ongoing or failed builds.Oct 26 2021, 6:23 AM
Closed by commit rG9bd5cfeb1fa4: [AMDGPU] Implement llvm.amdgcn.mulhi.[i,u]24 intrinsics. (authored by abinavpp). · Explain Why
This revision was automatically updated to reflect the committed changes.
abinavpp added a commit: rG9bd5cfeb1fa4: [AMDGPU] Implement llvm.amdgcn.mulhi.[i,u]24 intrinsics..
Harbormaster completed remote builds in B130677: Diff 382283.Oct 26 2021, 7:01 AM

Revision Contents

PathSize
llvm/
include/
llvm/
IR/
10 lines
lib/
Target/
AMDGPU/
12 lines
4 lines
test/
CodeGen/
AMDGPU/
14 lines
14 lines

Diff 381824

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

Show First 20 Lines • Show All 1,529 Lines • ▼ Show 20 Linesdef int_amdgcn_mul_i24 : Intrinsic<[llvm_i32_ty],
[IntrNoMem, IntrSpeculatable, IntrWillReturn] [IntrNoMem, IntrSpeculatable, IntrWillReturn]
>; >;
def int_amdgcn_mul_u24 : Intrinsic<[llvm_i32_ty], def int_amdgcn_mul_u24 : Intrinsic<[llvm_i32_ty],
[llvm_i32_ty, llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
[IntrNoMem, IntrSpeculatable, IntrWillReturn] [IntrNoMem, IntrSpeculatable, IntrWillReturn]
>; >;
def int_amdgcn_mulhi_i24 : Intrinsic<[llvm_i32_ty],
[llvm_i32_ty, llvm_i32_ty],
[IntrNoMem, IntrSpeculatable, IntrWillReturn]
>;
def int_amdgcn_mulhi_u24 : Intrinsic<[llvm_i32_ty],
[llvm_i32_ty, llvm_i32_ty],
[IntrNoMem, IntrSpeculatable, IntrWillReturn]
>;
// llvm.amdgcn.ds.gws.init(i32 bar_val, i32 resource_id) // llvm.amdgcn.ds.gws.init(i32 bar_val, i32 resource_id)
// //
// bar_val is the total number of waves that will wait on this // bar_val is the total number of waves that will wait on this
// barrier, minus 1. // barrier, minus 1.
def int_amdgcn_ds_gws_init : def int_amdgcn_ds_gws_init :
GCCBuiltin<"__builtin_amdgcn_ds_gws_init">, GCCBuiltin<"__builtin_amdgcn_ds_gws_init">,
Intrinsic<[], Intrinsic<[],
[llvm_i32_ty, llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
▲ Show 20 Lines • Show All 448 Lines • Show Last 20 Lines

llvm/lib/Target/AMDGPU/AMDGPUInstrInfo.td

Show First 20 Lines • Show All 273 Lines • ▼ Show 20 Lines
// when performing the multiply. The result is a 32-bit value. // when performing the multiply. The result is a 32-bit value.
def AMDGPUmul_u24_impl : SDNode<"AMDGPUISD::MUL_U24", SDTIntBinOp, def AMDGPUmul_u24_impl : SDNode<"AMDGPUISD::MUL_U24", SDTIntBinOp,
[SDNPCommutative, SDNPAssociative] [SDNPCommutative, SDNPAssociative]
>; >;
def AMDGPUmul_i24_impl : SDNode<"AMDGPUISD::MUL_I24", SDTIntBinOp, def AMDGPUmul_i24_impl : SDNode<"AMDGPUISD::MUL_I24", SDTIntBinOp,
[SDNPCommutative, SDNPAssociative] [SDNPCommutative, SDNPAssociative]
>; >;
def AMDGPUmulhi_u24 : SDNode<"AMDGPUISD::MULHI_U24", SDTIntBinOp, def AMDGPUmulhi_u24_impl : SDNode<"AMDGPUISD::MULHI_U24", SDTIntBinOp,
[SDNPCommutative, SDNPAssociative] [SDNPCommutative, SDNPAssociative]
foadUnsubmitted
Done
ReplyInline Actions

Not your fault, but are these mulhi nodes really associative??

foad: Not your fault, but are these mulhi nodes really associative??
>; >;
def AMDGPUmulhi_i24 : SDNode<"AMDGPUISD::MULHI_I24", SDTIntBinOp, def AMDGPUmulhi_i24_impl : SDNode<"AMDGPUISD::MULHI_I24", SDTIntBinOp,
[SDNPCommutative, SDNPAssociative] [SDNPCommutative, SDNPAssociative]
>; >;
def AMDGPUmad_u24 : SDNode<"AMDGPUISD::MAD_U24", AMDGPUDTIntTernaryOp, def AMDGPUmad_u24 : SDNode<"AMDGPUISD::MAD_U24", AMDGPUDTIntTernaryOp,
[] []
>; >;
def AMDGPUmad_i24 : SDNode<"AMDGPUISD::MAD_I24", AMDGPUDTIntTernaryOp, def AMDGPUmad_i24 : SDNode<"AMDGPUISD::MAD_I24", AMDGPUDTIntTernaryOp,
[] []
▲ Show 20 Lines • Show All 136 Lines • ▼ Show 20 Lines
def AMDGPUmul_u24 : PatFrags<(ops node:$src0, node:$src1), def AMDGPUmul_u24 : PatFrags<(ops node:$src0, node:$src1),
[(int_amdgcn_mul_u24 node:$src0, node:$src1), [(int_amdgcn_mul_u24 node:$src0, node:$src1),
(AMDGPUmul_u24_impl node:$src0, node:$src1)]>; (AMDGPUmul_u24_impl node:$src0, node:$src1)]>;
def AMDGPUmul_i24 : PatFrags<(ops node:$src0, node:$src1), def AMDGPUmul_i24 : PatFrags<(ops node:$src0, node:$src1),
[(int_amdgcn_mul_i24 node:$src0, node:$src1), [(int_amdgcn_mul_i24 node:$src0, node:$src1),
(AMDGPUmul_i24_impl node:$src0, node:$src1)]>; (AMDGPUmul_i24_impl node:$src0, node:$src1)]>;
def AMDGPUmulhi_u24 : PatFrags<(ops node:$src0, node:$src1),
[(int_amdgcn_mulhi_u24 node:$src0, node:$src1),
(AMDGPUmulhi_u24_impl node:$src0, node:$src1)]>;
def AMDGPUmulhi_i24 : PatFrags<(ops node:$src0, node:$src1),
[(int_amdgcn_mulhi_i24 node:$src0, node:$src1),
(AMDGPUmulhi_i24_impl node:$src0, node:$src1)]>;
def AMDGPUbfe_i32 : PatFrags<(ops node:$src0, node:$src1, node:$src2), def AMDGPUbfe_i32 : PatFrags<(ops node:$src0, node:$src1, node:$src2),
[(int_amdgcn_sbfe node:$src0, node:$src1, node:$src2), [(int_amdgcn_sbfe node:$src0, node:$src1, node:$src2),
(AMDGPUbfe_i32_impl node:$src0, node:$src1, node:$src2)]>; (AMDGPUbfe_i32_impl node:$src0, node:$src1, node:$src2)]>;
def AMDGPUbfe_u32 : PatFrags<(ops node:$src0, node:$src1, node:$src2), def AMDGPUbfe_u32 : PatFrags<(ops node:$src0, node:$src1, node:$src2),
[(int_amdgcn_ubfe node:$src0, node:$src1, node:$src2), [(int_amdgcn_ubfe node:$src0, node:$src1, node:$src2),
(AMDGPUbfe_u32_impl node:$src0, node:$src1, node:$src2)]>; (AMDGPUbfe_u32_impl node:$src0, node:$src1, node:$src2)]>;
Show All 15 Lines

llvm/lib/Target/AMDGPU/VOP2Instructions.td

Show First 20 Lines • Show All 486 Lines • ▼ Show 20 Lines
let isCommutable = 1 in { let isCommutable = 1 in {
let isReMaterializable = 1 in { let isReMaterializable = 1 in {
defm V_ADD_F32 : VOP2Inst <"v_add_f32", VOP_F32_F32_F32, any_fadd>; defm V_ADD_F32 : VOP2Inst <"v_add_f32", VOP_F32_F32_F32, any_fadd>;
defm V_SUB_F32 : VOP2Inst <"v_sub_f32", VOP_F32_F32_F32, any_fsub>; defm V_SUB_F32 : VOP2Inst <"v_sub_f32", VOP_F32_F32_F32, any_fsub>;
defm V_SUBREV_F32 : VOP2Inst <"v_subrev_f32", VOP_F32_F32_F32, null_frag, "v_sub_f32">; defm V_SUBREV_F32 : VOP2Inst <"v_subrev_f32", VOP_F32_F32_F32, null_frag, "v_sub_f32">;
defm V_MUL_LEGACY_F32 : VOP2Inst <"v_mul_legacy_f32", VOP_F32_F32_F32, AMDGPUfmul_legacy>; defm V_MUL_LEGACY_F32 : VOP2Inst <"v_mul_legacy_f32", VOP_F32_F32_F32, AMDGPUfmul_legacy>;
defm V_MUL_F32 : VOP2Inst <"v_mul_f32", VOP_F32_F32_F32, any_fmul>; defm V_MUL_F32 : VOP2Inst <"v_mul_f32", VOP_F32_F32_F32, any_fmul>;
defm V_MUL_I32_I24 : VOP2Inst <"v_mul_i32_i24", VOP_I32_I32_I32_ARITH, AMDGPUmul_i24>; defm V_MUL_I32_I24 : VOP2Inst <"v_mul_i32_i24", VOP_I32_I32_I32_ARITH, AMDGPUmul_i24>;
defm V_MUL_HI_I32_I24 : VOP2Inst <"v_mul_hi_i32_i24", VOP_PAT_GEN<VOP_I32_I32_I32, 2>, AMDGPUmulhi_i24>; defm V_MUL_HI_I32_I24 : VOP2Inst <"v_mul_hi_i32_i24", VOP_I32_I32_I32, AMDGPUmulhi_i24>;
defm V_MUL_U32_U24 : VOP2Inst <"v_mul_u32_u24", VOP_I32_I32_I32_ARITH, AMDGPUmul_u24>; defm V_MUL_U32_U24 : VOP2Inst <"v_mul_u32_u24", VOP_I32_I32_I32_ARITH, AMDGPUmul_u24>;
defm V_MUL_HI_U32_U24 : VOP2Inst <"v_mul_hi_u32_u24", VOP_PAT_GEN<VOP_I32_I32_I32, 2>, AMDGPUmulhi_u24>; defm V_MUL_HI_U32_U24 : VOP2Inst <"v_mul_hi_u32_u24", VOP_I32_I32_I32, AMDGPUmulhi_u24>;
defm V_MIN_F32 : VOP2Inst <"v_min_f32", VOP_F32_F32_F32, fminnum_like>; defm V_MIN_F32 : VOP2Inst <"v_min_f32", VOP_F32_F32_F32, fminnum_like>;
defm V_MAX_F32 : VOP2Inst <"v_max_f32", VOP_F32_F32_F32, fmaxnum_like>; defm V_MAX_F32 : VOP2Inst <"v_max_f32", VOP_F32_F32_F32, fmaxnum_like>;
defm V_MIN_I32 : VOP2Inst <"v_min_i32", VOP_PAT_GEN<VOP_I32_I32_I32>, smin>; defm V_MIN_I32 : VOP2Inst <"v_min_i32", VOP_PAT_GEN<VOP_I32_I32_I32>, smin>;
defm V_MAX_I32 : VOP2Inst <"v_max_i32", VOP_PAT_GEN<VOP_I32_I32_I32>, smax>; defm V_MAX_I32 : VOP2Inst <"v_max_i32", VOP_PAT_GEN<VOP_I32_I32_I32>, smax>;
defm V_MIN_U32 : VOP2Inst <"v_min_u32", VOP_PAT_GEN<VOP_I32_I32_I32>, umin>; defm V_MIN_U32 : VOP2Inst <"v_min_u32", VOP_PAT_GEN<VOP_I32_I32_I32>, umin>;
defm V_MAX_U32 : VOP2Inst <"v_max_u32", VOP_PAT_GEN<VOP_I32_I32_I32>, umax>; defm V_MAX_U32 : VOP2Inst <"v_max_u32", VOP_PAT_GEN<VOP_I32_I32_I32>, umax>;
defm V_LSHRREV_B32 : VOP2Inst <"v_lshrrev_b32", VOP_I32_I32_I32, lshr_rev, "v_lshr_b32">; defm V_LSHRREV_B32 : VOP2Inst <"v_lshrrev_b32", VOP_I32_I32_I32, lshr_rev, "v_lshr_b32">;
defm V_ASHRREV_I32 : VOP2Inst <"v_ashrrev_i32", VOP_I32_I32_I32, ashr_rev, "v_ashr_i32">; defm V_ASHRREV_I32 : VOP2Inst <"v_ashrrev_i32", VOP_I32_I32_I32, ashr_rev, "v_ashr_i32">;
▲ Show 20 Lines • Show All 1,240 Lines • Show Last 20 Lines

llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mulhi.i24.ll

  • This file was added.
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
; RUN: llc -march=amdgcn -mcpu=gfx900 -verify-machineinstrs < %s | FileCheck %s
foadUnsubmitted
Not Done
ReplyInline Actions

Can you add a -global-isel run line? Does it Just Work?

foad: Can you add a -global-isel run line? Does it Just Work?
abinavppAuthorUnsubmitted
Done
ReplyInline Actions

It's not working at the moment. I've added the mir version for now.

abinavpp: It's not working at the moment. I've added the mir version for now.
foadUnsubmitted
Done
ReplyInline Actions

It should be trivial to implement, just handle the mulhi intrinsics the same as the mul intrinsics in AMDGPURegisterBankInfo.

foad: It should be trivial to implement, just handle the mulhi intrinsics the same as the mul…
define i32 @basic(i32 %a, i32 %b) {
; CHECK-LABEL: basic:
; CHECK: ; %bb.0:
; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
; CHECK-NEXT: v_mul_hi_i32_i24_e32 v0, v0, v1
; CHECK-NEXT: s_setpc_b64 s[30:31]
%mul = call i32 @llvm.amdgcn.mulhi.i24(i32 %a, i32 %b)
ret i32 %mul
}
declare i32 @llvm.amdgcn.mulhi.i24(i32, i32)

llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mulhi.u24.ll

  • This file was added.
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
; RUN: llc -march=amdgcn -mcpu=gfx900 -verify-machineinstrs < %s | FileCheck %s
define i32 @basic(i32 %a, i32 %b) {
; CHECK-LABEL: basic:
; CHECK: ; %bb.0:
; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
; CHECK-NEXT: v_mul_hi_u32_u24_e32 v0, v0, v1
; CHECK-NEXT: s_setpc_b64 s[30:31]
%mul = call i32 @llvm.amdgcn.mulhi.u24(i32 %a, i32 %b)
ret i32 %mul
}
declare i32 @llvm.amdgcn.mulhi.u24(i32, i32)