|
| 1 | +# RUN: llvm-mca -mtriple=amdgcn -mcpu=gfx940 --timeline --iterations=1 --timeline-max-cycles=0 < %s | FileCheck %s |
| 2 | + |
| 3 | +# CHECK: Iterations: 1 |
| 4 | +# CHECK: Instructions: 71 |
| 5 | +# CHECK: Total Cycles: 562 |
| 6 | +# CHECK: Total uOps: 77 |
| 7 | + |
| 8 | +# CHECK: Resources: |
| 9 | +# CHECK: [0] - HWBranch |
| 10 | +# CHECK: [1] - HWExport |
| 11 | +# CHECK: [2] - HWLGKM |
| 12 | +# CHECK: [3] - HWSALU |
| 13 | +# CHECK: [4] - HWVALU |
| 14 | +# CHECK: [5] - HWVMEM |
| 15 | +# CHECK: [6] - HWXDL |
| 16 | + |
| 17 | +v_pk_fma_f32 v[0:1], v[0:1], v[0:1], v[0:1] |
| 18 | +v_pk_mov_b32 v[0:1], v[2:3], v[4:5] |
| 19 | +v_pk_add_f32 v[0:1], v[0:1], v[0:1] |
| 20 | +v_pk_mul_f32 v[0:1], v[0:1], v[0:1] |
| 21 | +v_add_co_u32 v5, s[0:1], v1, v2 |
| 22 | +v_sub_co_u32 v5, s[0:1], v1, v2 |
| 23 | +v_subrev_co_u32 v5, s[0:1], v1, v2 |
| 24 | +v_addc_co_u32 v5, s[0:1], v1, v2, s[2:3] |
| 25 | +v_subb_co_u32 v5, s[0:1], v1, v2, s[2:3] |
| 26 | +v_subbrev_co_u32 v5, s[0:1], v1, v2, s[2:3] |
| 27 | +v_add_u32 v5, v1, v2 |
| 28 | +v_sub_u32 v5, v1, v2 |
| 29 | +v_subrev_u32 v5, v1, v2 |
| 30 | + |
| 31 | +v_mfma_f32_16x16x4_f32 a[0:3], v0, v1, a[2:5] |
| 32 | +v_mfma_f32_16x16x4_f32 v[0:3], v0, v1, v[2:5] |
| 33 | + |
| 34 | +v_mfma_f32_32x32x2_f32 a[0:15], v0, v1, a[18:33] |
| 35 | +v_mfma_f32_32x32x2_f32 v[0:15], v0, v1, v[18:33] |
| 36 | + |
| 37 | +v_mfma_f64_4x4x4_4b_f64 a[0:1], v[0:1], a[2:3], a[2:3] |
| 38 | +v_mfma_f64_4x4x4_4b_f64 v[0:1], v[0:1], v[2:3], v[2:3] |
| 39 | + |
| 40 | +v_mfma_f64_16x16x4_f64 a[0:7], v[0:1], v[2:3], a[0:7] |
| 41 | +v_mfma_f64_16x16x4_f64 v[0:7], v[0:1], v[2:3], v[0:7] |
| 42 | + |
| 43 | +v_mfma_f32_16x16x16_f16 v[0:3], v[4:5], v[6:7], v[0:3] |
| 44 | +v_mfma_f32_16x16x16_f16 a[0:3], v[4:5], v[6:7], a[0:3] |
| 45 | + |
| 46 | +v_mfma_f32_32x32x8_f16 v[0:15], v[4:5], v[6:7], v[0:15] |
| 47 | +v_mfma_f32_32x32x8_f16 a[0:15], v[4:5], v[6:7], a[0:15] |
| 48 | + |
| 49 | +v_mfma_f32_16x16x16_bf16 v[0:3], v[4:5], v[6:7], v[0:3] |
| 50 | +v_mfma_f32_16x16x16_bf16 a[0:3], v[4:5], v[6:7], a[0:3] |
| 51 | + |
| 52 | +v_mfma_f32_32x32x8_bf16 v[0:15], v[4:5], v[6:7], v[0:15] |
| 53 | +v_mfma_f32_32x32x8_bf16 a[0:15], v[4:5], v[6:7], a[0:15] |
| 54 | + |
| 55 | +v_mfma_i32_16x16x32_i8 v[0:3], v[4:5], v[6:7], v[0:3] |
| 56 | +v_mfma_i32_16x16x32_i8 a[0:3], v[4:5], v[6:7], a[0:3] |
| 57 | + |
| 58 | +v_mfma_i32_32x32x16_i8 v[0:15], v[2:3], v[4:5], v[0:15] |
| 59 | +v_mfma_i32_32x32x16_i8 a[0:15], v[2:3], v[4:5], a[0:15] |
| 60 | + |
| 61 | +v_mfma_f32_4x4x4_16b_f16 v[0:3], v[0:1], v[2:3], v[2:5] |
| 62 | +v_mfma_f32_4x4x4_16b_f16 a[0:3], v[0:1], v[2:3], a[2:5] |
| 63 | + |
| 64 | +v_mfma_f32_16x16x4_4b_f16 v[0:15], v[2:3], v[4:5], v[18:33] |
| 65 | +v_mfma_f32_16x16x4_4b_f16 a[0:15], v[2:3], v[4:5], a[18:33] |
| 66 | + |
| 67 | +v_mfma_f32_32x32x4_2b_f16 v[0:31], v[0:1], v[2:3], v[34:65] |
| 68 | +v_mfma_f32_32x32x4_2b_f16 a[0:31], v[0:1], v[2:3], a[34:65] |
| 69 | + |
| 70 | +v_mfma_f32_4x4x4_16b_bf16 v[0:3], v[0:1], v[2:3], v[2:5] |
| 71 | +v_mfma_f32_4x4x4_16b_bf16 a[0:3], v[0:1], v[2:3], a[2:5] |
| 72 | + |
| 73 | +v_mfma_f32_16x16x4_4b_bf16 v[0:15], v[2:3], v[4:5], v[18:33] |
| 74 | +v_mfma_f32_16x16x4_4b_bf16 a[0:15], v[2:3], v[4:5], a[18:33] |
| 75 | + |
| 76 | +v_mfma_f32_32x32x4_2b_bf16 v[0:31], v[0:1], v[2:3], v[34:65] |
| 77 | +v_mfma_f32_32x32x4_2b_bf16 a[0:31], v[0:1], v[2:3], a[34:65] |
| 78 | + |
| 79 | +v_mfma_f32_4x4x1_16b_f32 v[0:3], v0, v1, v[2:5] |
| 80 | +v_mfma_f32_4x4x1_16b_f32 a[0:3], v0, v1, a[2:5] |
| 81 | + |
| 82 | +v_mfma_f32_16x16x1_4b_f32 v[0:15], v0, v1, v[18:33] |
| 83 | +v_mfma_f32_16x16x1_4b_f32 a[0:15], v0, v1, a[18:33] |
| 84 | + |
| 85 | +v_mfma_f32_16x16x4_f32 v[0:3], v0, v1, v[2:5] |
| 86 | +v_mfma_f32_16x16x4_f32 a[0:3], v0, v1, a[2:5] |
| 87 | + |
| 88 | +v_mfma_f32_32x32x1_2b_f32 v[0:31], v0, v1, v[34:65] blgp:7 |
| 89 | +v_mfma_f32_32x32x1_2b_f32 a[0:31], v0, v1, a[34:65] blgp:7 |
| 90 | + |
| 91 | +v_mfma_f32_32x32x2_f32 v[0:15], v0, v1, v[18:33] |
| 92 | +v_mfma_f32_32x32x2_f32 a[0:15], v0, v1, a[18:33] |
| 93 | + |
| 94 | +v_mfma_i32_4x4x4_16b_i8 v[0:3], v0, v1, v[2:5] |
| 95 | +v_mfma_i32_4x4x4_16b_i8 a[0:3], v0, v1, a[2:5] |
| 96 | + |
| 97 | +v_mfma_i32_16x16x4_4b_i8 v[0:15], v0, v1, v[18:33] |
| 98 | +v_mfma_i32_16x16x4_4b_i8 a[0:15], v0, v1, a[18:33] |
| 99 | + |
| 100 | +v_mfma_i32_32x32x4_2b_i8 v[0:31], v0, v1, v[34:65] |
| 101 | +v_mfma_i32_32x32x4_2b_i8 a[0:31], v0, v1, a[34:65] |
| 102 | + |
| 103 | +v_smfmac_f32_16x16x32_f16 v[10:13], a[2:3], v[4:7], v0 cbsz:3 abid:1 |
| 104 | +v_smfmac_f32_16x16x32_f16 a[10:13], v[2:3], a[4:7], v1 |
| 105 | + |
| 106 | +v_smfmac_f32_32x32x16_f16 v[10:25], a[2:3], v[4:7], v2 cbsz:3 abid:1 |
| 107 | +v_smfmac_f32_32x32x16_f16 a[10:25], v[2:3], a[4:7], v3 |
| 108 | + |
| 109 | +v_smfmac_f32_16x16x32_bf16 v[10:13], a[2:3], v[4:7], v4 cbsz:3 abid:1 |
| 110 | +v_smfmac_f32_16x16x32_bf16 a[10:13], v[2:3], a[4:7], v5 |
| 111 | + |
| 112 | +v_smfmac_i32_16x16x64_i8 v[10:13], a[2:3], v[4:7], v8 cbsz:3 abid:1 |
| 113 | +v_smfmac_i32_16x16x64_i8 a[10:13], v[2:3], a[4:7], v9 |
| 114 | + |
| 115 | +v_smfmac_i32_32x32x32_i8 v[10:25], a[2:3], v[4:7], v10 cbsz:3 abid:1 |
| 116 | +v_smfmac_i32_32x32x32_i8 a[10:25], v[2:3], a[4:7], v11 |
| 117 | + |
| 118 | +# CHECK: [0] [1] [2] [3] [4] [5] [6] Instructions: |
| 119 | +# CHECK-NEXT: - - - - 1.00 - - v_pk_fma_f32 v[0:1], v[0:1], v[0:1], v[0:1] |
| 120 | +# CHECK-NEXT: - - - - 1.00 - - v_pk_mov_b32 v[0:1], v[2:3], v[4:5] |
| 121 | +# CHECK-NEXT: - - - - 1.00 - - v_pk_add_f32 v[0:1], v[0:1], v[0:1] |
| 122 | +# CHECK-NEXT: - - - - 1.00 - - v_pk_mul_f32 v[0:1], v[0:1], v[0:1] |
| 123 | +# CHECK-NEXT: - - - 1.00 1.00 - - v_add_co_u32_e64 v5, s[0:1], v1, v2 |
| 124 | +# CHECK-NEXT: - - - 1.00 1.00 - - v_sub_co_u32_e64 v5, s[0:1], v1, v2 |
| 125 | +# CHECK-NEXT: - - - 1.00 1.00 - - v_subrev_co_u32_e64 v5, s[0:1], v1, v2 |
| 126 | +# CHECK-NEXT: - - - 1.00 1.00 - - v_addc_co_u32_e64 v5, s[0:1], v1, v2, s[2:3] |
| 127 | +# CHECK-NEXT: - - - 1.00 1.00 - - v_subb_co_u32_e64 v5, s[0:1], v1, v2, s[2:3] |
| 128 | +# CHECK-NEXT: - - - 1.00 1.00 - - v_subbrev_co_u32_e64 v5, s[0:1], v1, v2, s[2:3] |
| 129 | +# CHECK-NEXT: - - - - 1.00 - - v_add_u32_e32 v5, v1, v2 |
| 130 | +# CHECK-NEXT: - - - - 1.00 - - v_sub_u32_e32 v5, v1, v2 |
| 131 | +# CHECK-NEXT: - - - - 1.00 - - v_subrev_u32_e32 v5, v1, v2 |
| 132 | +# CHECK-NEXT: - - - - - - 8.00 v_mfma_f32_16x16x4_f32 a[0:3], v0, v1, a[2:5] |
| 133 | +# CHECK-NEXT: - - - - - - 8.00 v_mfma_f32_16x16x4_f32 v[0:3], v0, v1, v[2:5] |
| 134 | +# CHECK-NEXT: - - - - - - 16.00 v_mfma_f32_32x32x2_f32 a[0:15], v0, v1, a[18:33] |
| 135 | +# CHECK-NEXT: - - - - - - 16.00 v_mfma_f32_32x32x2_f32 v[0:15], v0, v1, v[18:33] |
| 136 | +# CHECK-NEXT: - - - - 1.00 - - v_mfma_f64_4x4x4_4b_f64 a[0:1], v[0:1], a[2:3], a[2:3] |
| 137 | +# CHECK-NEXT: - - - - 1.00 - - v_mfma_f64_4x4x4_4b_f64 v[0:1], v[0:1], v[2:3], v[2:3] |
| 138 | +# CHECK-NEXT: - - - - 1.00 - - v_mfma_f64_16x16x4_f64 a[0:7], v[0:1], v[2:3], a[0:7] |
| 139 | +# CHECK-NEXT: - - - - 1.00 - - v_mfma_f64_16x16x4_f64 v[0:7], v[0:1], v[2:3], v[0:7] |
| 140 | +# CHECK-NEXT: - - - - - - 4.00 v_mfma_f32_16x16x16_f16 v[0:3], v[4:5], v[6:7], v[0:3] |
| 141 | +# CHECK-NEXT: - - - - - - 4.00 v_mfma_f32_16x16x16_f16 a[0:3], v[4:5], v[6:7], a[0:3] |
| 142 | +# CHECK-NEXT: - - - - - - 8.00 v_mfma_f32_32x32x8_f16 v[0:15], v[4:5], v[6:7], v[0:15] |
| 143 | +# CHECK-NEXT: - - - - - - 8.00 v_mfma_f32_32x32x8_f16 a[0:15], v[4:5], v[6:7], a[0:15] |
| 144 | +# CHECK-NEXT: - - - - - - 4.00 v_mfma_f32_16x16x16_bf16 v[0:3], v[4:5], v[6:7], v[0:3] |
| 145 | +# CHECK-NEXT: - - - - - - 4.00 v_mfma_f32_16x16x16_bf16 a[0:3], v[4:5], v[6:7], a[0:3] |
| 146 | +# CHECK-NEXT: - - - - - - 8.00 v_mfma_f32_32x32x8_bf16 v[0:15], v[4:5], v[6:7], v[0:15] |
| 147 | +# CHECK-NEXT: - - - - - - 8.00 v_mfma_f32_32x32x8_bf16 a[0:15], v[4:5], v[6:7], a[0:15] |
| 148 | +# CHECK-NEXT: - - - - - - 4.00 v_mfma_i32_16x16x32_i8 v[0:3], v[4:5], v[6:7], v[0:3] |
| 149 | +# CHECK-NEXT: - - - - - - 4.00 v_mfma_i32_16x16x32_i8 a[0:3], v[4:5], v[6:7], a[0:3] |
| 150 | +# CHECK-NEXT: - - - - - - 8.00 v_mfma_i32_32x32x16_i8 v[0:15], v[2:3], v[4:5], v[0:15] |
| 151 | +# CHECK-NEXT: - - - - - - 8.00 v_mfma_i32_32x32x16_i8 a[0:15], v[2:3], v[4:5], a[0:15] |
| 152 | +# CHECK-NEXT: - - - - - - 2.00 v_mfma_f32_4x4x4_16b_f16 v[0:3], v[0:1], v[2:3], v[2:5] |
| 153 | +# CHECK-NEXT: - - - - - - 2.00 v_mfma_f32_4x4x4_16b_f16 a[0:3], v[0:1], v[2:3], a[2:5] |
| 154 | +# CHECK-NEXT: - - - - - - 8.00 v_mfma_f32_16x16x4_4b_f16 v[0:15], v[2:3], v[4:5], v[18:33] |
| 155 | +# CHECK-NEXT: - - - - - - 8.00 v_mfma_f32_16x16x4_4b_f16 a[0:15], v[2:3], v[4:5], a[18:33] |
| 156 | +# CHECK-NEXT: - - - - - - 16.00 v_mfma_f32_32x32x4_2b_f16 v[0:31], v[0:1], v[2:3], v[34:65] |
| 157 | +# CHECK-NEXT: - - - - - - 16.00 v_mfma_f32_32x32x4_2b_f16 a[0:31], v[0:1], v[2:3], a[34:65] |
| 158 | +# CHECK-NEXT: - - - - - - 2.00 v_mfma_f32_4x4x4_16b_bf16 v[0:3], v[0:1], v[2:3], v[2:5] |
| 159 | +# CHECK-NEXT: - - - - - - 2.00 v_mfma_f32_4x4x4_16b_bf16 a[0:3], v[0:1], v[2:3], a[2:5] |
| 160 | +# CHECK-NEXT: - - - - - - 8.00 v_mfma_f32_16x16x4_4b_bf16 v[0:15], v[2:3], v[4:5], v[18:33] |
| 161 | +# CHECK-NEXT: - - - - - - 8.00 v_mfma_f32_16x16x4_4b_bf16 a[0:15], v[2:3], v[4:5], a[18:33] |
| 162 | +# CHECK-NEXT: - - - - - - 16.00 v_mfma_f32_32x32x4_2b_bf16 v[0:31], v[0:1], v[2:3], v[34:65] |
| 163 | +# CHECK-NEXT: - - - - - - 16.00 v_mfma_f32_32x32x4_2b_bf16 a[0:31], v[0:1], v[2:3], a[34:65] |
| 164 | +# CHECK-NEXT: - - - - - - 2.00 v_mfma_f32_4x4x1_16b_f32 v[0:3], v0, v1, v[2:5] |
| 165 | +# CHECK-NEXT: - - - - - - 2.00 v_mfma_f32_4x4x1_16b_f32 a[0:3], v0, v1, a[2:5] |
| 166 | +# CHECK-NEXT: - - - - - - 8.00 v_mfma_f32_16x16x1_4b_f32 v[0:15], v0, v1, v[18:33] |
| 167 | +# CHECK-NEXT: - - - - - - 8.00 v_mfma_f32_16x16x1_4b_f32 a[0:15], v0, v1, a[18:33] |
| 168 | +# CHECK-NEXT: - - - - - - 8.00 v_mfma_f32_16x16x4_f32 v[0:3], v0, v1, v[2:5] |
| 169 | +# CHECK-NEXT: - - - - - - 8.00 v_mfma_f32_16x16x4_f32 a[0:3], v0, v1, a[2:5] |
| 170 | +# CHECK-NEXT: - - - - - - 16.00 v_mfma_f32_32x32x1_2b_f32 v[0:31], v0, v1, v[34:65] blgp:7 |
| 171 | +# CHECK-NEXT: - - - - - - 16.00 v_mfma_f32_32x32x1_2b_f32 a[0:31], v0, v1, a[34:65] blgp:7 |
| 172 | +# CHECK-NEXT: - - - - - - 16.00 v_mfma_f32_32x32x2_f32 v[0:15], v0, v1, v[18:33] |
| 173 | +# CHECK-NEXT: - - - - - - 16.00 v_mfma_f32_32x32x2_f32 a[0:15], v0, v1, a[18:33] |
| 174 | +# CHECK-NEXT: - - - - - - 2.00 v_mfma_i32_4x4x4_16b_i8 v[0:3], v0, v1, v[2:5] |
| 175 | +# CHECK-NEXT: - - - - - - 2.00 v_mfma_i32_4x4x4_16b_i8 a[0:3], v0, v1, a[2:5] |
| 176 | +# CHECK-NEXT: - - - - - - 8.00 v_mfma_i32_16x16x4_4b_i8 v[0:15], v0, v1, v[18:33] |
| 177 | +# CHECK-NEXT: - - - - - - 8.00 v_mfma_i32_16x16x4_4b_i8 a[0:15], v0, v1, a[18:33] |
| 178 | +# CHECK-NEXT: - - - - - - 16.00 v_mfma_i32_32x32x4_2b_i8 v[0:31], v0, v1, v[34:65] |
| 179 | +# CHECK-NEXT: - - - - - - 16.00 v_mfma_i32_32x32x4_2b_i8 a[0:31], v0, v1, a[34:65] |
| 180 | +# CHECK-NEXT: - - - - - - 4.00 v_smfmac_f32_16x16x32_f16 v[10:13], a[2:3], v[4:7], v0 cbsz:3 abid:1 |
| 181 | +# CHECK-NEXT: - - - - - - 4.00 v_smfmac_f32_16x16x32_f16 a[10:13], v[2:3], a[4:7], v1 |
| 182 | +# CHECK-NEXT: - - - - - - 8.00 v_smfmac_f32_32x32x16_f16 v[10:25], a[2:3], v[4:7], v2 cbsz:3 abid:1 |
| 183 | +# CHECK-NEXT: - - - - - - 8.00 v_smfmac_f32_32x32x16_f16 a[10:25], v[2:3], a[4:7], v3 |
| 184 | +# CHECK-NEXT: - - - - - - 4.00 v_smfmac_f32_16x16x32_bf16 v[10:13], a[2:3], v[4:7], v4 cbsz:3 abid:1 |
| 185 | +# CHECK-NEXT: - - - - - - 4.00 v_smfmac_f32_16x16x32_bf16 a[10:13], v[2:3], a[4:7], v5 |
| 186 | +# CHECK-NEXT: - - - - - - 4.00 v_smfmac_i32_16x16x64_i8 v[10:13], a[2:3], v[4:7], v8 cbsz:3 abid:1 |
| 187 | +# CHECK-NEXT: - - - - - - 4.00 v_smfmac_i32_16x16x64_i8 a[10:13], v[2:3], a[4:7], v9 |
| 188 | +# CHECK-NEXT: - - - - - - 8.00 v_smfmac_i32_32x32x32_i8 v[10:25], a[2:3], v[4:7], v10 cbsz:3 abid:1 |
| 189 | +# CHECK-NEXT: - - - - - - 8.00 v_smfmac_i32_32x32x32_i8 a[10:25], v[2:3], a[4:7], v11 |
0 commit comments