llvm/test/MC/AMDGPU/sym_kernel_scope_agpr.s

   1 // RUN: llvm-mc -triple=amdgcn -mcpu=gfx90a %s 2>&1 | FileCheck -check-prefixes=GFX90A %s
   2 // RUN: llvm-mc -triple=amdgcn -mcpu=gfx908 %s 2>&1 | FileCheck -check-prefixes=GFX908 %s
   3 // Based on sym_kernel_scope.s
   4
   5 .byte .kernel.agpr_count
   6 // CHECK: .byte 0
   7 .byte .kernel.vgpr_count
   8 // CHECK: .byte 0
   9
  10     v_accvgpr_write_b32 a0, v6
  11     v_accvgpr_read_b32 v3, a3
  12     s_endpgm
  13 .byte .kernel.agpr_count
  14 // GFX90A: .byte 4
  15 // GFX908: .byte 4
  16 .byte .kernel.vgpr_count
  17 // GFX90A: .byte 12
  18 // GFX908: .byte 7
  19
  20 .amdgpu_hsa_kernel K1
  21 K1:
  22 .byte .kernel.agpr_count
  23 // CHECK: .byte 0
  24 .byte .kernel.vgpr_count
  25 // CHECK: .byte 0
  26     v_accvgpr_write_b32 a44, v6
  27     s_endpgm
  28 .byte .kernel.agpr_count
  29 // GFX90A: .byte 45
  30 // GFX908: .byte 45
  31 .byte .kernel.vgpr_count
  32 // GFX90A: .byte 53
  33 // GFX908: .byte 45
  34
  35 .amdgpu_hsa_kernel K2
  36 .byte .kernel.agpr_count
  37 // CHECK: .byte 0
  38 .byte .kernel.vgpr_count
  39 // CHECK: .byte 0
  40 K2:
  41     v_mfma_f32_4x4x1f32 a[0:3], v1, v0, a[0:3] cbsz:1 abid:2 blgp:3
  42     s_endpgm
  43 .byte .kernel.agpr_count
  44 // GFX90A: .byte 4
  45 // GFX908: .byte 4
  46 .byte .kernel.vgpr_count
  47 // GFX90A: .byte 8
  48 // GFX908: .byte 4
  49
  50 .text
  51 .amdgpu_hsa_kernel K3
  52 K3:
  53     v_accvgpr_read_b32 v[0], a0
  54     v_mfma_f32_16x16x1f32 a[0:15], v1, v0, a[0:15] cbsz:1 abid:2 blgp:3
  55     s_endpgm
  56
  57 .byte .kernel.agpr_count
  58 // GFX90A: .byte 16
  59 // GFX908: .byte 16
  60 .byte .kernel.vgpr_count
  61 // GFX90A: .byte 20
  62 // GFX908: .byte 16