This is an archive of the discontinued LLVM Phabricator instance.

Differential D142575

[AMDGPU] Remove predicates from real dot instructions. NFCI.
ClosedPublic

Authored by rampitec on Jan 25 2023, 12:05 PM.

Details

Reviewers
foad
Commits
rGccdd4ae1db23: [AMDGPU] Remove predicates from real dot instructions. NFCI.
Summary

These are copied from pseudos automatically.

Diff Detail

Unit TestsFailed

TimeTest
60,070 msx64 debian > libFuzzer.libFuzzer::fuzzer-leak.test
60,160 msx64 debian > libFuzzer.libFuzzer::minimize_crash.test
60,060 msx64 debian > libFuzzer.libFuzzer::value-profile-load.test

Event Timeline

rampitec created this revision.Jan 25 2023, 12:05 PM
Herald added a project: Restricted Project. · View Herald TranscriptJan 25 2023, 12:05 PM
Herald added subscribers: kosarev, StephenFan, kerbowa and 7 others. · View Herald Transcript
rampitec requested review of this revision.Jan 25 2023, 12:05 PM
Herald added a project: Restricted Project. · View Herald TranscriptJan 25 2023, 12:05 PM
Herald added a subscriber: wdng. · View Herald Transcript
rampitec added a child revision: D142507: [AMDGPU] Split dot7 feature.Jan 25 2023, 12:09 PM
rampitec mentioned this in D142507: [AMDGPU] Split dot7 feature.
foad accepted this revision.Jan 25 2023, 12:35 PM

LGTM, thanks!

This revision is now accepted and ready to land.Jan 25 2023, 12:35 PM
This revision was landed with ongoing or failed builds.Jan 25 2023, 12:50 PM
Closed by commit rGccdd4ae1db23: [AMDGPU] Remove predicates from real dot instructions. NFCI. (authored by rampitec). · Explain Why
This revision was automatically updated to reflect the committed changes.
rampitec added a commit: rGccdd4ae1db23: [AMDGPU] Remove predicates from real dot instructions. NFCI..
Harbormaster completed remote builds in B209947: Diff 492218.Jan 25 2023, 1:09 PM

Revision Contents

PathSize
llvm/
lib/
Target/
AMDGPU/
24 lines

Diff 492218

llvm/lib/Target/AMDGPU/VOP3PInstructions.td

