Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

AMDGPU: Add scheduling test for gfx940 #83220

Merged
merged 3 commits into from
Feb 29, 2024
Merged

Conversation

arsenm
Copy link
Contributor

@arsenm arsenm commented Feb 28, 2024

I'm not sure the f64 fma cases are correct.

I'm not sure the f64 fma cases are correct.
@llvmbot
Copy link
Collaborator

llvmbot commented Feb 28, 2024

@llvm/pr-subscribers-backend-amdgpu

Author: Matt Arsenault (arsenm)

Changes

I'm not sure the f64 fma cases are correct.


Full diff: https://github.com/llvm/llvm-project/pull/83220.diff

1 Files Affected:

  • (added) llvm/test/tools/llvm-mca/AMDGPU/gfx940.s (+55)
diff --git a/llvm/test/tools/llvm-mca/AMDGPU/gfx940.s b/llvm/test/tools/llvm-mca/AMDGPU/gfx940.s
new file mode 100644
index 00000000000000..0c723859b74577
--- /dev/null
+++ b/llvm/test/tools/llvm-mca/AMDGPU/gfx940.s
@@ -0,0 +1,55 @@
+# RUN: llvm-mca -mtriple=amdgcn -mcpu=gfx940 --timeline --iterations=1 --timeline-max-cycles=0 < %s | FileCheck %s
+
+# CHECK: Iterations:        1
+# CHECK: Instructions:      21
+# CHECK: Total Cycles:      102
+# CHECK: Total uOps:        27
+
+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]
+
+# 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]

@arsenm arsenm merged commit f0484e0 into llvm:main Feb 29, 2024
3 of 4 checks passed
@arsenm arsenm deleted the gfx940-sched-test branch February 29, 2024 03:43
@metaflow
Copy link
Contributor

Hi Matt, this test fails with -DLLVM_USE_SANITIZER=Address

ninja check-llvm-tools-llvm-mca
[4176/4177] Running lit suite [...]/llvm-project/llvm/test/tools/llvm-mca
FAIL: LLVM :: tools/llvm-mca/AMDGPU/gfx940.s (247 of 1125)
******************** TEST 'LLVM :: tools/llvm-mca/AMDGPU/gfx940.s' FAILED ********************
Exit Code: 2

Command Output (stderr):
--
RUN: at line 1: [...]/llvm-project/build/bin/llvm-mca -mtriple=amdgcn -mcpu=gfx940 --timeline --iterations=1 --timeline-max-cycles=0 < [...]/llvm-project/llvm/test/tools/llvm-mca/AMDGPU/gfx940.s | [...]/llvm-project/build/bin/FileCheck [...]/llvm-project/llvm/test/tools/llvm-mca/AMDGPU/gfx940.s
+ [...]/llvm-project/build/bin/llvm-mca -mtriple=amdgcn -mcpu=gfx940 --timeline --iterations=1 --timeline-max-cycles=0
+ [...]/llvm-project/build/bin/FileCheck [...]/llvm-project/llvm/test/tools/llvm-mca/AMDGPU/gfx940.s
=================================================================
==2365057==ERROR: AddressSanitizer: heap-use-after-free on address 0x61600000ef9c at pc 0x558b60ea614a bp 0x7ffee1971e90 sp 0x7ffee1971e88
READ of size 1 at 0x61600000ef9c thread T0
    #0 0x558b60ea6149 in getEndGroup [...]/llvm-project/llvm/include/llvm/MCA/Instruction.h:592:37
    #1 0x558b60ea6149 in llvm::mca::InOrderIssueStage::updateCarriedOver() [...]/llvm-project/llvm/lib/MCA/Stages/InOrderIssueStage.cpp:327:37
    #2 0x558b60ea65a6 in llvm::mca::InOrderIssueStage::cycleStart() [...]/llvm-project/llvm/lib/MCA/Stages/InOrderIssueStage.cpp:395:3
    #3 0x558b60e94390 in llvm::mca::Pipeline::runCycle() [...]/llvm-project/llvm/lib/MCA/Pipeline.cpp:60:16
    #4 0x558b60e93b17 in llvm::mca::Pipeline::run() [...]/llvm-project/llvm/lib/MCA/Pipeline.cpp:43:21
    #5 0x558b6037278a in runPipeline(llvm::mca::Pipeline&) [...]/llvm-project/llvm/tools/llvm-mca/llvm-mca.cpp:308:33
    #6 0x558b6036b35f in main [...]/llvm-project/llvm/tools/llvm-mca/llvm-mca.cpp:750:10
    #7 0x7f2c63e456c9 in __libc_start_call_main csu/../sysdeps/nptl/libc_start_call_main.h:58:16
    #8 0x7f2c63e45784 in __libc_start_main csu/../csu/libc-start.c:360:3
    #9 0x558b60285e20 in _start ([...]/llvm-project/build/bin/llvm-mca+0xa87e20) (BuildId: 27cf574469c56298a8763364567a5afdf281f38f)

