Skip to content

Commit

Permalink
AMDGPU: Add scheduling test for gfx940 (#83220)
Browse files Browse the repository at this point in the history 
I'm not sure the f64 fma cases are correct.
  • Loading branch information
@arsenm
arsenm committed Feb 29, 2024
1 parent 7d8b50a commit f0484e0
Showing 1 changed file with 189 additions and 0 deletions.
189 changes: 189 additions & 0 deletions llvm/test/tools/llvm-mca/AMDGPU/gfx940.s
View file
Original file line number Diff line number Diff line change
@@ -0,0 +1,189 @@
# RUN: llvm-mca -mtriple=amdgcn -mcpu=gfx940 --timeline --iterations=1 --timeline-max-cycles=0 < %s | FileCheck %s

# CHECK: Iterations: 1
# CHECK: Instructions: 71
# CHECK: Total Cycles: 562
# CHECK: Total uOps: 77

# CHECK: Resources:
# CHECK: [0] - HWBranch
# CHECK: [1] - HWExport
# CHECK: [2] - HWLGKM
# CHECK: [3] - HWSALU
# CHECK: [4] - HWVALU
# CHECK: [5] - HWVMEM
# CHECK: [6] - HWXDL

v_pk_fma_f32 v[0:1], v[0:1], v[0:1], v[0:1]
v_pk_mov_b32 v[0:1], v[2:3], v[4:5]
v_pk_add_f32 v[0:1], v[0:1], v[0:1]
v_pk_mul_f32 v[0:1], v[0:1], v[0:1]
v_add_co_u32 v5, s[0:1], v1, v2
v_sub_co_u32 v5, s[0:1], v1, v2
v_subrev_co_u32 v5, s[0:1], v1, v2
v_addc_co_u32 v5, s[0:1], v1, v2, s[2:3]
v_subb_co_u32 v5, s[0:1], v1, v2, s[2:3]
v_subbrev_co_u32 v5, s[0:1], v1, v2, s[2:3]
v_add_u32 v5, v1, v2
v_sub_u32 v5, v1, v2
v_subrev_u32 v5, v1, v2

v_mfma_f32_16x16x4_f32 a[0:3], v0, v1, a[2:5]
v_mfma_f32_16x16x4_f32 v[0:3], v0, v1, v[2:5]

v_mfma_f32_32x32x2_f32 a[0:15], v0, v1, a[18:33]
v_mfma_f32_32x32x2_f32 v[0:15], v0, v1, v[18:33]

v_mfma_f64_4x4x4_4b_f64 a[0:1], v[0:1], a[2:3], a[2:3]
v_mfma_f64_4x4x4_4b_f64 v[0:1], v[0:1], v[2:3], v[2:3]

v_mfma_f64_16x16x4_f64 a[0:7], v[0:1], v[2:3], a[0:7]
v_mfma_f64_16x16x4_f64 v[0:7], v[0:1], v[2:3], v[0:7]

v_mfma_f32_16x16x16_f16 v[0:3], v[4:5], v[6:7], v[0:3]
v_mfma_f32_16x16x16_f16 a[0:3], v[4:5], v[6:7], a[0:3]

v_mfma_f32_32x32x8_f16 v[0:15], v[4:5], v[6:7], v[0:15]
v_mfma_f32_32x32x8_f16 a[0:15], v[4:5], v[6:7], a[0:15]

v_mfma_f32_16x16x16_bf16 v[0:3], v[4:5], v[6:7], v[0:3]
v_mfma_f32_16x16x16_bf16 a[0:3], v[4:5], v[6:7], a[0:3]

v_mfma_f32_32x32x8_bf16 v[0:15], v[4:5], v[6:7], v[0:15]
v_mfma_f32_32x32x8_bf16 a[0:15], v[4:5], v[6:7], a[0:15]

v_mfma_i32_16x16x32_i8 v[0:3], v[4:5], v[6:7], v[0:3]
v_mfma_i32_16x16x32_i8 a[0:3], v[4:5], v[6:7], a[0:3]

v_mfma_i32_32x32x16_i8 v[0:15], v[2:3], v[4:5], v[0:15]
v_mfma_i32_32x32x16_i8 a[0:15], v[2:3], v[4:5], a[0:15]

v_mfma_f32_4x4x4_16b_f16 v[0:3], v[0:1], v[2:3], v[2:5]
v_mfma_f32_4x4x4_16b_f16 a[0:3], v[0:1], v[2:3], a[2:5]

v_mfma_f32_16x16x4_4b_f16 v[0:15], v[2:3], v[4:5], v[18:33]
v_mfma_f32_16x16x4_4b_f16 a[0:15], v[2:3], v[4:5], a[18:33]

v_mfma_f32_32x32x4_2b_f16 v[0:31], v[0:1], v[2:3], v[34:65]
v_mfma_f32_32x32x4_2b_f16 a[0:31], v[0:1], v[2:3], a[34:65]

v_mfma_f32_4x4x4_16b_bf16 v[0:3], v[0:1], v[2:3], v[2:5]
v_mfma_f32_4x4x4_16b_bf16 a[0:3], v[0:1], v[2:3], a[2:5]

v_mfma_f32_16x16x4_4b_bf16 v[0:15], v[2:3], v[4:5], v[18:33]
v_mfma_f32_16x16x4_4b_bf16 a[0:15], v[2:3], v[4:5], a[18:33]

v_mfma_f32_32x32x4_2b_bf16 v[0:31], v[0:1], v[2:3], v[34:65]
v_mfma_f32_32x32x4_2b_bf16 a[0:31], v[0:1], v[2:3], a[34:65]

v_mfma_f32_4x4x1_16b_f32 v[0:3], v0, v1, v[2:5]
v_mfma_f32_4x4x1_16b_f32 a[0:3], v0, v1, a[2:5]

v_mfma_f32_16x16x1_4b_f32 v[0:15], v0, v1, v[18:33]
v_mfma_f32_16x16x1_4b_f32 a[0:15], v0, v1, a[18:33]

v_mfma_f32_16x16x4_f32 v[0:3], v0, v1, v[2:5]
v_mfma_f32_16x16x4_f32 a[0:3], v0, v1, a[2:5]

v_mfma_f32_32x32x1_2b_f32 v[0:31], v0, v1, v[34:65] blgp:7
v_mfma_f32_32x32x1_2b_f32 a[0:31], v0, v1, a[34:65] blgp:7

v_mfma_f32_32x32x2_f32 v[0:15], v0, v1, v[18:33]
v_mfma_f32_32x32x2_f32 a[0:15], v0, v1, a[18:33]

v_mfma_i32_4x4x4_16b_i8 v[0:3], v0, v1, v[2:5]
v_mfma_i32_4x4x4_16b_i8 a[0:3], v0, v1, a[2:5]

v_mfma_i32_16x16x4_4b_i8 v[0:15], v0, v1, v[18:33]
v_mfma_i32_16x16x4_4b_i8 a[0:15], v0, v1, a[18:33]

v_mfma_i32_32x32x4_2b_i8 v[0:31], v0, v1, v[34:65]
v_mfma_i32_32x32x4_2b_i8 a[0:31], v0, v1, a[34:65]

v_smfmac_f32_16x16x32_f16 v[10:13], a[2:3], v[4:7], v0 cbsz:3 abid:1
v_smfmac_f32_16x16x32_f16 a[10:13], v[2:3], a[4:7], v1

v_smfmac_f32_32x32x16_f16 v[10:25], a[2:3], v[4:7], v2 cbsz:3 abid:1
v_smfmac_f32_32x32x16_f16 a[10:25], v[2:3], a[4:7], v3

v_smfmac_f32_16x16x32_bf16 v[10:13], a[2:3], v[4:7], v4 cbsz:3 abid:1
v_smfmac_f32_16x16x32_bf16 a[10:13], v[2:3], a[4:7], v5

v_smfmac_i32_16x16x64_i8 v[10:13], a[2:3], v[4:7], v8 cbsz:3 abid:1
v_smfmac_i32_16x16x64_i8 a[10:13], v[2:3], a[4:7], v9

v_smfmac_i32_32x32x32_i8 v[10:25], a[2:3], v[4:7], v10 cbsz:3 abid:1
v_smfmac_i32_32x32x32_i8 a[10:25], v[2:3], a[4:7], v11

# CHECK: [0] [1] [2] [3] [4] [5] [6] Instructions:
# CHECK-NEXT: - - - - 1.00 - - v_pk_fma_f32 v[0:1], v[0:1], v[0:1], v[0:1]
# CHECK-NEXT: - - - - 1.00 - - v_pk_mov_b32 v[0:1], v[2:3], v[4:5]
# CHECK-NEXT: - - - - 1.00 - - v_pk_add_f32 v[0:1], v[0:1], v[0:1]
# CHECK-NEXT: - - - - 1.00 - - v_pk_mul_f32 v[0:1], v[0:1], v[0:1]
# CHECK-NEXT: - - - 1.00 1.00 - - v_add_co_u32_e64 v5, s[0:1], v1, v2
# CHECK-NEXT: - - - 1.00 1.00 - - v_sub_co_u32_e64 v5, s[0:1], v1, v2
# CHECK-NEXT: - - - 1.00 1.00 - - v_subrev_co_u32_e64 v5, s[0:1], v1, v2
# CHECK-NEXT: - - - 1.00 1.00 - - v_addc_co_u32_e64 v5, s[0:1], v1, v2, s[2:3]
# CHECK-NEXT: - - - 1.00 1.00 - - v_subb_co_u32_e64 v5, s[0:1], v1, v2, s[2:3]
# CHECK-NEXT: - - - 1.00 1.00 - - v_subbrev_co_u32_e64 v5, s[0:1], v1, v2, s[2:3]
# CHECK-NEXT: - - - - 1.00 - - v_add_u32_e32 v5, v1, v2
# CHECK-NEXT: - - - - 1.00 - - v_sub_u32_e32 v5, v1, v2
# CHECK-NEXT: - - - - 1.00 - - v_subrev_u32_e32 v5, v1, v2
# CHECK-NEXT: - - - - - - 8.00 v_mfma_f32_16x16x4_f32 a[0:3], v0, v1, a[2:5]
# CHECK-NEXT: - - - - - - 8.00 v_mfma_f32_16x16x4_f32 v[0:3], v0, v1, v[2:5]
# CHECK-NEXT: - - - - - - 16.00 v_mfma_f32_32x32x2_f32 a[0:15], v0, v1, a[18:33]
# CHECK-NEXT: - - - - - - 16.00 v_mfma_f32_32x32x2_f32 v[0:15], v0, v1, v[18:33]
# CHECK-NEXT: - - - - 1.00 - - v_mfma_f64_4x4x4_4b_f64 a[0:1], v[0:1], a[2:3], a[2:3]
# CHECK-NEXT: - - - - 1.00 - - v_mfma_f64_4x4x4_4b_f64 v[0:1], v[0:1], v[2:3], v[2:3]
# CHECK-NEXT: - - - - 1.00 - - v_mfma_f64_16x16x4_f64 a[0:7], v[0:1], v[2:3], a[0:7]
# CHECK-NEXT: - - - - 1.00 - - v_mfma_f64_16x16x4_f64 v[0:7], v[0:1], v[2:3], v[0:7]
# CHECK-NEXT: - - - - - - 4.00 v_mfma_f32_16x16x16_f16 v[0:3], v[4:5], v[6:7], v[0:3]
# CHECK-NEXT: - - - - - - 4.00 v_mfma_f32_16x16x16_f16 a[0:3], v[4:5], v[6:7], a[0:3]
# CHECK-NEXT: - - - - - - 8.00 v_mfma_f32_32x32x8_f16 v[0:15], v[4:5], v[6:7], v[0:15]
# CHECK-NEXT: - - - - - - 8.00 v_mfma_f32_32x32x8_f16 a[0:15], v[4:5], v[6:7], a[0:15]
# CHECK-NEXT: - - - - - - 4.00 v_mfma_f32_16x16x16_bf16 v[0:3], v[4:5], v[6:7], v[0:3]
# CHECK-NEXT: - - - - - - 4.00 v_mfma_f32_16x16x16_bf16 a[0:3], v[4:5], v[6:7], a[0:3]
# CHECK-NEXT: - - - - - - 8.00 v_mfma_f32_32x32x8_bf16 v[0:15], v[4:5], v[6:7], v[0:15]
# CHECK-NEXT: - - - - - - 8.00 v_mfma_f32_32x32x8_bf16 a[0:15], v[4:5], v[6:7], a[0:15]
# CHECK-NEXT: - - - - - - 4.00 v_mfma_i32_16x16x32_i8 v[0:3], v[4:5], v[6:7], v[0:3]
# CHECK-NEXT: - - - - - - 4.00 v_mfma_i32_16x16x32_i8 a[0:3], v[4:5], v[6:7], a[0:3]
# CHECK-NEXT: - - - - - - 8.00 v_mfma_i32_32x32x16_i8 v[0:15], v[2:3], v[4:5], v[0:15]
# CHECK-NEXT: - - - - - - 8.00 v_mfma_i32_32x32x16_i8 a[0:15], v[2:3], v[4:5], a[0:15]
# CHECK-NEXT: - - - - - - 2.00 v_mfma_f32_4x4x4_16b_f16 v[0:3], v[0:1], v[2:3], v[2:5]
# CHECK-NEXT: - - - - - - 2.00 v_mfma_f32_4x4x4_16b_f16 a[0:3], v[0:1], v[2:3], a[2:5]
# CHECK-NEXT: - - - - - - 8.00 v_mfma_f32_16x16x4_4b_f16 v[0:15], v[2:3], v[4:5], v[18:33]
# CHECK-NEXT: - - - - - - 8.00 v_mfma_f32_16x16x4_4b_f16 a[0:15], v[2:3], v[4:5], a[18:33]
# CHECK-NEXT: - - - - - - 16.00 v_mfma_f32_32x32x4_2b_f16 v[0:31], v[0:1], v[2:3], v[34:65]
# CHECK-NEXT: - - - - - - 16.00 v_mfma_f32_32x32x4_2b_f16 a[0:31], v[0:1], v[2:3], a[34:65]
# CHECK-NEXT: - - - - - - 2.00 v_mfma_f32_4x4x4_16b_bf16 v[0:3], v[0:1], v[2:3], v[2:5]
# CHECK-NEXT: - - - - - - 2.00 v_mfma_f32_4x4x4_16b_bf16 a[0:3], v[0:1], v[2:3], a[2:5]
# CHECK-NEXT: - - - - - - 8.00 v_mfma_f32_16x16x4_4b_bf16 v[0:15], v[2:3], v[4:5], v[18:33]
# CHECK-NEXT: - - - - - - 8.00 v_mfma_f32_16x16x4_4b_bf16 a[0:15], v[2:3], v[4:5], a[18:33]
# CHECK-NEXT: - - - - - - 16.00 v_mfma_f32_32x32x4_2b_bf16 v[0:31], v[0:1], v[2:3], v[34:65]
# CHECK-NEXT: - - - - - - 16.00 v_mfma_f32_32x32x4_2b_bf16 a[0:31], v[0:1], v[2:3], a[34:65]
# CHECK-NEXT: - - - - - - 2.00 v_mfma_f32_4x4x1_16b_f32 v[0:3], v0, v1, v[2:5]
# CHECK-NEXT: - - - - - - 2.00 v_mfma_f32_4x4x1_16b_f32 a[0:3], v0, v1, a[2:5]
# CHECK-NEXT: - - - - - - 8.00 v_mfma_f32_16x16x1_4b_f32 v[0:15], v0, v1, v[18:33]
# CHECK-NEXT: - - - - - - 8.00 v_mfma_f32_16x16x1_4b_f32 a[0:15], v0, v1, a[18:33]
# CHECK-NEXT: - - - - - - 8.00 v_mfma_f32_16x16x4_f32 v[0:3], v0, v1, v[2:5]
# CHECK-NEXT: - - - - - - 8.00 v_mfma_f32_16x16x4_f32 a[0:3], v0, v1, a[2:5]
# CHECK-NEXT: - - - - - - 16.00 v_mfma_f32_32x32x1_2b_f32 v[0:31], v0, v1, v[34:65] blgp:7
# CHECK-NEXT: - - - - - - 16.00 v_mfma_f32_32x32x1_2b_f32 a[0:31], v0, v1, a[34:65] blgp:7
# CHECK-NEXT: - - - - - - 16.00 v_mfma_f32_32x32x2_f32 v[0:15], v0, v1, v[18:33]
# CHECK-NEXT: - - - - - - 16.00 v_mfma_f32_32x32x2_f32 a[0:15], v0, v1, a[18:33]
# CHECK-NEXT: - - - - - - 2.00 v_mfma_i32_4x4x4_16b_i8 v[0:3], v0, v1, v[2:5]
# CHECK-NEXT: - - - - - - 2.00 v_mfma_i32_4x4x4_16b_i8 a[0:3], v0, v1, a[2:5]
# CHECK-NEXT: - - - - - - 8.00 v_mfma_i32_16x16x4_4b_i8 v[0:15], v0, v1, v[18:33]
# CHECK-NEXT: - - - - - - 8.00 v_mfma_i32_16x16x4_4b_i8 a[0:15], v0, v1, a[18:33]
# CHECK-NEXT: - - - - - - 16.00 v_mfma_i32_32x32x4_2b_i8 v[0:31], v0, v1, v[34:65]
# CHECK-NEXT: - - - - - - 16.00 v_mfma_i32_32x32x4_2b_i8 a[0:31], v0, v1, a[34:65]
# CHECK-NEXT: - - - - - - 4.00 v_smfmac_f32_16x16x32_f16 v[10:13], a[2:3], v[4:7], v0 cbsz:3 abid:1
# CHECK-NEXT: - - - - - - 4.00 v_smfmac_f32_16x16x32_f16 a[10:13], v[2:3], a[4:7], v1
# CHECK-NEXT: - - - - - - 8.00 v_smfmac_f32_32x32x16_f16 v[10:25], a[2:3], v[4:7], v2 cbsz:3 abid:1
# CHECK-NEXT: - - - - - - 8.00 v_smfmac_f32_32x32x16_f16 a[10:25], v[2:3], a[4:7], v3
# CHECK-NEXT: - - - - - - 4.00 v_smfmac_f32_16x16x32_bf16 v[10:13], a[2:3], v[4:7], v4 cbsz:3 abid:1
# CHECK-NEXT: - - - - - - 4.00 v_smfmac_f32_16x16x32_bf16 a[10:13], v[2:3], a[4:7], v5
# CHECK-NEXT: - - - - - - 4.00 v_smfmac_i32_16x16x64_i8 v[10:13], a[2:3], v[4:7], v8 cbsz:3 abid:1
# CHECK-NEXT: - - - - - - 4.00 v_smfmac_i32_16x16x64_i8 a[10:13], v[2:3], a[4:7], v9
# CHECK-NEXT: - - - - - - 8.00 v_smfmac_i32_32x32x32_i8 v[10:25], a[2:3], v[4:7], v10 cbsz:3 abid:1
# CHECK-NEXT: - - - - - - 8.00 v_smfmac_i32_32x32x32_i8 a[10:25], v[2:3], a[4:7], v11

2 comments on commit f0484e0

@hctim
Copy link
Collaborator

@hctim hctim commented on f0484e0 Feb 29, 2024

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

G'day Matt,

Looks like this test has a heap-use-after-free bug that's tripped on the sanitizer buildbots:

https://lab.llvm.org/buildbot/#/builders/5/builds/41382/steps/9/logs/stdio

******************** TEST 'LLVM :: tools/llvm-mca/AMDGPU/gfx940.s' FAILED ********************
Exit Code: 2
Command Output (stderr):
--
RUN: at line 1: /b/sanitizer-x86_64-linux-fast/build/llvm_build_asan_ubsan/bin/llvm-mca -mtriple=amdgcn -mcpu=gfx940 --timeline --iterations=1 --timeline-max-cycles=0 < /b/sanitizer-x86_64-linux-fast/build/llvm-project/llvm/test/tools/llvm-mca/AMDGPU/gfx940.s | /b/sanitizer-x86_64-linux-fast/build/llvm_build_asan_ubsan/bin/FileCheck /b/sanitizer-x86_64-linux-fast/build/llvm-project/llvm/test/tools/llvm-mca/AMDGPU/gfx940.s
+ /b/sanitizer-x86_64-linux-fast/build/llvm_build_asan_ubsan/bin/llvm-mca -mtriple=amdgcn -mcpu=gfx940 --timeline --iterations=1 --timeline-max-cycles=0
+ /b/sanitizer-x86_64-linux-fast/build/llvm_build_asan_ubsan/bin/FileCheck /b/sanitizer-x86_64-linux-fast/build/llvm-project/llvm/test/tools/llvm-mca/AMDGPU/gfx940.s
=================================================================
==2072808==ERROR: AddressSanitizer: heap-use-after-free on address 0x51600000ef9c at pc 0x55cb5d218d94 bp 0x7ffe73ec4f90 sp 0x7ffe73ec4f88
READ of size 1 at 0x51600000ef9c thread T0
    #0 0x55cb5d218d93 in getEndGroup /b/sanitizer-x86_64-linux-fast/build/llvm-project/llvm/include/llvm/MCA/Instruction.h:592:37
    #1 0x55cb5d218d93 in llvm::mca::InOrderIssueStage::updateCarriedOver() /b/sanitizer-x86_64-linux-fast/build/llvm-project/llvm/lib/MCA/Stages/InOrderIssueStage.cpp:327:37
    #2 0x55cb5d2192e1 in llvm::mca::InOrderIssueStage::cycleStart() /b/sanitizer-x86_64-linux-fast/build/llvm-project/llvm/lib/MCA/Stages/InOrderIssueStage.cpp:395:3
    #3 0x55cb5d204389 in llvm::mca::Pipeline::runCycle() /b/sanitizer-x86_64-linux-fast/build/llvm-project/llvm/lib/MCA/Pipeline.cpp:60:16
    #4 0x55cb5d20370b in llvm::mca::Pipeline::run() /b/sanitizer-x86_64-linux-fast/build/llvm-project/llvm/lib/MCA/Pipeline.cpp:43:21
    #5 0x55cb5c872d5b in runPipeline(llvm::mca::Pipeline&) /b/sanitizer-x86_64-linux-fast/build/llvm-project/llvm/tools/llvm-mca/llvm-mca.cpp:308:33
    #6 0x55cb5c86ae33 in main /b/sanitizer-x86_64-linux-fast/build/llvm-project/llvm/tools/llvm-mca/llvm-mca.cpp:750:10
    #7 0x7fe4daa23a8f  (/lib/x86_64-linux-gnu/libc.so.6+0x23a8f) (BuildId: d320ce4e63925d698610ed423fc4b1f0e8ed51f1)
    #8 0x7fe4daa23b48 in __libc_start_main (/lib/x86_64-linux-gnu/libc.so.6+0x23b48) (BuildId: d320ce4e63925d698610ed423fc4b1f0e8ed51f1)
    #9 0x55cb5c787724 in _start (/b/sanitizer-x86_64-linux-fast/build/llvm_build_asan_ubsan/bin/llvm-mca+0x2862724)
0x51600000ef9c is located 540 bytes inside of 608-byte region [0x51600000ed80,0x51600000efe0)
freed by thread T0 here:
    #0 0x55cb5c85681d in operator delete(void*) /b/sanitizer-x86_64-linux-fast/build/llvm-project/compiler-rt/lib/asan/asan_new_delete.cpp:152:3
    #1 0x55cb5c888fcf in operator() /b/sanitizer-x86_64-linux-fast/build/libcxx_build_asan_ubsan/include/c++/v1/__memory/unique_ptr.h:68:5
    #2 0x55cb5c888fcf in reset /b/sanitizer-x86_64-linux-fast/build/libcxx_build_asan_ubsan/include/c++/v1/__memory/unique_ptr.h:279:7
    #3 0x55cb5c888fcf in ~unique_ptr /b/sanitizer-x86_64-linux-fast/build/libcxx_build_asan_ubsan/include/c++/v1/__memory/unique_ptr.h:249:71
    #4 0x55cb5c888fcf in llvm::SmallVectorTemplateBase<std::__1::unique_ptr<llvm::mca::Instruction, std::__1::default_delete<llvm::mca::Instruction>>, false>::destroy_range(std::__1::unique_ptr<llvm::mca::Instruction, std::__1::default_delete<llvm::mca::Instruction>>*, std::__1::unique_ptr<llvm::mca::Instruction, std::__1::default_delete<llvm::mca::Instruction>>*) /b/sanitizer-x86_64-linux-fast/build/llvm-project/llvm/include/llvm/ADT/SmallVector.h:354:11
    #5 0x55cb5d20a3e5 in llvm::SmallVectorImpl<std::__1::unique_ptr<llvm::mca::Instruction, std::__1::default_delete<llvm::mca::Instruction>>>::erase(std::__1::unique_ptr<llvm::mca::Instruction, std::__1::default_delete<llvm::mca::Instruction>> const*, std::__1::unique_ptr<llvm::mca::Instruction, std::__1::default_delete<llvm::mca::Instruction>> const*) /b/sanitizer-x86_64-linux-fast/build/llvm-project/llvm/include/llvm/ADT/SmallVector.h:775:5
    #6 0x55cb5d20a071 in llvm::mca::EntryStage::cycleEnd() /b/sanitizer-x86_64-linux-fast/build/llvm-project/llvm/lib/MCA/Stages/EntryStage.cpp:78:18
    #7 0x55cb5d2047e4 in llvm::mca::Pipeline::runCycle() /b/sanitizer-x86_64-linux-fast/build/llvm-project/llvm/lib/MCA/Pipeline.cpp:78:14
    #8 0x55cb5d20370b in llvm::mca::Pipeline::run() /b/sanitizer-x86_64-linux-fast/build/llvm-project/llvm/lib/MCA/Pipeline.cpp:43:21
    #9 0x55cb5c872d5b in runPipeline(llvm::mca::Pipeline&) /b/sanitizer-x86_64-linux-fast/build/llvm-project/llvm/tools/llvm-mca/llvm-mca.cpp:308:33
    #10 0x55cb5c86ae33 in main /b/sanitizer-x86_64-linux-fast/build/llvm-project/llvm/tools/llvm-mca/llvm-mca.cpp:750:10
    #11 0x7fe4daa23a8f  (/lib/x86_64-linux-gnu/libc.so.6+0x23a8f) (BuildId: d320ce4e63925d698610ed423fc4b1f0e8ed51f1)
previously allocated by thread T0 here:
    #0 0x55cb5c855fbd in operator new(unsigned long) /b/sanitizer-x86_64-linux-fast/build/llvm-project/compiler-rt/lib/asan/asan_new_delete.cpp:95:3
    #1 0x55cb5d209762 in make_unique<llvm::mca::Instruction, const llvm::mca::Instruction &> /b/sanitizer-x86_64-linux-fast/build/libcxx_build_asan_ubsan/include/c++/v1/__memory/unique_ptr.h:621:26
    #2 0x55cb5d209762 in llvm::mca::EntryStage::getNextInstruction() /b/sanitizer-x86_64-linux-fast/build/llvm-project/llvm/lib/MCA/Stages/EntryStage.cpp:40:39
    #3 0x55cb5d209d7c in llvm::mca::EntryStage::execute(llvm::mca::InstRef&) /b/sanitizer-x86_64-linux-fast/build/llvm-project/llvm/lib/MCA/Stages/EntryStage.cpp:54:10
    #4 0x55cb5d2045cd in llvm::mca::Pipeline::runCycle() /b/sanitizer-x86_64-linux-fast/build/llvm-project/llvm/lib/MCA/Pipeline.cpp:69:22
    #5 0x55cb5d20370b in llvm::mca::Pipeline::run() /b/sanitizer-x86_64-linux-fast/build/llvm-project/llvm/lib/MCA/Pipeline.cpp:43:21
    #6 0x55cb5c872d5b in runPipeline(llvm::mca::Pipeline&) /b/sanitizer-x86_64-linux-fast/build/llvm-project/llvm/tools/llvm-mca/llvm-mca.cpp:308:33
    #7 0x55cb5c86ae33 in main /b/sanitizer-x86_64-linux-fast/build/llvm-project/llvm/tools/llvm-mca/llvm-mca.cpp:750:10
    #8 0x7fe4daa23a8f  (/lib/x86_64-linux-gnu/libc.so.6+0x23a8f) (BuildId: d320ce4e63925d698610ed423fc4b1f0e8ed51f1)
SUMMARY: AddressSanitizer: heap-use-after-free /b/sanitizer-x86_64-linux-fast/build/llvm-project/llvm/include/llvm/MCA/Instruction.h:592:37 in getEndGroup
Shadow bytes around the buggy address:
  0x51600000ed00: fa fa fa fa fa fa fa fa fa fa fa fa fa fa fa fa
  0x51600000ed80: fd fd fd fd fd fd fd fd fd fd fd fd fd fd fd fd
  0x51600000ee00: fd fd fd fd fd fd fd fd fd fd fd fd fd fd fd fd
  0x51600000ee80: fd fd fd fd fd fd fd fd fd fd fd fd fd fd fd fd
  0x51600000ef00: fd fd fd fd fd fd fd fd fd fd fd fd fd fd fd fd
=>0x51600000ef80: fd fd fd[fd]fd fd fd fd fd fd fd fd fa fa fa fa
  0x51600000f000: fa fa fa fa fa fa fa fa fa fa fa fa fa fa fa fa
  0x51600000f080: 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00
  0x51600000f100: 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00
  0x51600000f180: 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00
  0x51600000f200: 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00
Shadow byte legend (one shadow byte represents 8 application bytes):
  Addressable:           00
  Partially addressable: 01 02 03 04 05 06 07 
  Heap left redzone:       fa
  Freed heap region:       fd
  Stack left redzone:      f1
  Stack mid redzone:       f2
  Stack right redzone:     f3
  Stack after return:      f5
  Stack use after scope:   f8
  Global redzone:          f9
  Global init order:       f6
  Poisoned by user:        f7
  Container overflow:      fc
  Array cookie:            ac
  Intra object redzone:    bb
  ASan internal:           fe
  Left alloca redzone:     ca
  Right alloca redzone:    cb
==2072808==ABORTING
FileCheck error: '<stdin>' is empty.
FileCheck command line:  /b/sanitizer-x86_64-linux-fast/build/llvm_build_asan_ubsan/bin/FileCheck /b/sanitizer-x86_64-linux-fast/build/llvm-project/llvm/test/tools/llvm-mca/AMDGPU/gfx940.s
--
********************

@hctim
Copy link
Collaborator

@hctim hctim commented on f0484e0 Feb 29, 2024

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Should be reproducible without the full sanitizer build script, using:

$ cmake \
-DCMAKE_C_COMPILER=clang \
-DCMAKE_CXX_COMPILER=clang++ \
-GNinja \
-DCMAKE_BUILD_TYPE=Release \
-DCMAKE_C_FLAGS="-fsanitize=address" \
-DCMAKE_CXX_FLAGS="-fsanitize=address" \
-DLLVM_USE_SANITIZER=Address \
-DLLVM_ENABLE_ASSERTIONS=On \
/path/to/llvm
$ LIT_OPTS='--filter gfx940' ninja check-llvm

Please sign in to comment.