Skip to content

[AMDGPU] Add more gfx1250 MC tests. NFC. #152388

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.

Sign up for GitHub

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

Jump to bottom
Merged
merged 1 commit into from
Aug 6, 2025
Merged

[AMDGPU] Add more gfx1250 MC tests. NFC. #152388

merged 1 commit into from
Aug 6, 2025

Conversation

@rampitec
Copy link
Collaborator

@rampitec rampitec commented Aug 6, 2025

These are already working, but left downstream.

@rampitec
[AMDGPU] Add more gfx1250 MC tests. NFC.
5c97ea7 
These are already working, but left downstream.
@rampitec rampitec requested a review from changpeng August 6, 2025 21:08
@rampitec Graphite App
Copy link
Collaborator Author

rampitec commented Aug 6, 2025

This stack of pull requests is managed by Graphite. Learn more about stacking.

@rampitec rampitec requested a review from shiltian August 6, 2025 21:08
@rampitec rampitec marked this pull request as ready for review August 6, 2025 21:08
@llvmbot llvmbot added backend:AMDGPU llvm:mc Machine (object) code labels Aug 6, 2025
@llvmbot
Copy link
Member

llvmbot commented Aug 6, 2025 •
edited
Loading

@llvm/pr-subscribers-mc

@llvm/pr-subscribers-backend-amdgpu

Author: Stanislav Mekhanoshin (rampitec)

Changes

These are already working, but left downstream.

Patch is 687.51 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/152388.diff

