63 lines · plain
1// RUN: llvm-mc -triple=amdgcn -mcpu=gfx90a %s 2>&1 | FileCheck -check-prefixes=GFX90A %s2// RUN: llvm-mc -triple=amdgcn -mcpu=gfx908 %s 2>&1 | FileCheck -check-prefixes=GFX908 %s3// Based on sym_kernel_scope.s4 5.byte .kernel.agpr_count6// CHECK: .byte 07.byte .kernel.vgpr_count8// CHECK: .byte 09 10 v_accvgpr_write_b32 a0, v611 v_accvgpr_read_b32 v3, a312 s_endpgm13.byte .kernel.agpr_count14// GFX90A: .byte 415// GFX908: .byte 416.byte .kernel.vgpr_count17// GFX90A: .byte 1218// GFX908: .byte 719 20.amdgpu_hsa_kernel K121K1:22.byte .kernel.agpr_count23// CHECK: .byte 024.byte .kernel.vgpr_count25// CHECK: .byte 026 v_accvgpr_write_b32 a44, v627 s_endpgm28.byte .kernel.agpr_count29// GFX90A: .byte 4530// GFX908: .byte 4531.byte .kernel.vgpr_count32// GFX90A: .byte 5333// GFX908: .byte 4534 35.amdgpu_hsa_kernel K236.byte .kernel.agpr_count37// CHECK: .byte 038.byte .kernel.vgpr_count39// CHECK: .byte 040K2:41 v_mfma_f32_4x4x1f32 a[0:3], v1, v0, a[0:3] cbsz:1 abid:2 blgp:342 s_endpgm43.byte .kernel.agpr_count44// GFX90A: .byte 445// GFX908: .byte 446.byte .kernel.vgpr_count47// GFX90A: .byte 848// GFX908: .byte 449 50.text51.amdgpu_hsa_kernel K352K3:53 v_accvgpr_read_b32 v[0], a054 v_mfma_f32_16x16x1f32 a[0:15], v1, v0, a[0:15] cbsz:1 abid:2 blgp:355 s_endpgm56 57.byte .kernel.agpr_count58// GFX90A: .byte 1659// GFX908: .byte 1660.byte .kernel.vgpr_count61// GFX90A: .byte 2062// GFX908: .byte 1663