0x61600000ef9c is located 540 bytes inside of 608-byte region [0x61600000ed80,0x61600000efe0)
freed by thread T0 here:
    #0 0x558b60350b31 in operator delete(void*) ([...]/llvm-project/build/bin/llvm-mca+0xb52b31) (BuildId: 27cf574469c56298a8763364567a5afdf281f38f)
    #1 0x558b6037e871 in operator() /usr/bin/../lib/gcc/x86_64-linux-gnu/13/../../../../include/c++/13/bits/unique_ptr.h:99:2
    #2 0x558b6037e871 in std::unique_ptr<llvm::mca::Instruction, std::default_delete<llvm::mca::Instruction>>::~unique_ptr() /usr/bin/../lib/gcc/x86_64-linux-gnu/13/../../../../include/c++/13/bits/unique_ptr.h:404:4
    #3 0x558b60e999db in destroy_range [...]/llvm-project/llvm/include/llvm/ADT/SmallVector.h:354:11
    #4 0x558b60e999db in llvm::SmallVectorImpl<std::unique_ptr<llvm::mca::Instruction, std::default_delete<llvm::mca::Instruction>>>::erase(std::unique_ptr<llvm::mca::Instruction, std::default_delete<llvm::mca::Instruction>> const*, std::unique_ptr<llvm::mca::Instruction, std::default_delete<llvm::mca::Instruction>> const*) [...]/llvm-project/llvm/include/llvm/ADT/SmallVector.h:775:5
    #5 0x558b60e99723 in llvm::mca::EntryStage::cycleEnd() [...]/llvm-project/llvm/lib/MCA/Stages/EntryStage.cpp:78:18
    #6 0x558b60e9472e in llvm::mca::Pipeline::runCycle() [...]/llvm-project/llvm/lib/MCA/Pipeline.cpp:78:14
    #7 0x558b60e93b17 in llvm::mca::Pipeline::run() [...]/llvm-project/llvm/lib/MCA/Pipeline.cpp:43:21
    #8 0x558b6037278a in runPipeline(llvm::mca::Pipeline&) [...]/llvm-project/llvm/tools/llvm-mca/llvm-mca.cpp:308:33
    #9 0x558b6036b35f in main [...]/llvm-project/llvm/tools/llvm-mca/llvm-mca.cpp:750:10
    #10 0x7f2c63e456c9 in __libc_start_call_main csu/../sysdeps/nptl/libc_start_call_main.h:58:16

previously allocated by thread T0 here:
    #0 0x558b603502b1 in operator new(unsigned long) ([...]/llvm-project/build/bin/llvm-mca+0xb522b1) (BuildId: 27cf574469c56298a8763364567a5afdf281f38f)
    #1 0x558b60e98783 in make_unique<llvm::mca::Instruction, const llvm::mca::Instruction &> /usr/bin/../lib/gcc/x86_64-linux-gnu/13/../../../../include/c++/13/bits/unique_ptr.h:1070:30
    #2 0x558b60e98783 in llvm::mca::EntryStage::getNextInstruction() [...]/llvm-project/llvm/lib/MCA/Stages/EntryStage.cpp:40:39
    #3 0x558b60e99224 in llvm::mca::EntryStage::execute(llvm::mca::InstRef&) [...]/llvm-project/llvm/lib/MCA/Stages/EntryStage.cpp:54:10
    #4 0x558b60e94587 in llvm::mca::Pipeline::runCycle() [...]/llvm-project/llvm/lib/MCA/Pipeline.cpp:69:22
    #5 0x558b60e93b17 in llvm::mca::Pipeline::run() [...]/llvm-project/llvm/lib/MCA/Pipeline.cpp:43:21
    #6 0x558b6037278a in runPipeline(llvm::mca::Pipeline&) [...]/llvm-project/llvm/tools/llvm-mca/llvm-mca.cpp:308:33
    #7 0x558b6036b35f in main [...]/llvm-project/llvm/tools/llvm-mca/llvm-mca.cpp:750:10
    #8 0x7f2c63e456c9 in __libc_start_call_main csu/../sysdeps/nptl/libc_start_call_main.h:58:16

SUMMARY: AddressSanitizer: heap-use-after-free [...]/llvm-project/llvm/include/llvm/MCA/Instruction.h:592:37 in getEndGroup
Shadow bytes around the buggy address:
  0x61600000ed00: fa fa fa fa fa fa fa fa fa fa fa fa fa fa fa fa
  0x61600000ed80: fd fd fd fd fd fd fd fd fd fd fd fd fd fd fd fd
  0x61600000ee00: fd fd fd fd fd fd fd fd fd fd fd fd fd fd fd fd
  0x61600000ee80: fd fd fd fd fd fd fd fd fd fd fd fd fd fd fd fd
  0x61600000ef00: fd fd fd fd fd fd fd fd fd fd fd fd fd fd fd fd
=>0x61600000ef80: fd fd fd[fd]fd fd fd fd fd fd fd fd fa fa fa fa
  0x61600000f000: fa fa fa fa fa fa fa fa fa fa fa fa fa fa fa fa
  0x61600000f080: 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00
  0x61600000f100: 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00
  0x61600000f180: 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00 00
  0x61600000f200: 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
==2365057==ABORTING
FileCheck error: '<stdin>' is empty.
FileCheck command line:  [...]/llvm-project/build/bin/FileCheck [...]/llvm-project/llvm/test/tools/llvm-mca/AMDGPU/gfx940.s

--

********************
********************
Failed Tests (1):
  LLVM :: tools/llvm-mca/AMDGPU/gfx940.s

Probably the issue is not with test itself but in underlying code but I don't know much about it.

@arsenm
Copy link
Contributor Author

arsenm commented Feb 29, 2024

Probably the issue is not with test itself but in underlying code but I don't know much about it.

This is obviously true, no test should be capable of crashing a tool

hctim added a commit that referenced this pull request Feb 29, 2024
This reverts commit f0484e0.

Reason: Broke the sanitizer build bots. See the github comments on
f0484e0
for more information
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Projects
None yet
Development

Successfully merging this pull request may close these issues.

None yet

5 participants