14 Files Affected:

  • (added) llvm/test/MC/AMDGPU/gfx1250_asm_features.s (+32)
  • (modified) llvm/test/MC/AMDGPU/gfx1250_asm_unsupported.s (+14)
  • (modified) llvm/test/MC/AMDGPU/gfx1250_asm_vbuffer_mubuf.s (+2304)
  • (modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop3_err.s (+165)
  • (added) llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop2_err.s (+13)
  • (added) llvm/test/MC/AMDGPU/gfx1250_asm_vop3cx.s (+3413)
  • (added) llvm/test/MC/AMDGPU/gfx1250_asm_vop3p_dpp16.s (+14)
  • (added) llvm/test/MC/AMDGPU/gfx1250_asm_vop3p_dpp8.s (+26)
  • (added) llvm/test/MC/AMDGPU/gfx1250_asm_vsample_err.s (+175)
  • (modified) llvm/test/MC/AMDGPU/gfx1250_err.s (+30)
  • (modified) llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vbuffer_mubuf.txt (+2133)
  • (added) llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3cx.txt (+3413)
  • (added) llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3p_dpp16.txt (+10)
  • (added) llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3p_dpp8.txt (+19)
diff --git a/llvm/test/MC/AMDGPU/gfx1250_asm_features.s b/llvm/test/MC/AMDGPU/gfx1250_asm_features.s
new file mode 100644
index 0000000000000..013b790b2b4b6
--- /dev/null
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_features.s
@@ -0,0 +1,32 @@
+// RUN: llvm-mc -triple=amdgcn -mcpu=gfx1250 -show-encoding %s | FileCheck --check-prefixes=GFX1250 %s
+
+//
+// Elements of CPol operand can be given in any order
+//
+
+s_load_b32 s4, s[2:3], 10 th:TH_LOAD_NT scope:SCOPE_SE nv
+// GFX1250: encoding: [0x01,0x01,0xb0,0xf4,0x0a,0x00,0x00,0xf8]
+
+s_load_b32 s4, s[2:3], 10 scope:SCOPE_SE nv th:TH_LOAD_NT
+// GFX1250: encoding: [0x01,0x01,0xb0,0xf4,0x0a,0x00,0x00,0xf8]
+
+s_load_b32 s4, s[2:3], 10 nv scope:SCOPE_SE th:TH_LOAD_NT
+// GFX1250: encoding: [0x01,0x01,0xb0,0xf4,0x0a,0x00,0x00,0xf8]
+
+buffer_load_b32 v5, v1, s[8:11], s3 offen offset:4095 th:TH_LOAD_NT scope:SCOPE_SE nv
+// GFX1250: encoding: [0x83,0x00,0x05,0xc4,0x05,0x10,0x94,0x40,0x01,0xff,0x0f,0x00]
+
+buffer_load_b32 v5, v1, s[8:11], s3 offen offset:4095 scope:SCOPE_SE nv th:TH_LOAD_NT
+// GFX1250: encoding: [0x83,0x00,0x05,0xc4,0x05,0x10,0x94,0x40,0x01,0xff,0x0f,0x00]
+
+buffer_load_b32 v5, v1, s[8:11], s3 offen offset:4095 nv scope:SCOPE_SE th:TH_LOAD_NT
+// GFX1250: encoding: [0x83,0x00,0x05,0xc4,0x05,0x10,0x94,0x40,0x01,0xff,0x0f,0x00]
+
+global_load_b32 v0, v[2:3], off th:TH_LOAD_NT scope:SCOPE_SE nv
+// GFX1250: encoding: [0xfc,0x00,0x05,0xee,0x00,0x00,0x14,0x00,0x02,0x00,0x00,0x00]
+
+global_load_b32 v0, v[2:3], off scope:SCOPE_SE nv th:TH_LOAD_NT
+// GFX1250: encoding: [0xfc,0x00,0x05,0xee,0x00,0x00,0x14,0x00,0x02,0x00,0x00,0x00]
+
+global_load_b32 v0, v[2:3], off nv scope:SCOPE_SE th:TH_LOAD_NT
+// GFX1250: encoding: [0xfc,0x00,0x05,0xee,0x00,0x00,0x14,0x00,0x02,0x00,0x00,0x00]
diff --git a/llvm/test/MC/AMDGPU/gfx1250_asm_unsupported.s b/llvm/test/MC/AMDGPU/gfx1250_asm_unsupported.s
index 89bd507942a22..7681a32dce7dc 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_unsupported.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_unsupported.s
@@ -97,6 +97,20 @@ v_interp_p10_rtz_f16_f32 v0, v1, v2, v3
 v_interp_p2_rtz_f16_f32 v0, v1, v2, v3
 // GFX1250-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
 
+;; *xf32
+
+v_mfma_f32_16x16x8_xf32 a[0:3], v[2:3], v[4:5], a[2:5]
+// GFX1250-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
+
+v_mfma_f32_16x16x8xf32 a[0:3], v[2:3], v[4:5], a[2:5]
+// GFX1250-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
+
+v_mfma_f32_32x32x4_xf32 a[0:15], v[2:3], v[4:5], a[18:33]
+// GFX1250-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
+
+v_mfma_f32_32x32x4xf32 a[0:15], v[2:3], v[4:5], a[18:33]
+// GFX1250-ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
+
 ;; Export, S_WAIT_EXPCNT and S_WAIT_EVENT
 
 export mrt0 off, off, off, off
diff --git a/llvm/test/MC/AMDGPU/gfx1250_asm_vbuffer_mubuf.s b/llvm/test/MC/AMDGPU/gfx1250_asm_vbuffer_mubuf.s
index 7a4da255b5594..0b8f190a7ae0d 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_vbuffer_mubuf.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_vbuffer_mubuf.s
@@ -1,6 +1,2310 @@
 // RUN: llvm-mc -triple=amdgcn -mcpu=gfx1250 -show-encoding %s | FileCheck --check-prefix=GFX1250 %s
 // RUN: not llvm-mc -triple=amdgcn -mcpu=gfx1200 -show-encoding %s 2>&1 | FileCheck --check-prefix=GFX12-ERR --implicit-check-not=error: --strict-whitespace %s
 
+buffer_load_b32 v5, off, s[8:11], s3 offset:8388607
+// GFX1250: buffer_load_b32 v5, off, s[8:11], s3 offset:8388607 ; encoding: [0x03,0x00,0x05,0xc4,0x05,0x10,0x80,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_b32 v255, off, s[8:11], s3 offset:8388607
+// GFX1250: buffer_load_b32 v255, off, s[8:11], s3 offset:8388607 ; encoding: [0x03,0x00,0x05,0xc4,0xff,0x10,0x80,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_b32 v5, off, s[12:15], s3 offset:8388607
+// GFX1250: buffer_load_b32 v5, off, s[12:15], s3 offset:8388607 ; encoding: [0x03,0x00,0x05,0xc4,0x05,0x18,0x80,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_b32 v5, off, s[96:99], s3 offset:8388607
+// GFX1250: buffer_load_b32 v5, off, s[96:99], s3 offset:8388607 ; encoding: [0x03,0x00,0x05,0xc4,0x05,0xc0,0x80,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_b32 v5, off, s[8:11], s101 offset:8388607
+// GFX1250: buffer_load_b32 v5, off, s[8:11], s101 offset:8388607 ; encoding: [0x65,0x00,0x05,0xc4,0x05,0x10,0x80,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_b32 v5, off, s[8:11], m0 offset:8388607
+// GFX1250: buffer_load_b32 v5, off, s[8:11], m0 offset:8388607 ; encoding: [0x7d,0x00,0x05,0xc4,0x05,0x10,0x80,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_b32 v5, v0, s[8:11], s3 idxen offset:8388607
+// GFX1250: buffer_load_b32 v5, v0, s[8:11], s3 idxen offset:8388607 ; encoding: [0x03,0x00,0x05,0xc4,0x05,0x10,0x80,0x80,0x00,0xff,0xff,0x7f]
+
+buffer_load_b32 v5, v0, s[8:11], s3 offen offset:8388607
+// GFX1250: buffer_load_b32 v5, v0, s[8:11], s3 offen offset:8388607 ; encoding: [0x03,0x00,0x05,0xc4,0x05,0x10,0x80,0x40,0x00,0xff,0xff,0x7f]
+
+buffer_load_b32 v5, off, s[8:11], s3
+// GFX1250: buffer_load_b32 v5, off, s[8:11], s3    ; encoding: [0x03,0x00,0x05,0xc4,0x05,0x10,0x80,0x00,0x00,0x00,0x00,0x00]
+
+buffer_load_b32 v5, off, s[8:11], s3 offset:0
+// GFX1250: buffer_load_b32 v5, off, s[8:11], s3    ; encoding: [0x03,0x00,0x05,0xc4,0x05,0x10,0x80,0x00,0x00,0x00,0x00,0x00]
+
+buffer_load_b32 v5, off, s[8:11], s3 offset:7
+// GFX1250: buffer_load_b32 v5, off, s[8:11], s3 offset:7 ; encoding: [0x03,0x00,0x05,0xc4,0x05,0x10,0x80,0x00,0x00,0x07,0x00,0x00]
+
+buffer_load_b32 v5, off, s[8:11], s3 offset:8388607 th:TH_LOAD_NT_HT scope:SCOPE_DEV
+// GFX1250: buffer_load_b32 v5, off, s[8:11], s3 offset:8388607 th:TH_LOAD_NT_HT scope:SCOPE_DEV ; encoding: [0x03,0x00,0x05,0xc4,0x05,0x10,0xe8,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_b32 v5, off, s[8:11], s3 offset:8388607 th:TH_LOAD_BYPASS scope:SCOPE_SYS
+// GFX1250: buffer_load_b32 v5, off, s[8:11], s3 offset:8388607 th:TH_LOAD_BYPASS scope:SCOPE_SYS ; encoding: [0x03,0x00,0x05,0xc4,0x05,0x10,0xbc,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_b64 v[6:7], off, s[8:11], s3 offset:8388607
+// GFX1250: buffer_load_b64 v[6:7], off, s[8:11], s3 offset:8388607 ; encoding: [0x03,0x40,0x05,0xc4,0x06,0x10,0x80,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_b64 v[254:255], off, s[8:11], s3 offset:8388607
+// GFX1250: buffer_load_b64 v[254:255], off, s[8:11], s3 offset:8388607 ; encoding: [0x03,0x40,0x05,0xc4,0xfe,0x10,0x80,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_b64 v[6:7], off, s[12:15], s3 offset:8388607
+// GFX1250: buffer_load_b64 v[6:7], off, s[12:15], s3 offset:8388607 ; encoding: [0x03,0x40,0x05,0xc4,0x06,0x18,0x80,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_b64 v[6:7], off, s[96:99], s3 offset:8388607
+// GFX1250: buffer_load_b64 v[6:7], off, s[96:99], s3 offset:8388607 ; encoding: [0x03,0x40,0x05,0xc4,0x06,0xc0,0x80,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_b64 v[6:7], off, s[8:11], s101 offset:8388607
+// GFX1250: buffer_load_b64 v[6:7], off, s[8:11], s101 offset:8388607 ; encoding: [0x65,0x40,0x05,0xc4,0x06,0x10,0x80,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_b64 v[6:7], off, s[8:11], m0 offset:8388607
+// GFX1250: buffer_load_b64 v[6:7], off, s[8:11], m0 offset:8388607 ; encoding: [0x7d,0x40,0x05,0xc4,0x06,0x10,0x80,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_b64 v[6:7], v0, s[8:11], s3 idxen offset:8388607
+// GFX1250: buffer_load_b64 v[6:7], v0, s[8:11], s3 idxen offset:8388607 ; encoding: [0x03,0x40,0x05,0xc4,0x06,0x10,0x80,0x80,0x00,0xff,0xff,0x7f]
+
+buffer_load_b64 v[6:7], v0, s[8:11], s3 offen offset:8388607
+// GFX1250: buffer_load_b64 v[6:7], v0, s[8:11], s3 offen offset:8388607 ; encoding: [0x03,0x40,0x05,0xc4,0x06,0x10,0x80,0x40,0x00,0xff,0xff,0x7f]
+
+buffer_load_b64 v[6:7], off, s[8:11], s3
+// GFX1250: buffer_load_b64 v[6:7], off, s[8:11], s3 ; encoding: [0x03,0x40,0x05,0xc4,0x06,0x10,0x80,0x00,0x00,0x00,0x00,0x00]
+
+buffer_load_b64 v[6:7], off, s[8:11], s3 offset:0
+// GFX1250: buffer_load_b64 v[6:7], off, s[8:11], s3 ; encoding: [0x03,0x40,0x05,0xc4,0x06,0x10,0x80,0x00,0x00,0x00,0x00,0x00]
+
+buffer_load_b64 v[6:7], off, s[8:11], s3 offset:7
+// GFX1250: buffer_load_b64 v[6:7], off, s[8:11], s3 offset:7 ; encoding: [0x03,0x40,0x05,0xc4,0x06,0x10,0x80,0x00,0x00,0x07,0x00,0x00]
+
+buffer_load_b64 v[6:7], off, s[8:11], s3 offset:8388607 th:TH_LOAD_NT_HT scope:SCOPE_DEV
+// GFX1250: buffer_load_b64 v[6:7], off, s[8:11], s3 offset:8388607 th:TH_LOAD_NT_HT scope:SCOPE_DEV ; encoding: [0x03,0x40,0x05,0xc4,0x06,0x10,0xe8,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_b64 v[6:7], off, s[8:11], s3 offset:8388607 th:TH_LOAD_BYPASS scope:SCOPE_SYS
+// GFX1250: buffer_load_b64 v[6:7], off, s[8:11], s3 offset:8388607 th:TH_LOAD_BYPASS scope:SCOPE_SYS ; encoding: [0x03,0x40,0x05,0xc4,0x06,0x10,0xbc,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_b96 v[6:8], off, s[8:11], s3 offset:8388607
+// GFX1250: buffer_load_b96 v[6:8], off, s[8:11], s3 offset:8388607 ; encoding: [0x03,0x80,0x05,0xc4,0x06,0x10,0x80,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_b96 v[252:254], off, s[8:11], s3 offset:8388607
+// GFX1250: buffer_load_b96 v[252:254], off, s[8:11], s3 offset:8388607 ; encoding: [0x03,0x80,0x05,0xc4,0xfc,0x10,0x80,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_b96 v[6:8], off, s[12:15], s3 offset:8388607
+// GFX1250: buffer_load_b96 v[6:8], off, s[12:15], s3 offset:8388607 ; encoding: [0x03,0x80,0x05,0xc4,0x06,0x18,0x80,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_b96 v[6:8], off, s[96:99], s3 offset:8388607
+// GFX1250: buffer_load_b96 v[6:8], off, s[96:99], s3 offset:8388607 ; encoding: [0x03,0x80,0x05,0xc4,0x06,0xc0,0x80,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_b96 v[6:8], off, s[8:11], s101 offset:8388607
+// GFX1250: buffer_load_b96 v[6:8], off, s[8:11], s101 offset:8388607 ; encoding: [0x65,0x80,0x05,0xc4,0x06,0x10,0x80,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_b96 v[6:8], off, s[8:11], m0 offset:8388607
+// GFX1250: buffer_load_b96 v[6:8], off, s[8:11], m0 offset:8388607 ; encoding: [0x7d,0x80,0x05,0xc4,0x06,0x10,0x80,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_b96 v[6:8], v0, s[8:11], s3 idxen offset:8388607
+// GFX1250: buffer_load_b96 v[6:8], v0, s[8:11], s3 idxen offset:8388607 ; encoding: [0x03,0x80,0x05,0xc4,0x06,0x10,0x80,0x80,0x00,0xff,0xff,0x7f]
+
+buffer_load_b96 v[6:8], v0, s[8:11], s3 offen offset:8388607
+// GFX1250: buffer_load_b96 v[6:8], v0, s[8:11], s3 offen offset:8388607 ; encoding: [0x03,0x80,0x05,0xc4,0x06,0x10,0x80,0x40,0x00,0xff,0xff,0x7f]
+
+buffer_load_b96 v[6:8], off, s[8:11], s3
+// GFX1250: buffer_load_b96 v[6:8], off, s[8:11], s3 ; encoding: [0x03,0x80,0x05,0xc4,0x06,0x10,0x80,0x00,0x00,0x00,0x00,0x00]
+
+buffer_load_b96 v[6:8], off, s[8:11], s3 offset:0
+// GFX1250: buffer_load_b96 v[6:8], off, s[8:11], s3 ; encoding: [0x03,0x80,0x05,0xc4,0x06,0x10,0x80,0x00,0x00,0x00,0x00,0x00]
+
+buffer_load_b96 v[6:8], off, s[8:11], s3 offset:7
+// GFX1250: buffer_load_b96 v[6:8], off, s[8:11], s3 offset:7 ; encoding: [0x03,0x80,0x05,0xc4,0x06,0x10,0x80,0x00,0x00,0x07,0x00,0x00]
+
+buffer_load_b96 v[6:8], off, s[8:11], s3 offset:8388607 th:TH_LOAD_NT_HT scope:SCOPE_DEV
+// GFX1250: buffer_load_b96 v[6:8], off, s[8:11], s3 offset:8388607 th:TH_LOAD_NT_HT scope:SCOPE_DEV ; encoding: [0x03,0x80,0x05,0xc4,0x06,0x10,0xe8,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_b96 v[6:8], off, s[8:11], s3 offset:8388607 th:TH_LOAD_BYPASS scope:SCOPE_SYS
+// GFX1250: buffer_load_b96 v[6:8], off, s[8:11], s3 offset:8388607 th:TH_LOAD_BYPASS scope:SCOPE_SYS ; encoding: [0x03,0x80,0x05,0xc4,0x06,0x10,0xbc,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_b128 v[6:9], off, s[8:11], s3 offset:8388607
+// GFX1250: buffer_load_b128 v[6:9], off, s[8:11], s3 offset:8388607 ; encoding: [0x03,0xc0,0x05,0xc4,0x06,0x10,0x80,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_b128 v[252:255], off, s[8:11], s3 offset:8388607
+// GFX1250: buffer_load_b128 v[252:255], off, s[8:11], s3 offset:8388607 ; encoding: [0x03,0xc0,0x05,0xc4,0xfc,0x10,0x80,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_b128 v[6:9], off, s[12:15], s3 offset:8388607
+// GFX1250: buffer_load_b128 v[6:9], off, s[12:15], s3 offset:8388607 ; encoding: [0x03,0xc0,0x05,0xc4,0x06,0x18,0x80,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_b128 v[6:9], off, s[96:99], s3 offset:8388607
+// GFX1250: buffer_load_b128 v[6:9], off, s[96:99], s3 offset:8388607 ; encoding: [0x03,0xc0,0x05,0xc4,0x06,0xc0,0x80,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_b128 v[6:9], off, s[8:11], s101 offset:8388607
+// GFX1250: buffer_load_b128 v[6:9], off, s[8:11], s101 offset:8388607 ; encoding: [0x65,0xc0,0x05,0xc4,0x06,0x10,0x80,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_b128 v[6:9], off, s[8:11], m0 offset:8388607
+// GFX1250: buffer_load_b128 v[6:9], off, s[8:11], m0 offset:8388607 ; encoding: [0x7d,0xc0,0x05,0xc4,0x06,0x10,0x80,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_b128 v[6:9], v0, s[8:11], s3 idxen offset:8388607
+// GFX1250: buffer_load_b128 v[6:9], v0, s[8:11], s3 idxen offset:8388607 ; encoding: [0x03,0xc0,0x05,0xc4,0x06,0x10,0x80,0x80,0x00,0xff,0xff,0x7f]
+
+buffer_load_b128 v[6:9], v0, s[8:11], s3 offen offset:8388607
+// GFX1250: buffer_load_b128 v[6:9], v0, s[8:11], s3 offen offset:8388607 ; encoding: [0x03,0xc0,0x05,0xc4,0x06,0x10,0x80,0x40,0x00,0xff,0xff,0x7f]
+
+buffer_load_b128 v[6:9], off, s[8:11], s3
+// GFX1250: buffer_load_b128 v[6:9], off, s[8:11], s3 ; encoding: [0x03,0xc0,0x05,0xc4,0x06,0x10,0x80,0x00,0x00,0x00,0x00,0x00]
+
+buffer_load_b128 v[6:9], off, s[8:11], s3 offset:0
+// GFX1250: buffer_load_b128 v[6:9], off, s[8:11], s3 ; encoding: [0x03,0xc0,0x05,0xc4,0x06,0x10,0x80,0x00,0x00,0x00,0x00,0x00]
+
+buffer_load_b128 v[6:9], off, s[8:11], s3 offset:7
+// GFX1250: buffer_load_b128 v[6:9], off, s[8:11], s3 offset:7 ; encoding: [0x03,0xc0,0x05,0xc4,0x06,0x10,0x80,0x00,0x00,0x07,0x00,0x00]
+
+buffer_load_b128 v[6:9], off, s[8:11], s3 offset:8388607 th:TH_LOAD_NT_HT scope:SCOPE_DEV
+// GFX1250: buffer_load_b128 v[6:9], off, s[8:11], s3 offset:8388607 th:TH_LOAD_NT_HT scope:SCOPE_DEV ; encoding: [0x03,0xc0,0x05,0xc4,0x06,0x10,0xe8,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_b128 v[6:9], off, s[8:11], s3 offset:8388607 th:TH_LOAD_BYPASS scope:SCOPE_SYS
+// GFX1250: buffer_load_b128 v[6:9], off, s[8:11], s3 offset:8388607 th:TH_LOAD_BYPASS scope:SCOPE_SYS ; encoding: [0x03,0xc0,0x05,0xc4,0x06,0x10,0xbc,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_d16_b16 v5, off, s[8:11], s3 offset:8388607
+// GFX1250: buffer_load_d16_b16 v5, off, s[8:11], s3 offset:8388607 ; encoding: [0x03,0x00,0x08,0xc4,0x05,0x10,0x80,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_d16_b16 v255, off, s[8:11], s3 offset:8388607
+// GFX1250: buffer_load_d16_b16 v255, off, s[8:11], s3 offset:8388607 ; encoding: [0x03,0x00,0x08,0xc4,0xff,0x10,0x80,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_d16_b16 v5, off, s[12:15], s3 offset:8388607
+// GFX1250: buffer_load_d16_b16 v5, off, s[12:15], s3 offset:8388607 ; encoding: [0x03,0x00,0x08,0xc4,0x05,0x18,0x80,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_d16_b16 v5, off, s[96:99], s3 offset:8388607
+// GFX1250: buffer_load_d16_b16 v5, off, s[96:99], s3 offset:8388607 ; encoding: [0x03,0x00,0x08,0xc4,0x05,0xc0,0x80,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_d16_b16 v5, off, s[8:11], s101 offset:8388607
+// GFX1250: buffer_load_d16_b16 v5, off, s[8:11], s101 offset:8388607 ; encoding: [0x65,0x00,0x08,0xc4,0x05,0x10,0x80,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_d16_b16 v5, off, s[8:11], m0 offset:8388607
+// GFX1250: buffer_load_d16_b16 v5, off, s[8:11], m0 offset:8388607 ; encoding: [0x7d,0x00,0x08,0xc4,0x05,0x10,0x80,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_d16_b16 v5, v0, s[8:11], s3 idxen offset:8388607
+// GFX1250: buffer_load_d16_b16 v5, v0, s[8:11], s3 idxen offset:8388607 ; encoding: [0x03,0x00,0x08,0xc4,0x05,0x10,0x80,0x80,0x00,0xff,0xff,0x7f]
+
+buffer_load_d16_b16 v5, v0, s[8:11], s3 offen offset:8388607
+// GFX1250: buffer_load_d16_b16 v5, v0, s[8:11], s3 offen offset:8388607 ; encoding: [0x03,0x00,0x08,0xc4,0x05,0x10,0x80,0x40,0x00,0xff,0xff,0x7f]
+
+buffer_load_d16_b16 v5, off, s[8:11], s3
+// GFX1250: buffer_load_d16_b16 v5, off, s[8:11], s3 ; encoding: [0x03,0x00,0x08,0xc4,0x05,0x10,0x80,0x00,0x00,0x00,0x00,0x00]
+
+buffer_load_d16_b16 v5, off, s[8:11], s3 offset:0
+// GFX1250: buffer_load_d16_b16 v5, off, s[8:11], s3 ; encoding: [0x03,0x00,0x08,0xc4,0x05,0x10,0x80,0x00,0x00,0x00,0x00,0x00]
+
+buffer_load_d16_b16 v5, off, s[8:11], s3 offset:7
+// GFX1250: buffer_load_d16_b16 v5, off, s[8:11], s3 offset:7 ; encoding: [0x03,0x00,0x08,0xc4,0x05,0x10,0x80,0x00,0x00,0x07,0x00,0x00]
+
+buffer_load_d16_b16 v5, off, s[8:11], s3 offset:8388607 th:TH_LOAD_NT_HT scope:SCOPE_DEV
+// GFX1250: buffer_load_d16_b16 v5, off, s[8:11], s3 offset:8388607 th:TH_LOAD_NT_HT scope:SCOPE_DEV ; encoding: [0x03,0x00,0x08,0xc4,0x05,0x10,0xe8,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_d16_b16 v5, off, s[8:11], s3 offset:8388607 th:TH_LOAD_BYPASS scope:SCOPE_SYS
+// GFX1250: buffer_load_d16_b16 v5, off, s[8:11], s3 offset:8388607 th:TH_LOAD_BYPASS scope:SCOPE_SYS ; encoding: [0x03,0x00,0x08,0xc4,0x05,0x10,0xbc,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_d16_hi_b16 v5, off, s[8:11], s3 offset:8388607
+// GFX1250: buffer_load_d16_hi_b16 v5, off, s[8:11], s3 offset:8388607 ; encoding: [0x03,0xc0,0x08,0xc4,0x05,0x10,0x80,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_d16_hi_b16 v255, off, s[8:11], s3 offset:8388607
+// GFX1250: buffer_load_d16_hi_b16 v255, off, s[8:11], s3 offset:8388607 ; encoding: [0x03,0xc0,0x08,0xc4,0xff,0x10,0x80,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_d16_hi_b16 v5, off, s[12:15], s3 offset:8388607
+// GFX1250: buffer_load_d16_hi_b16 v5, off, s[12:15], s3 offset:8388607 ; encoding: [0x03,0xc0,0x08,0xc4,0x05,0x18,0x80,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_d16_hi_b16 v5, off, s[96:99], s3 offset:8388607
+// GFX1250: buffer_load_d16_hi_b16 v5, off, s[96:99], s3 offset:8388607 ; encoding: [0x03,0xc0,0x08,0xc4,0x05,0xc0,0x80,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_d16_hi_b16 v5, off, s[8:11], s101 offset:8388607
+// GFX1250: buffer_load_d16_hi_b16 v5, off, s[8:11], s101 offset:8388607 ; encoding: [0x65,0xc0,0x08,0xc4,0x05,0x10,0x80,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_d16_hi_b16 v5, off, s[8:11], m0 offset:8388607
+// GFX1250: buffer_load_d16_hi_b16 v5, off, s[8:11], m0 offset:8388607 ; encoding: [0x7d,0xc0,0x08,0xc4,0x05,0x10,0x80,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_d16_hi_b16 v5, v0, s[8:11], s3 idxen offset:8388607
+// GFX1250: buffer_load_d16_hi_b16 v5, v0, s[8:11], s3 idxen offset:8388607 ; encoding: [0x03,0xc0,0x08,0xc4,0x05,0x10,0x80,0x80,0x00,0xff,0xff,0x7f]
+
+buffer_load_d16_hi_b16 v5, v0, s[8:11], s3 offen offset:8388607
+// GFX1250: buffer_load_d16_hi_b16 v5, v0, s[8:11], s3 offen offset:8388607 ; encoding: [0x03,0xc0,0x08,0xc4,0x05,0x10,0x80,0x40,0x00,0xff,0xff,0x7f]
+
+buffer_load_d16_hi_b16 v5, off, s[8:11], s3
+// GFX1250: buffer_load_d16_hi_b16 v5, off, s[8:11], s3 ; encoding: [0x03,0xc0,0x08,0xc4,0x05,0x10,0x80,0x00,0x00,0x00,0x00,0x00]
+
+buffer_load_d16_hi_b16 v5, off, s[8:11], s3 offset:0
+// GFX1250: buffer_load_d16_hi_b16 v5, off, s[8:11], s3 ; encoding: [0x03,0xc0,0x08,0xc4,0x05,0x10,0x80,0x00,0x00,0x00,0x00,0x00]
+
+buffer_load_d16_hi_b16 v5, off, s[8:11], s3 offset:7
+// GFX1250: buffer_load_d16_hi_b16 v5, off, s[8:11], s3 offset:7 ; encoding: [0x03,0xc0,0x08,0xc4,0x05,0x10,0x80,0x00,0x00,0x07,0x00,0x00]
+
+buffer_load_d16_hi_b16 v5, off, s[8:11], s3 offset:8388607 th:TH_LOAD_NT_HT scope:SCOPE_DEV
+// GFX1250: buffer_load_d16_hi_b16 v5, off, s[8:11], s3 offset:8388607 th:TH_LOAD_NT_HT scope:SCOPE_DEV ; encoding: [0x03,0xc0,0x08,0xc4,0x05,0x10,0xe8,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_d16_hi_b16 v5, off, s[8:11], s3 offset:8388607 th:TH_LOAD_BYPASS scope:SCOPE_SYS
+// GFX1250: buffer_load_d16_hi_b16 v5, off, s[8:11], s3 offset:8388607 th:TH_LOAD_BYPASS scope:SCOPE_SYS ; encoding: [0x03,0xc0,0x08,0xc4,0x05,0x10,0xbc,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_d16_hi_i8 v5, off, s[8:11], s3 offset:8388607
+// GFX1250: buffer_load_d16_hi_i8 v5, off, s[8:11], s3 offset:8388607 ; encoding: [0x03,0x80,0x08,0xc4,0x05,0x10,0x80,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_d16_hi_i8 v255, off, s[8:11], s3 offset:8388607
+// GFX1250: buffer_load_d16_hi_i8 v255, off, s[8:11], s3 offset:8388607 ; encoding: [0x03,0x80,0x08,0xc4,0xff,0x10,0x80,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_d16_hi_i8 v5, off, s[12:15], s3 offset:8388607
+// GFX1250: buffer_load_d16_hi_i8 v5, off, s[12:15], s3 offset:8388607 ; encoding: [0x03,0x80,0x08,0xc4,0x05,0x18,0x80,0x00,0x00,0xff,0xff,0x7f]
+
+buffer_load_d16_hi_i8 v5, off, s[96:99], s...
[truncated]

@rampitec rampitec merged commit c310306 into main Aug 6, 2025
14 checks passed
@rampitec rampitec deleted the users/rampitec/08-06-_amdgpu_add_more_gfx1250_mc_tests._nfc branch August 6, 2025 21:38
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

@shiltian shiltian shiltian approved these changes

@changpeng changpeng Awaiting requested review from changpeng

Assignees

No one assigned

Labels

backend:AMDGPU llvm:mc Machine (object) code

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

4 participants

@rampitec @llvmbot @shiltian