Show First 20 Lines • Show All 1,075 Lines • ▼ Show 20 Lines
// different DecoderNamespace to avoid the ambiguity. // different DecoderNamespace to avoid the ambiguity.
defm V_FMA_MIX_F32 : VOP3P_Real_vi <0x20>; defm V_FMA_MIX_F32 : VOP3P_Real_vi <0x20>;
defm V_FMA_MIXLO_F16 : VOP3P_Real_vi <0x21>; defm V_FMA_MIXLO_F16 : VOP3P_Real_vi <0x21>;
defm V_FMA_MIXHI_F16 : VOP3P_Real_vi <0x22>; defm V_FMA_MIXHI_F16 : VOP3P_Real_vi <0x22>;
} }
} }
let SubtargetPredicate = HasDot2Insts in {
defm V_DOT2_I32_I16 : VOP3P_Real_vi <0x26>; defm V_DOT2_I32_I16 : VOP3P_Real_vi <0x26>;
defm V_DOT2_U32_U16 : VOP3P_Real_vi <0x27>; defm V_DOT2_U32_U16 : VOP3P_Real_vi <0x27>;
} // End SubtargetPredicate = HasDot2Insts
let SubtargetPredicate = HasDot7Insts in {
defm V_DOT2_F32_F16 : VOP3P_Real_vi <0x23>; defm V_DOT2_F32_F16 : VOP3P_Real_vi <0x23>;
defm V_DOT4_U32_U8 : VOP3P_Real_vi <0x29>; defm V_DOT4_U32_U8 : VOP3P_Real_vi <0x29>;
defm V_DOT8_U32_U4 : VOP3P_Real_vi <0x2b>; defm V_DOT8_U32_U4 : VOP3P_Real_vi <0x2b>;
} // End SubtargetPredicate = HasDot7Insts
let SubtargetPredicate = HasDot1Insts in {
defm V_DOT4_I32_I8 : VOP3P_Real_vi <0x28>; defm V_DOT4_I32_I8 : VOP3P_Real_vi <0x28>;
defm V_DOT8_I32_I4 : VOP3P_Real_vi <0x2a>; defm V_DOT8_I32_I4 : VOP3P_Real_vi <0x2a>;
} // End SubtargetPredicate = HasDot1Insts
let SubtargetPredicate = HasMAIInsts in { let SubtargetPredicate = HasMAIInsts in {
defm V_ACCVGPR_READ_B32 : VOP3P_Real_MAI <0x58>; defm V_ACCVGPR_READ_B32 : VOP3P_Real_MAI <0x58>;
defm V_ACCVGPR_WRITE_B32 : VOP3P_Real_MAI <0x59>; defm V_ACCVGPR_WRITE_B32 : VOP3P_Real_MAI <0x59>;
defm V_MFMA_F32_32X32X1F32 : VOP3P_Real_MFMA <0x40, "v_mfma_f32_32x32x1_2b_f32">; defm V_MFMA_F32_32X32X1F32 : VOP3P_Real_MFMA <0x40, "v_mfma_f32_32x32x1_2b_f32">;
defm V_MFMA_F32_16X16X1F32 : VOP3P_Real_MFMA <0x41, "v_mfma_f32_16x16x1_4b_f32">; defm V_MFMA_F32_16X16X1F32 : VOP3P_Real_MFMA <0x41, "v_mfma_f32_16x16x1_4b_f32">;
defm V_MFMA_F32_4X4X1F32 : VOP3P_Real_MFMA <0x42, "v_mfma_f32_4x4x1_16b_f32">; defm V_MFMA_F32_4X4X1F32 : VOP3P_Real_MFMA <0x42, "v_mfma_f32_4x4x1_16b_f32">;
defm V_MFMA_F32_32X32X2F32 : VOP3P_Real_MFMA <0x44, "v_mfma_f32_32x32x2_f32">; defm V_MFMA_F32_32X32X2F32 : VOP3P_Real_MFMA <0x44, "v_mfma_f32_32x32x2_f32">;
▲ Show 20 Lines • Show All 106 Lines • ▼ Show 20 Lines
defm V_PK_ADD_F16 : VOP3P_Real_gfx10_gfx11<0x0f>; defm V_PK_ADD_F16 : VOP3P_Real_gfx10_gfx11<0x0f>;
defm V_PK_MUL_F16 : VOP3P_Real_gfx10_gfx11<0x10>; defm V_PK_MUL_F16 : VOP3P_Real_gfx10_gfx11<0x10>;
defm V_PK_MIN_F16 : VOP3P_Real_gfx10_gfx11<0x11>; defm V_PK_MIN_F16 : VOP3P_Real_gfx10_gfx11<0x11>;
defm V_PK_MAX_F16 : VOP3P_Real_gfx10_gfx11<0x12>; defm V_PK_MAX_F16 : VOP3P_Real_gfx10_gfx11<0x12>;
defm V_FMA_MIX_F32 : VOP3P_Real_gfx10_gfx11_Triple <0x20>; defm V_FMA_MIX_F32 : VOP3P_Real_gfx10_gfx11_Triple <0x20>;
defm V_FMA_MIXLO_F16 : VOP3P_Real_gfx10_gfx11_Triple <0x21>; defm V_FMA_MIXLO_F16 : VOP3P_Real_gfx10_gfx11_Triple <0x21>;
defm V_FMA_MIXHI_F16 : VOP3P_Real_gfx10_gfx11_Triple <0x22>; defm V_FMA_MIXHI_F16 : VOP3P_Real_gfx10_gfx11_Triple <0x22>;
let SubtargetPredicate = HasDot2Insts in {
defm V_DOT2_I32_I16 : VOP3P_Real_gfx10 <0x14>; defm V_DOT2_I32_I16 : VOP3P_Real_gfx10 <0x14>;
defm V_DOT2_U32_U16 : VOP3P_Real_gfx10 <0x15>; defm V_DOT2_U32_U16 : VOP3P_Real_gfx10 <0x15>;
} // End SubtargetPredicate = HasDot2Insts
let SubtargetPredicate = HasDot7Insts in {
defm V_DOT2_F32_F16 : VOP3P_Real_gfx10_gfx11_Triple <0x13>; defm V_DOT2_F32_F16 : VOP3P_Real_gfx10_gfx11_Triple <0x13>;
defm V_DOT4_U32_U8 : VOP3P_Real_gfx10_gfx11 <0x17>; defm V_DOT4_U32_U8 : VOP3P_Real_gfx10_gfx11 <0x17>;
defm V_DOT8_U32_U4 : VOP3P_Real_gfx10_gfx11 <0x19>; defm V_DOT8_U32_U4 : VOP3P_Real_gfx10_gfx11 <0x19>;
} // End SubtargetPredicate = HasDot7Insts
let SubtargetPredicate = HasDot1Insts in {
defm V_DOT4_I32_I8 : VOP3P_Real_gfx10 <0x16>; defm V_DOT4_I32_I8 : VOP3P_Real_gfx10 <0x16>;
defm V_DOT8_I32_I4 : VOP3P_Real_gfx10 <0x18>; defm V_DOT8_I32_I4 : VOP3P_Real_gfx10 <0x18>;
} // End SubtargetPredicate = HasDot1Insts