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

References

gen/lib/Target/AMDGPU/AMDGPUGenAsmMatcher.inc
 4268   MCK_LAST_REGISTER = MCK_AV_32,
 5152     return B == MCK_AV_32;
 5159     case MCK_AV_32: return true;
10043   case MCK_AV_32: return "MCK_AV_32";
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 }, },
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 }, },
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 }, },
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 }, },
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 }, },
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 }, },
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 }, },
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 }, },
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 }, },
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 }, },
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 }, },
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 }, },
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 }, },
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 }, },
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 }, },
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 }, },
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 }, },
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 }, },
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 }, },