reference, declarationdefinition
definition → references, declarations, derived classes, virtual overrides
reference to multiple definitions → definitions
unreferenced

References

gen/lib/Target/AMDGPU/AMDGPUGenAsmMatcher.inc
21436   { 13172 /* v_accvgpr_read_b32 */, AMDGPU::V_ACCVGPR_READ_B32_vi, Convert__Reg1_0__Reg1_1, AMFBS_HasMAIInsts_HasMAIInsts, { MCK_VGPR_32, MCK_AGPR_32 }, },
21437   { 13191 /* v_accvgpr_write_b32 */, AMDGPU::V_ACCVGPR_WRITE_B32_vi, Convert__Reg1_0__VISrcB321_1, AMFBS_HasMAIInsts_HasMAIInsts, { MCK_AGPR_32, MCK_VISrcB32 }, },
22484   { 24275 /* v_mfma_f32_16x16x16f16 */, AMDGPU::V_MFMA_F32_16X16X16F16_vi, Convert__Reg1_0__Reg1_1__Reg1_2__AISrc_128B321_3__ImmCBSZ1_4__ImmABID1_5__ImmBLGP1_6, AMFBS_HasMAIInsts_HasMAIInsts, { MCK_AReg_128, MCK_AV_64, MCK_AV_64, MCK_AISrc_128B32, MCK_ImmCBSZ, MCK_ImmABID, MCK_ImmBLGP }, },
22485   { 24298 /* v_mfma_f32_16x16x1f32 */, AMDGPU::V_MFMA_F32_16X16X1F32_vi, Convert__Reg1_0__Reg1_1__Reg1_2__AISrc_512F321_3__ImmCBSZ1_4__ImmABID1_5__ImmBLGP1_6, AMFBS_HasMAIInsts_HasMAIInsts, { MCK_AReg_512, MCK_AV_32, MCK_AV_32, MCK_AISrc_512F32, MCK_ImmCBSZ, MCK_ImmABID, MCK_ImmBLGP }, },
22486   { 24320 /* v_mfma_f32_16x16x2bf16 */, AMDGPU::V_MFMA_F32_16X16X2BF16_vi, Convert__Reg1_0__Reg1_1__Reg1_2__AISrc_512B321_3__ImmCBSZ1_4__ImmABID1_5__ImmBLGP1_6, AMFBS_HasMAIInsts_HasMAIInsts, { MCK_AReg_512, MCK_AV_32, MCK_AV_32, MCK_AISrc_512B32, MCK_ImmCBSZ, MCK_ImmABID, MCK_ImmBLGP }, },
22487   { 24343 /* v_mfma_f32_16x16x4f16 */, AMDGPU::V_MFMA_F32_16X16X4F16_vi, Convert__Reg1_0__Reg1_1__Reg1_2__AISrc_512B321_3__ImmCBSZ1_4__ImmABID1_5__ImmBLGP1_6, AMFBS_HasMAIInsts_HasMAIInsts, { MCK_AReg_512, MCK_AV_64, MCK_AV_64, MCK_AISrc_512B32, MCK_ImmCBSZ, MCK_ImmABID, MCK_ImmBLGP }, },
22488   { 24365 /* v_mfma_f32_16x16x4f32 */, AMDGPU::V_MFMA_F32_16X16X4F32_vi, Convert__Reg1_0__Reg1_1__Reg1_2__AISrc_128F321_3__ImmCBSZ1_4__ImmABID1_5__ImmBLGP1_6, AMFBS_HasMAIInsts_HasMAIInsts, { MCK_AReg_128, MCK_AV_32, MCK_AV_32, MCK_AISrc_128F32, MCK_ImmCBSZ, MCK_ImmABID, MCK_ImmBLGP }, },
22489   { 24387 /* v_mfma_f32_16x16x8bf16 */, AMDGPU::V_MFMA_F32_16X16X8BF16_vi, Convert__Reg1_0__Reg1_1__Reg1_2__AISrc_128B321_3__ImmCBSZ1_4__ImmABID1_5__ImmBLGP1_6, AMFBS_HasMAIInsts_HasMAIInsts, { MCK_AReg_128, MCK_AV_32, MCK_AV_32, MCK_AISrc_128B32, MCK_ImmCBSZ, MCK_ImmABID, MCK_ImmBLGP }, },
22490   { 24410 /* v_mfma_f32_32x32x1f32 */, AMDGPU::V_MFMA_F32_32X32X1F32_vi, Convert__Reg1_0__Reg1_1__Reg1_2__AISrc_1024F321_3__ImmCBSZ1_4__ImmABID1_5__ImmBLGP1_6, AMFBS_HasMAIInsts_HasMAIInsts, { MCK_AReg_1024, MCK_AV_32, MCK_AV_32, MCK_AISrc_1024F32, MCK_ImmCBSZ, MCK_ImmABID, MCK_ImmBLGP }, },
22491   { 24432 /* v_mfma_f32_32x32x2bf16 */, AMDGPU::V_MFMA_F32_32X32X2BF16_vi, Convert__Reg1_0__Reg1_1__Reg1_2__AISrc_1024B321_3__ImmCBSZ1_4__ImmABID1_5__ImmBLGP1_6, AMFBS_HasMAIInsts_HasMAIInsts, { MCK_AReg_1024, MCK_AV_32, MCK_AV_32, MCK_AISrc_1024B32, MCK_ImmCBSZ, MCK_ImmABID, MCK_ImmBLGP }, },
22492   { 24455 /* v_mfma_f32_32x32x2f32 */, AMDGPU::V_MFMA_F32_32X32X2F32_vi, Convert__Reg1_0__Reg1_1__Reg1_2__AISrc_512F321_3__ImmCBSZ1_4__ImmABID1_5__ImmBLGP1_6, AMFBS_HasMAIInsts_HasMAIInsts, { MCK_AReg_512, MCK_AV_32, MCK_AV_32, MCK_AISrc_512F32, MCK_ImmCBSZ, MCK_ImmABID, MCK_ImmBLGP }, },
22493   { 24477 /* v_mfma_f32_32x32x4bf16 */, AMDGPU::V_MFMA_F32_32X32X4BF16_vi, Convert__Reg1_0__Reg1_1__Reg1_2__AISrc_512B321_3__ImmCBSZ1_4__ImmABID1_5__ImmBLGP1_6, AMFBS_HasMAIInsts_HasMAIInsts, { MCK_AReg_512, MCK_AV_32, MCK_AV_32, MCK_AISrc_512B32, MCK_ImmCBSZ, MCK_ImmABID, MCK_ImmBLGP }, },
22494   { 24500 /* v_mfma_f32_32x32x4f16 */, AMDGPU::V_MFMA_F32_32X32X4F16_vi, Convert__Reg1_0__Reg1_1__Reg1_2__AISrc_1024B321_3__ImmCBSZ1_4__ImmABID1_5__ImmBLGP1_6, AMFBS_HasMAIInsts_HasMAIInsts, { MCK_AReg_1024, MCK_AV_64, MCK_AV_64, MCK_AISrc_1024B32, MCK_ImmCBSZ, MCK_ImmABID, MCK_ImmBLGP }, },
22495   { 24522 /* v_mfma_f32_32x32x8f16 */, AMDGPU::V_MFMA_F32_32X32X8F16_vi, Convert__Reg1_0__Reg1_1__Reg1_2__AISrc_512B321_3__ImmCBSZ1_4__ImmABID1_5__ImmBLGP1_6, AMFBS_HasMAIInsts_HasMAIInsts, { MCK_AReg_512, MCK_AV_64, MCK_AV_64, MCK_AISrc_512B32, MCK_ImmCBSZ, MCK_ImmABID, MCK_ImmBLGP }, },
22496   { 24544 /* v_mfma_f32_4x4x1f32 */, AMDGPU::V_MFMA_F32_4X4X1F32_vi, Convert__Reg1_0__Reg1_1__Reg1_2__AISrc_128F321_3__ImmCBSZ1_4__ImmABID1_5__ImmBLGP1_6, AMFBS_HasMAIInsts_HasMAIInsts, { MCK_AReg_128, MCK_AV_32, MCK_AV_32, MCK_AISrc_128F32, MCK_ImmCBSZ, MCK_ImmABID, MCK_ImmBLGP }, },
22497   { 24564 /* v_mfma_f32_4x4x2bf16 */, AMDGPU::V_MFMA_F32_4X4X2BF16_vi, Convert__Reg1_0__Reg1_1__Reg1_2__AISrc_128B321_3__ImmCBSZ1_4__ImmABID1_5__ImmBLGP1_6, AMFBS_HasMAIInsts_HasMAIInsts, { MCK_AReg_128, MCK_AV_32, MCK_AV_32, MCK_AISrc_128B32, MCK_ImmCBSZ, MCK_ImmABID, MCK_ImmBLGP }, },
22498   { 24585 /* v_mfma_f32_4x4x4f16 */, AMDGPU::V_MFMA_F32_4X4X4F16_vi, Convert__Reg1_0__Reg1_1__Reg1_2__AISrc_128B321_3__ImmCBSZ1_4__ImmABID1_5__ImmBLGP1_6, AMFBS_HasMAIInsts_HasMAIInsts, { MCK_AReg_128, MCK_AV_64, MCK_AV_64, MCK_AISrc_128B32, MCK_ImmCBSZ, MCK_ImmABID, MCK_ImmBLGP }, },
22499   { 24605 /* v_mfma_i32_16x16x16i8 */, AMDGPU::V_MFMA_I32_16X16X16I8_vi, Convert__Reg1_0__Reg1_1__Reg1_2__AISrc_128B321_3__ImmCBSZ1_4__ImmABID1_5__ImmBLGP1_6, AMFBS_HasMAIInsts_HasMAIInsts, { MCK_AReg_128, MCK_AV_32, MCK_AV_32, MCK_AISrc_128B32, MCK_ImmCBSZ, MCK_ImmABID, MCK_ImmBLGP }, },
22500   { 24627 /* v_mfma_i32_16x16x4i8 */, AMDGPU::V_MFMA_I32_16X16X4I8_vi, Convert__Reg1_0__Reg1_1__Reg1_2__AISrc_512B321_3__ImmCBSZ1_4__ImmABID1_5__ImmBLGP1_6, AMFBS_HasMAIInsts_HasMAIInsts, { MCK_AReg_512, MCK_AV_32, MCK_AV_32, MCK_AISrc_512B32, MCK_ImmCBSZ, MCK_ImmABID, MCK_ImmBLGP }, },
22501   { 24648 /* v_mfma_i32_32x32x4i8 */, AMDGPU::V_MFMA_I32_32X32X4I8_vi, Convert__Reg1_0__Reg1_1__Reg1_2__AISrc_1024B321_3__ImmCBSZ1_4__ImmABID1_5__ImmBLGP1_6, AMFBS_HasMAIInsts_HasMAIInsts, { MCK_AReg_1024, MCK_AV_32, MCK_AV_32, MCK_AISrc_1024B32, MCK_ImmCBSZ, MCK_ImmABID, MCK_ImmBLGP }, },
22502   { 24669 /* v_mfma_i32_32x32x8i8 */, AMDGPU::V_MFMA_I32_32X32X8I8_vi, Convert__Reg1_0__Reg1_1__Reg1_2__AISrc_512B321_3__ImmCBSZ1_4__ImmABID1_5__ImmBLGP1_6, AMFBS_HasMAIInsts_HasMAIInsts, { MCK_AReg_512, MCK_AV_32, MCK_AV_32, MCK_AISrc_512B32, MCK_ImmCBSZ, MCK_ImmABID, MCK_ImmBLGP }, },
22503   { 24690 /* v_mfma_i32_4x4x4i8 */, AMDGPU::V_MFMA_I32_4X4X4I8_vi, Convert__Reg1_0__Reg1_1__Reg1_2__AISrc_128B321_3__ImmCBSZ1_4__ImmABID1_5__ImmBLGP1_6, AMFBS_HasMAIInsts_HasMAIInsts, { MCK_AReg_128, MCK_AV_32, MCK_AV_32, MCK_AISrc_128B32, MCK_ImmCBSZ, MCK_ImmABID, MCK_ImmBLGP }, },
78282   { 24275 /* v_mfma_f32_16x16x16f16 */, 64 /* 6 */, MCK_ImmBLGP, AMFBS_HasMAIInsts_HasMAIInsts },
78283   { 24275 /* v_mfma_f32_16x16x16f16 */, 16 /* 4 */, MCK_ImmCBSZ, AMFBS_HasMAIInsts_HasMAIInsts },
78284   { 24275 /* v_mfma_f32_16x16x16f16 */, 32 /* 5 */, MCK_ImmABID, AMFBS_HasMAIInsts_HasMAIInsts },
78285   { 24298 /* v_mfma_f32_16x16x1f32 */, 64 /* 6 */, MCK_ImmBLGP, AMFBS_HasMAIInsts_HasMAIInsts },
78286   { 24298 /* v_mfma_f32_16x16x1f32 */, 16 /* 4 */, MCK_ImmCBSZ, AMFBS_HasMAIInsts_HasMAIInsts },
78287   { 24298 /* v_mfma_f32_16x16x1f32 */, 32 /* 5 */, MCK_ImmABID, AMFBS_HasMAIInsts_HasMAIInsts },
78288   { 24320 /* v_mfma_f32_16x16x2bf16 */, 64 /* 6 */, MCK_ImmBLGP, AMFBS_HasMAIInsts_HasMAIInsts },
78289   { 24320 /* v_mfma_f32_16x16x2bf16 */, 16 /* 4 */, MCK_ImmCBSZ, AMFBS_HasMAIInsts_HasMAIInsts },
78290   { 24320 /* v_mfma_f32_16x16x2bf16 */, 32 /* 5 */, MCK_ImmABID, AMFBS_HasMAIInsts_HasMAIInsts },
78291   { 24343 /* v_mfma_f32_16x16x4f16 */, 64 /* 6 */, MCK_ImmBLGP, AMFBS_HasMAIInsts_HasMAIInsts },
78292   { 24343 /* v_mfma_f32_16x16x4f16 */, 16 /* 4 */, MCK_ImmCBSZ, AMFBS_HasMAIInsts_HasMAIInsts },
78293   { 24343 /* v_mfma_f32_16x16x4f16 */, 32 /* 5 */, MCK_ImmABID, AMFBS_HasMAIInsts_HasMAIInsts },
78294   { 24365 /* v_mfma_f32_16x16x4f32 */, 64 /* 6 */, MCK_ImmBLGP, AMFBS_HasMAIInsts_HasMAIInsts },
78295   { 24365 /* v_mfma_f32_16x16x4f32 */, 16 /* 4 */, MCK_ImmCBSZ, AMFBS_HasMAIInsts_HasMAIInsts },
78296   { 24365 /* v_mfma_f32_16x16x4f32 */, 32 /* 5 */, MCK_ImmABID, AMFBS_HasMAIInsts_HasMAIInsts },
78297   { 24387 /* v_mfma_f32_16x16x8bf16 */, 64 /* 6 */, MCK_ImmBLGP, AMFBS_HasMAIInsts_HasMAIInsts },
78298   { 24387 /* v_mfma_f32_16x16x8bf16 */, 16 /* 4 */, MCK_ImmCBSZ, AMFBS_HasMAIInsts_HasMAIInsts },
78299   { 24387 /* v_mfma_f32_16x16x8bf16 */, 32 /* 5 */, MCK_ImmABID, AMFBS_HasMAIInsts_HasMAIInsts },
78300   { 24410 /* v_mfma_f32_32x32x1f32 */, 64 /* 6 */, MCK_ImmBLGP, AMFBS_HasMAIInsts_HasMAIInsts },
78301   { 24410 /* v_mfma_f32_32x32x1f32 */, 16 /* 4 */, MCK_ImmCBSZ, AMFBS_HasMAIInsts_HasMAIInsts },
78302   { 24410 /* v_mfma_f32_32x32x1f32 */, 32 /* 5 */, MCK_ImmABID, AMFBS_HasMAIInsts_HasMAIInsts },
78303   { 24432 /* v_mfma_f32_32x32x2bf16 */, 64 /* 6 */, MCK_ImmBLGP, AMFBS_HasMAIInsts_HasMAIInsts },
78304   { 24432 /* v_mfma_f32_32x32x2bf16 */, 16 /* 4 */, MCK_ImmCBSZ, AMFBS_HasMAIInsts_HasMAIInsts },
78305   { 24432 /* v_mfma_f32_32x32x2bf16 */, 32 /* 5 */, MCK_ImmABID, AMFBS_HasMAIInsts_HasMAIInsts },
78306   { 24455 /* v_mfma_f32_32x32x2f32 */, 64 /* 6 */, MCK_ImmBLGP, AMFBS_HasMAIInsts_HasMAIInsts },
78307   { 24455 /* v_mfma_f32_32x32x2f32 */, 16 /* 4 */, MCK_ImmCBSZ, AMFBS_HasMAIInsts_HasMAIInsts },
78308   { 24455 /* v_mfma_f32_32x32x2f32 */, 32 /* 5 */, MCK_ImmABID, AMFBS_HasMAIInsts_HasMAIInsts },
78309   { 24477 /* v_mfma_f32_32x32x4bf16 */, 64 /* 6 */, MCK_ImmBLGP, AMFBS_HasMAIInsts_HasMAIInsts },
78310   { 24477 /* v_mfma_f32_32x32x4bf16 */, 16 /* 4 */, MCK_ImmCBSZ, AMFBS_HasMAIInsts_HasMAIInsts },
78311   { 24477 /* v_mfma_f32_32x32x4bf16 */, 32 /* 5 */, MCK_ImmABID, AMFBS_HasMAIInsts_HasMAIInsts },
78312   { 24500 /* v_mfma_f32_32x32x4f16 */, 64 /* 6 */, MCK_ImmBLGP, AMFBS_HasMAIInsts_HasMAIInsts },
78313   { 24500 /* v_mfma_f32_32x32x4f16 */, 16 /* 4 */, MCK_ImmCBSZ, AMFBS_HasMAIInsts_HasMAIInsts },
78314   { 24500 /* v_mfma_f32_32x32x4f16 */, 32 /* 5 */, MCK_ImmABID, AMFBS_HasMAIInsts_HasMAIInsts },
78315   { 24522 /* v_mfma_f32_32x32x8f16 */, 64 /* 6 */, MCK_ImmBLGP, AMFBS_HasMAIInsts_HasMAIInsts },
78316   { 24522 /* v_mfma_f32_32x32x8f16 */, 16 /* 4 */, MCK_ImmCBSZ, AMFBS_HasMAIInsts_HasMAIInsts },
78317   { 24522 /* v_mfma_f32_32x32x8f16 */, 32 /* 5 */, MCK_ImmABID, AMFBS_HasMAIInsts_HasMAIInsts },
78318   { 24544 /* v_mfma_f32_4x4x1f32 */, 64 /* 6 */, MCK_ImmBLGP, AMFBS_HasMAIInsts_HasMAIInsts },
78319   { 24544 /* v_mfma_f32_4x4x1f32 */, 16 /* 4 */, MCK_ImmCBSZ, AMFBS_HasMAIInsts_HasMAIInsts },
78320   { 24544 /* v_mfma_f32_4x4x1f32 */, 32 /* 5 */, MCK_ImmABID, AMFBS_HasMAIInsts_HasMAIInsts },
78321   { 24564 /* v_mfma_f32_4x4x2bf16 */, 64 /* 6 */, MCK_ImmBLGP, AMFBS_HasMAIInsts_HasMAIInsts },
78322   { 24564 /* v_mfma_f32_4x4x2bf16 */, 16 /* 4 */, MCK_ImmCBSZ, AMFBS_HasMAIInsts_HasMAIInsts },
78323   { 24564 /* v_mfma_f32_4x4x2bf16 */, 32 /* 5 */, MCK_ImmABID, AMFBS_HasMAIInsts_HasMAIInsts },
78324   { 24585 /* v_mfma_f32_4x4x4f16 */, 64 /* 6 */, MCK_ImmBLGP, AMFBS_HasMAIInsts_HasMAIInsts },
78325   { 24585 /* v_mfma_f32_4x4x4f16 */, 16 /* 4 */, MCK_ImmCBSZ, AMFBS_HasMAIInsts_HasMAIInsts },
78326   { 24585 /* v_mfma_f32_4x4x4f16 */, 32 /* 5 */, MCK_ImmABID, AMFBS_HasMAIInsts_HasMAIInsts },
78327   { 24605 /* v_mfma_i32_16x16x16i8 */, 64 /* 6 */, MCK_ImmBLGP, AMFBS_HasMAIInsts_HasMAIInsts },
78328   { 24605 /* v_mfma_i32_16x16x16i8 */, 16 /* 4 */, MCK_ImmCBSZ, AMFBS_HasMAIInsts_HasMAIInsts },
78329   { 24605 /* v_mfma_i32_16x16x16i8 */, 32 /* 5 */, MCK_ImmABID, AMFBS_HasMAIInsts_HasMAIInsts },
78330   { 24627 /* v_mfma_i32_16x16x4i8 */, 64 /* 6 */, MCK_ImmBLGP, AMFBS_HasMAIInsts_HasMAIInsts },
78331   { 24627 /* v_mfma_i32_16x16x4i8 */, 16 /* 4 */, MCK_ImmCBSZ, AMFBS_HasMAIInsts_HasMAIInsts },
78332   { 24627 /* v_mfma_i32_16x16x4i8 */, 32 /* 5 */, MCK_ImmABID, AMFBS_HasMAIInsts_HasMAIInsts },
78333   { 24648 /* v_mfma_i32_32x32x4i8 */, 64 /* 6 */, MCK_ImmBLGP, AMFBS_HasMAIInsts_HasMAIInsts },
78334   { 24648 /* v_mfma_i32_32x32x4i8 */, 16 /* 4 */, MCK_ImmCBSZ, AMFBS_HasMAIInsts_HasMAIInsts },
78335   { 24648 /* v_mfma_i32_32x32x4i8 */, 32 /* 5 */, MCK_ImmABID, AMFBS_HasMAIInsts_HasMAIInsts },
78336   { 24669 /* v_mfma_i32_32x32x8i8 */, 64 /* 6 */, MCK_ImmBLGP, AMFBS_HasMAIInsts_HasMAIInsts },
78337   { 24669 /* v_mfma_i32_32x32x8i8 */, 16 /* 4 */, MCK_ImmCBSZ, AMFBS_HasMAIInsts_HasMAIInsts },
78338   { 24669 /* v_mfma_i32_32x32x8i8 */, 32 /* 5 */, MCK_ImmABID, AMFBS_HasMAIInsts_HasMAIInsts },
78339   { 24690 /* v_mfma_i32_4x4x4i8 */, 64 /* 6 */, MCK_ImmBLGP, AMFBS_HasMAIInsts_HasMAIInsts },
78340   { 24690 /* v_mfma_i32_4x4x4i8 */, 16 /* 4 */, MCK_ImmCBSZ, AMFBS_HasMAIInsts_HasMAIInsts },
78341   { 24690 /* v_mfma_i32_4x4x4i8 */, 32 /* 5 */, MCK_ImmABID, AMFBS_HasMAIInsts_HasMAIInsts },