[llvm] AMDGPU: Add scheduling test for gfx940 (PR #83220)
via llvm-commits
llvm-commits at lists.llvm.org
Tue Feb 27 20:57:03 PST 2024
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-backend-amdgpu
Author: Matt Arsenault (arsenm)
<details>
<summary>Changes</summary>
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
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]
``````````
</details>
https://github.com/llvm/llvm-project/pull/83220
More information about the llvm-commits
mailing list