This is an archive of the discontinued LLVM Phabricator instance.

Differential D94553

[NFC] Disallow unused prefixes in MC/AMDGPU
ClosedPublic

Authored by mtrofin on Jan 12 2021, 2:07 PM.

Details

Reviewers
arsenm
Commits
rG55f2eeebc96e: [NFC] Disallow unused prefixes in MC/AMDGPU
Summary

1 out of 2 patches.

Diff Detail

Event Timeline

mtrofin created this revision.Jan 12 2021, 2:07 PM
Herald added subscribers: kerbowa, t-tye, tpr and 5 others. · View Herald TranscriptJan 12 2021, 2:07 PM
mtrofin requested review of this revision.Jan 12 2021, 2:07 PM
Herald added a project: Restricted Project. · View Herald TranscriptJan 12 2021, 2:07 PM
Herald added subscribers: llvm-commits, wdng. · View Herald Transcript
arsenm accepted this revision.Jan 12 2021, 2:30 PM
This revision is now accepted and ready to land.Jan 12 2021, 2:30 PM
This revision was landed with ongoing or failed builds.Jan 12 2021, 2:31 PM
Closed by commit rG55f2eeebc96e: [NFC] Disallow unused prefixes in MC/AMDGPU (authored by mtrofin). · Explain Why
This revision was automatically updated to reflect the committed changes.
mtrofin added a commit: rG55f2eeebc96e: [NFC] Disallow unused prefixes in MC/AMDGPU.
Harbormaster completed remote builds in B84919: Diff 316235.Jan 12 2021, 2:52 PM

Revision Contents

Diff 316242

llvm/test/MC/AMDGPU/add-sub-no-carry.s

// RUN: llvm-mc -arch=amdgcn -mcpu=gfx900 -show-encoding %s | FileCheck -check-prefixes=GFX9 %s // RUN: llvm-mc -arch=amdgcn -mcpu=gfx900 -show-encoding %s | FileCheck --check-prefix=GFX9 %s
// RUN: not llvm-mc -arch=amdgcn -mcpu=fiji %s 2>&1 | FileCheck -check-prefixes=ERR-VI,ERR-SICIVI --implicit-check-not=error: %s // RUN: not llvm-mc -arch=amdgcn -mcpu=fiji %s 2>&1 | FileCheck --check-prefix=ERR-VI --implicit-check-not=error: %s
// RUN: not llvm-mc -arch=amdgcn -mcpu=bonaire %s 2>&1 | FileCheck -check-prefixes=ERR-SICI,ERR-SICIVI --implicit-check-not=error: %s // RUN: not llvm-mc -arch=amdgcn -mcpu=bonaire %s 2>&1 | FileCheck --check-prefix=ERR-SICI --implicit-check-not=error: %s
// FIXME: pre-gfx9 errors should be more useful // FIXME: pre-gfx9 errors should be more useful
v_add_u32 v1, v2, v3 v_add_u32 v1, v2, v3
// GFX9: v_add_u32_e32 v1, v2, v3 ; encoding: [0x02,0x07,0x02,0x68] // GFX9: v_add_u32_e32 v1, v2, v3 ; encoding: [0x02,0x07,0x02,0x68]
// ERR-SICI: error: instruction not supported on this GPU // ERR-SICI: error: instruction not supported on this GPU
// ERR-VI: error: operands are not valid for this GPU or mode // ERR-VI: error: operands are not valid for this GPU or mode
▲ Show 20 Lines • Show All 103 Lines • Show Last 20 Lines

llvm/test/MC/AMDGPU/buf-fmt-d16-packed.s

// RUN: llvm-mc -arch=amdgcn -mcpu=gfx810 -show-encoding %s | FileCheck -check-prefix=GCN -check-prefix=PACKED %s // RUN: llvm-mc -arch=amdgcn -mcpu=gfx810 -show-encoding %s | FileCheck -check-prefix=PACKED %s
// RUN: llvm-mc -arch=amdgcn -mcpu=gfx900 -show-encoding %s | FileCheck -check-prefix=GCN -check-prefix=PACKED %s // RUN: llvm-mc -arch=amdgcn -mcpu=gfx900 -show-encoding %s | FileCheck -check-prefix=PACKED %s
// RUN: not llvm-mc -arch=amdgcn -mcpu=fiji 2>&1 %s | FileCheck -check-prefix=UNPACKED-ERR -check-prefix=GCN-ERR --implicit-check-not=error: %s // RUN: not llvm-mc -arch=amdgcn -mcpu=fiji 2>&1 %s | FileCheck -check-prefix=UNPACKED-ERR --implicit-check-not=error: %s
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
// Buffer Format Instructions. // Buffer Format Instructions.
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
buffer_load_format_d16_x v1, off, s[4:7], s1 buffer_load_format_d16_x v1, off, s[4:7], s1
// PACKED: buffer_load_format_d16_x v1, off, s[4:7], s1 ; encoding: [0x00,0x00,0x20,0xe0,0x00,0x01,0x01,0x01] // PACKED: buffer_load_format_d16_x v1, off, s[4:7], s1 ; encoding: [0x00,0x00,0x20,0xe0,0x00,0x01,0x01,0x01]
▲ Show 20 Lines • Show All 62 Lines • Show Last 20 Lines

llvm/test/MC/AMDGPU/buf-fmt-d16-unpacked.s

// RUN: llvm-mc -arch=amdgcn -mcpu=fiji -show-encoding %s | FileCheck -check-prefix=GCN -check-prefix=UNPACKED %s // RUN: llvm-mc -arch=amdgcn -mcpu=fiji -show-encoding %s | FileCheck -check-prefix=UNPACKED %s
// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx810 2>&1 %s | FileCheck -check-prefix=PACKED-ERR -check-prefix=GCN-ERR --implicit-check-not=error: %s // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx810 2>&1 %s | FileCheck -check-prefix=PACKED-ERR --implicit-check-not=error: %s
// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx900 2>&1 %s | FileCheck -check-prefix=PACKED-ERR -check-prefix=GCN-ERR --implicit-check-not=error: %s // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx900 2>&1 %s | FileCheck -check-prefix=PACKED-ERR --implicit-check-not=error: %s
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
// Buffer Format Instructions. // Buffer Format Instructions.
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
buffer_load_format_d16_x v1, off, s[4:7], s1 buffer_load_format_d16_x v1, off, s[4:7], s1
// UNPACKED: buffer_load_format_d16_x v1, off, s[4:7], s1 ; encoding: [0x00,0x00,0x20,0xe0,0x00,0x01,0x01,0x01] // UNPACKED: buffer_load_format_d16_x v1, off, s[4:7], s1 ; encoding: [0x00,0x00,0x20,0xe0,0x00,0x01,0x01,0x01]
▲ Show 20 Lines • Show All 62 Lines • Show Last 20 Lines

llvm/test/MC/AMDGPU/ds-gfx9.s

// RUN: llvm-mc -arch=amdgcn -mcpu=gfx900 -show-encoding %s | FileCheck -check-prefix=GFX9 -check-prefix=GCN %s // RUN: llvm-mc -arch=amdgcn -mcpu=gfx900 -show-encoding %s | FileCheck -check-prefix=GFX9 %s
// RUN: not llvm-mc -arch=amdgcn -mcpu=tonga 2>&1 %s | FileCheck -check-prefix=VI-ERR -check-prefix=GCNERR --implicit-check-not=error: %s // RUN: not llvm-mc -arch=amdgcn -mcpu=tonga 2>&1 %s | FileCheck -check-prefix=VI-ERR --implicit-check-not=error: %s
ds_read_u8_d16 v8, v2 ds_read_u8_d16 v8, v2
// GFX9: ds_read_u8_d16 v8, v2 ; encoding: [0x00,0x00,0xac,0xd8,0x02,0x00,0x00,0x08] // GFX9: ds_read_u8_d16 v8, v2 ; encoding: [0x00,0x00,0xac,0xd8,0x02,0x00,0x00,0x08]
// VI-ERR: error: instruction not supported on this GPU // VI-ERR: error: instruction not supported on this GPU
ds_read_u8_d16_hi v8, v2 ds_read_u8_d16_hi v8, v2
// GFX9: ds_read_u8_d16_hi v8, v2 ; encoding: [0x00,0x00,0xae,0xd8,0x02,0x00,0x00,0x08] // GFX9: ds_read_u8_d16_hi v8, v2 ; encoding: [0x00,0x00,0xae,0xd8,0x02,0x00,0x00,0x08]
// VI-ERR: error: instruction not supported on this GPU // VI-ERR: error: instruction not supported on this GPU
Show All 32 Lines

llvm/test/MC/AMDGPU/ds.s

// RUN: not llvm-mc -arch=amdgcn -show-encoding %s | FileCheck %s --check-prefix=SI --check-prefix=SICI // RUN: not llvm-mc -arch=amdgcn -show-encoding %s | FileCheck %s --check-prefix=SICI
// RUN: not llvm-mc -arch=amdgcn -mcpu=tahiti -show-encoding %s | FileCheck %s --check-prefix=SI --check-prefix=SICI // RUN: not llvm-mc -arch=amdgcn -mcpu=tahiti -show-encoding %s | FileCheck %s --check-prefix=SICI
// RUN: not llvm-mc -arch=amdgcn -mcpu=bonaire -show-encoding %s | FileCheck %s --check-prefix=CI --check-prefix=SICI // RUN: not llvm-mc -arch=amdgcn -mcpu=bonaire -show-encoding %s | FileCheck %s --check-prefixes=CI,SICI
// RUN: llvm-mc -arch=amdgcn -mcpu=tonga -show-encoding %s | FileCheck %s --check-prefix=VI // RUN: llvm-mc -arch=amdgcn -mcpu=tonga -show-encoding %s | FileCheck %s --check-prefix=VI
// RUN: not llvm-mc -arch=amdgcn %s 2>&1 | FileCheck %s --check-prefix=NOSI --check-prefix=NOSICI --implicit-check-not=error: // RUN: not llvm-mc -arch=amdgcn %s 2>&1 | FileCheck %s --check-prefixes=NOSI,NOSICI --implicit-check-not=error:
// RUN: not llvm-mc -arch=amdgcn -mcpu=tahiti %s 2>&1 | FileCheck %s --check-prefix=NOSI --check-prefix=NOSICI --implicit-check-not=error: // RUN: not llvm-mc -arch=amdgcn -mcpu=tahiti %s 2>&1 | FileCheck %s --check-prefixes=NOSI,NOSICI --implicit-check-not=error:
// RUN: not llvm-mc -arch=amdgcn -mcpu=bonaire %s 2>&1 | FileCheck %s --check-prefix=NOCI --check-prefix=NOSICI --implicit-check-not=error: // RUN: not llvm-mc -arch=amdgcn -mcpu=bonaire %s 2>&1 | FileCheck %s --check-prefix=NOSICI --implicit-check-not=error:
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
// Checks for 16-bit Offsets // Checks for 16-bit Offsets
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
ds_add_u32 v2, v4 offset:16 ds_add_u32 v2, v4 offset:16
// SICI: ds_add_u32 v2, v4 offset:16 ; encoding: [0x10,0x00,0x00,0xd8,0x02,0x04,0x00,0x00] // SICI: ds_add_u32 v2, v4 offset:16 ; encoding: [0x10,0x00,0x00,0xd8,0x02,0x04,0x00,0x00]
// VI: ds_add_u32 v2, v4 offset:16 ; encoding: [0x10,0x00,0x00,0xd8,0x02,0x04,0x00,0x00] // VI: ds_add_u32 v2, v4 offset:16 ; encoding: [0x10,0x00,0x00,0xd8,0x02,0x04,0x00,0x00]
▲ Show 20 Lines • Show All 640 Lines • Show Last 20 Lines

llvm/test/MC/AMDGPU/flat-gfx10.s

// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -show-encoding %s | FileCheck --check-prefixes=GFX10,W32 %s // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -show-encoding %s | FileCheck --check-prefix=GFX10 %s
// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 %s 2>&1 | FileCheck --check-prefixes=GFX10-ERR,W32-ERR --implicit-check-not=error: %s // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 %s 2>&1 | FileCheck --check-prefix=GFX10-ERR --implicit-check-not=error: %s
flat_load_dword v1, v[3:4] flat_load_dword v1, v[3:4]
// GFX10: encoding: [0x00,0x00,0x30,0xdc,0x03,0x00,0x7d,0x01] // GFX10: encoding: [0x00,0x00,0x30,0xdc,0x03,0x00,0x7d,0x01]
flat_load_dword v1, v[3:4] offset:-1 flat_load_dword v1, v[3:4] offset:-1
// GFX10-ERR: :28: error: expected a 11-bit unsigned offset // GFX10-ERR: :28: error: expected a 11-bit unsigned offset
flat_load_dword v1, v[3:4] offset:2047 flat_load_dword v1, v[3:4] offset:2047
▲ Show 20 Lines • Show All 109 Lines • Show Last 20 Lines

llvm/test/MC/AMDGPU/flat-global.s

// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx900 -show-encoding %s | FileCheck -check-prefix=GFX9 -check-prefix=GCN %s // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx900 -show-encoding %s | FileCheck --check-prefixes=GFX9,GCN %s
// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx900 2>&1 %s | FileCheck -check-prefix=GFX9-ERR -check-prefix=GCNERR --implicit-check-not=error: %s // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx900 2>&1 %s | FileCheck -check-prefix=GFX9-ERR --implicit-check-not=error: %s
// RUN: not llvm-mc -arch=amdgcn -mcpu=tonga 2>&1 %s | FileCheck -check-prefix=VI-ERR -check-prefix=GCNERR --implicit-check-not=error: %s // RUN: not llvm-mc -arch=amdgcn -mcpu=tonga 2>&1 %s | FileCheck -check-prefix=VI-ERR --implicit-check-not=error: %s
// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -show-encoding %s | FileCheck --check-prefixes=GFX10 %s // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -show-encoding %s | FileCheck --check-prefix=GFX10 %s
// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 %s 2>&1 | FileCheck --check-prefixes=GFX10-ERR --implicit-check-not=error: %s // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 %s 2>&1 | FileCheck --check-prefix=GFX10-ERR --implicit-check-not=error: %s
global_load_ubyte v1, v[3:4], off global_load_ubyte v1, v[3:4], off
// GFX10: encoding: [0x00,0x80,0x20,0xdc,0x03,0x00,0x7d,0x01] // GFX10: encoding: [0x00,0x80,0x20,0xdc,0x03,0x00,0x7d,0x01]
// GFX9: global_load_ubyte v1, v[3:4], off ; encoding: [0x00,0x80,0x40,0xdc,0x03,0x00,0x7f,0x01] // GFX9: global_load_ubyte v1, v[3:4], off ; encoding: [0x00,0x80,0x40,0xdc,0x03,0x00,0x7f,0x01]
// VI-ERR: error: instruction not supported on this GPU // VI-ERR: error: instruction not supported on this GPU
global_load_ubyte v1, v[3:4], off dlc global_load_ubyte v1, v[3:4], off dlc
// GFX10: encoding: [0x00,0x90,0x20,0xdc,0x03,0x00,0x7d,0x01] // GFX10: encoding: [0x00,0x90,0x20,0xdc,0x03,0x00,0x7d,0x01]
▲ Show 20 Lines • Show All 598 Lines • Show Last 20 Lines

llvm/test/MC/AMDGPU/flat-scratch-instructions.s

// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx900 -show-encoding %s | FileCheck -check-prefix=GFX9 -check-prefix=GCN %s // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx900 -show-encoding %s | FileCheck -check-prefix=GFX9 %s
// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx900 2>&1 %s | FileCheck -check-prefix=GFX9-ERR -check-prefix=GCNERR --implicit-check-not=error: %s // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx900 2>&1 %s | FileCheck -check-prefix=GFX9-ERR --implicit-check-not=error: %s
// RUN: not llvm-mc -arch=amdgcn -mcpu=tonga 2>&1 %s | FileCheck -check-prefix=VI-ERR -check-prefix=GCNERR --implicit-check-not=error: %s // RUN: not llvm-mc -arch=amdgcn -mcpu=tonga 2>&1 %s | FileCheck -check-prefix=VI-ERR --implicit-check-not=error: %s
// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -show-encoding %s | FileCheck --check-prefixes=GFX10,W32 %s // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -show-encoding %s | FileCheck --check-prefix=GFX10 %s
// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 %s 2>&1 | FileCheck --check-prefixes=GFX10-ERR,W32-ERR --implicit-check-not=error: %s // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 %s 2>&1 | FileCheck --check-prefix=GFX10-ERR --implicit-check-not=error: %s
scratch_load_ubyte v1, v2, off scratch_load_ubyte v1, v2, off
// GFX10: encoding: [0x00,0x40,0x20,0xdc,0x02,0x00,0x7d,0x01] // GFX10: encoding: [0x00,0x40,0x20,0xdc,0x02,0x00,0x7d,0x01]
// GFX9: scratch_load_ubyte v1, v2, off ; encoding: [0x00,0x40,0x40,0xdc,0x02,0x00,0x7f,0x01] // GFX9: scratch_load_ubyte v1, v2, off ; encoding: [0x00,0x40,0x40,0xdc,0x02,0x00,0x7f,0x01]
// VI-ERR: error: instruction not supported on this GPU // VI-ERR: error: instruction not supported on this GPU
scratch_load_ubyte v1, v2, off dlc scratch_load_ubyte v1, v2, off dlc
// GFX10: encoding: [0x00,0x50,0x20,0xdc,0x02,0x00,0x7d,0x01] // GFX10: encoding: [0x00,0x50,0x20,0xdc,0x02,0x00,0x7d,0x01]
▲ Show 20 Lines • Show All 298 Lines • Show Last 20 Lines

llvm/test/MC/AMDGPU/gfx10_asm_dpp16.s

// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize32,-wavefrontsize64 -show-encoding %s | FileCheck --check-prefixes=GFX10,W32 %s // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize32,-wavefrontsize64 -show-encoding %s | FileCheck --check-prefixes=GFX10,W32 %s
// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=-wavefrontsize32,+wavefrontsize64 -show-encoding %s | FileCheck --check-prefixes=GFX10,W64 %s // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=-wavefrontsize32,+wavefrontsize64 -show-encoding %s | FileCheck --check-prefixes=GFX10,W64 %s
// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize32,-wavefrontsize64 %s 2>&1 | FileCheck --check-prefixes=GFX10-ERR,W32-ERR --implicit-check-not=error: %s // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize32,-wavefrontsize64 %s 2>&1 | FileCheck --check-prefix=W32-ERR --implicit-check-not=error: %s
// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=-wavefrontsize32,+wavefrontsize64 %s 2>&1 | FileCheck --check-prefixes=GFX10-ERR,W64-ERR --implicit-check-not=error: %s // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=-wavefrontsize32,+wavefrontsize64 %s 2>&1 | FileCheck --check-prefix=W64-ERR --implicit-check-not=error: %s
v_mov_b32_dpp v5, v1 quad_perm:[3,2,1,0] row_mask:0x0 bank_mask:0x0 v_mov_b32_dpp v5, v1 quad_perm:[3,2,1,0] row_mask:0x0 bank_mask:0x0
// GFX10: [0xfa,0x02,0x0a,0x7e,0x01,0x1b,0x00,0x00] // GFX10: [0xfa,0x02,0x0a,0x7e,0x01,0x1b,0x00,0x00]
v_mov_b32_dpp v5, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0 v_mov_b32_dpp v5, v1 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
// GFX10: [0xfa,0x02,0x0a,0x7e,0x01,0xe4,0x00,0x00] // GFX10: [0xfa,0x02,0x0a,0x7e,0x01,0xe4,0x00,0x00]
v_mov_b32_dpp v5, v1 row_mirror row_mask:0x0 bank_mask:0x0 v_mov_b32_dpp v5, v1 row_mirror row_mask:0x0 bank_mask:0x0
▲ Show 20 Lines • Show All 682 Lines • Show Last 20 Lines

llvm/test/MC/AMDGPU/gfx10_asm_dpp8.s

// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize32,-wavefrontsize64 -show-encoding %s | FileCheck --check-prefixes=GFX10,W32 %s // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize32,-wavefrontsize64 -show-encoding %s | FileCheck --check-prefixes=GFX10,W32 %s
// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=-wavefrontsize32,+wavefrontsize64 -show-encoding %s | FileCheck --check-prefixes=GFX10,W64 %s // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=-wavefrontsize32,+wavefrontsize64 -show-encoding %s | FileCheck --check-prefixes=GFX10,W64 %s
// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize32,-wavefrontsize64 %s 2>&1 | FileCheck --check-prefixes=GFX10-ERR,W32-ERR --implicit-check-not=error: %s // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize32,-wavefrontsize64 %s 2>&1 | FileCheck --check-prefix=W32-ERR --implicit-check-not=error: %s
// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=-wavefrontsize32,+wavefrontsize64 %s 2>&1 | FileCheck --check-prefixes=GFX10-ERR,W64-ERR --implicit-check-not=error: %s // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=-wavefrontsize32,+wavefrontsize64 %s 2>&1 | FileCheck --check-prefix=W64-ERR --implicit-check-not=error: %s
v_mov_b32_dpp v5, v1 dpp8:[0,1,2,3,4,5,6,7] v_mov_b32_dpp v5, v1 dpp8:[0,1,2,3,4,5,6,7]
// GFX10: encoding: [0xe9,0x02,0x0a,0x7e,0x01,0x88,0xc6,0xfa] // GFX10: encoding: [0xe9,0x02,0x0a,0x7e,0x01,0x88,0xc6,0xfa]
v_cvt_f32_i32_dpp v5, v1 dpp8:[0,1,2,3,4,5,6,7] v_cvt_f32_i32_dpp v5, v1 dpp8:[0,1,2,3,4,5,6,7]
// GFX10: encoding: [0xe9,0x0a,0x0a,0x7e,0x01,0x88,0xc6,0xfa] // GFX10: encoding: [0xe9,0x0a,0x0a,0x7e,0x01,0x88,0xc6,0xfa]
v_cvt_f32_u32_dpp v5, v1 dpp8:[0,1,2,3,4,5,6,7] v_cvt_f32_u32_dpp v5, v1 dpp8:[0,1,2,3,4,5,6,7]
▲ Show 20 Lines • Show All 573 Lines • Show Last 20 Lines

llvm/test/MC/AMDGPU/gfx10_asm_ds.s

// RUN: llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize32,-wavefrontsize64 -show-encoding %s | FileCheck --check-prefixes=GFX10,W32 %s // RUN: llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize32,-wavefrontsize64 -show-encoding %s | FileCheck --check-prefix=GFX10 %s
// RUN: llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=-wavefrontsize32,+wavefrontsize64 -show-encoding %s | FileCheck --check-prefixes=GFX10,W64 %s // RUN: llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=-wavefrontsize32,+wavefrontsize64 -show-encoding %s | FileCheck --check-prefix=GFX10 %s
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
// ENC_DS. // ENC_DS.
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
ds_add_u32 v0, v1 ds_add_u32 v0, v1
// GFX10: encoding: [0x00,0x00,0x00,0xd8,0x00,0x01,0x00,0x00] // GFX10: encoding: [0x00,0x00,0x00,0xd8,0x00,0x01,0x00,0x00]
▲ Show 20 Lines • Show All 6,977 Lines • Show Last 20 Lines

llvm/test/MC/AMDGPU/gfx10_asm_flat.s

// RUN: llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize32,-wavefrontsize64 -show-encoding %s | FileCheck --check-prefixes=GFX10,W32 %s // RUN: llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize32,-wavefrontsize64 -show-encoding %s | FileCheck --check-prefix=GFX10 %s
// RUN: llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=-wavefrontsize32,+wavefrontsize64 -show-encoding %s | FileCheck --check-prefixes=GFX10,W64 %s // RUN: llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=-wavefrontsize32,+wavefrontsize64 -show-encoding %s | FileCheck --check-prefix=GFX10 %s
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
// ENC_FLAT. // ENC_FLAT.
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
flat_load_ubyte v5, v[1:2] flat_load_ubyte v5, v[1:2]
// GFX10: encoding: [0x00,0x00,0x20,0xdc,0x01,0x00,0x7d,0x05] // GFX10: encoding: [0x00,0x00,0x20,0xdc,0x01,0x00,0x7d,0x05]
▲ Show 20 Lines • Show All 118 Lines • Show Last 20 Lines

llvm/test/MC/AMDGPU/gfx10_asm_mubuf.s

// RUN: llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize32,-wavefrontsize64 -show-encoding %s | FileCheck --check-prefixes=GFX10,W32 %s // RUN: llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize32,-wavefrontsize64 -show-encoding %s | FileCheck --check-prefix=GFX10 %s
// RUN: llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=-wavefrontsize32,+wavefrontsize64 -show-encoding %s | FileCheck --check-prefixes=GFX10,W64 %s // RUN: llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=-wavefrontsize32,+wavefrontsize64 -show-encoding %s | FileCheck --check-prefix=GFX10 %s
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
// ENC_MUBUF. // ENC_MUBUF.
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
buffer_load_format_x v5, off, s[8:11], s3 offset:4095 buffer_load_format_x v5, off, s[8:11], s3 offset:4095
// GFX10: encoding: [0xff,0x0f,0x00,0xe0,0x00,0x05,0x02,0x03] // GFX10: encoding: [0xff,0x0f,0x00,0xe0,0x00,0x05,0x02,0x03]
▲ Show 20 Lines • Show All 2,279 Lines • Show Last 20 Lines

llvm/test/MC/AMDGPU/gfx10_asm_smem.s

// RUN: llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize32,-wavefrontsize64 -show-encoding %s | FileCheck --check-prefixes=GFX10,W32 %s // RUN: llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize32,-wavefrontsize64 -show-encoding %s | FileCheck --check-prefix=GFX10 %s
// RUN: llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=-wavefrontsize32,+wavefrontsize64 -show-encoding %s | FileCheck --check-prefixes=GFX10,W64 %s // RUN: llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=-wavefrontsize32,+wavefrontsize64 -show-encoding %s | FileCheck --check-prefix=GFX10 %s
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
// ENC_SMEM. // ENC_SMEM.
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
s_load_dword s5, s[2:3], s0 s_load_dword s5, s[2:3], s0
// GFX10: encoding: [0x41,0x01,0x00,0xf4,0x00,0x00,0x00,0x00] // GFX10: encoding: [0x41,0x01,0x00,0xf4,0x00,0x00,0x00,0x00]
▲ Show 20 Lines • Show All 1,184 Lines • Show Last 20 Lines

llvm/test/MC/AMDGPU/gfx10_asm_sop.s

// RUN: llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize32,-wavefrontsize64 -show-encoding %s | FileCheck --check-prefixes=GFX10,W32 %s // RUN: llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize32,-wavefrontsize64 -show-encoding %s | FileCheck --check-prefix=GFX10 %s
// RUN: llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=-wavefrontsize32,+wavefrontsize64 -show-encoding %s | FileCheck --check-prefixes=GFX10,W64 %s // RUN: llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=-wavefrontsize32,+wavefrontsize64 -show-encoding %s | FileCheck --check-prefix=GFX10 %s
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
// ENC_SOP1. // ENC_SOP1.
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
s_mov_b32 s0, s1 s_mov_b32 s0, s1
// GFX10: encoding: [0x01,0x03,0x80,0xbe] // GFX10: encoding: [0x01,0x03,0x80,0xbe]
▲ Show 20 Lines • Show All 10,062 Lines • Show Last 20 Lines

llvm/test/MC/AMDGPU/gfx10_asm_vop1.s

  • This file is larger than 256 KB, so syntax highlighting is disabled by default.
// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize32,-wavefrontsize64 -show-encoding %s | FileCheck --check-prefixes=GFX10,W32 %s // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize32,-wavefrontsize64 -show-encoding %s | FileCheck --check-prefix=GFX10 %s
// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=-wavefrontsize32,+wavefrontsize64 -show-encoding %s | FileCheck --check-prefixes=GFX10,W64 %s // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=-wavefrontsize32,+wavefrontsize64 -show-encoding %s | FileCheck --check-prefix=GFX10 %s
// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize32,-wavefrontsize64 %s 2>&1 | FileCheck --check-prefixes=GFX10-ERR,W32-ERR --implicit-check-not=error: %s // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize32,-wavefrontsize64 %s 2>&1 | FileCheck --check-prefix=GFX10-ERR --implicit-check-not=error: %s
// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=-wavefrontsize32,+wavefrontsize64 %s 2>&1 | FileCheck --check-prefixes=GFX10-ERR,W64-ERR --implicit-check-not=error: %s // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=-wavefrontsize32,+wavefrontsize64 %s 2>&1 | FileCheck --check-prefix=GFX10-ERR --implicit-check-not=error: %s
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
// ENC_VOP1. // ENC_VOP1.
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
v_nop_e32 v_nop_e32
// GFX10: encoding: [0x00,0x00,0x00,0x7e] // GFX10: encoding: [0x00,0x00,0x00,0x7e]
▲ Show 20 Lines • Show All 16,217 Lines • Show Last 20 Lines

llvm/test/MC/AMDGPU/gfx10_asm_vop2.s

  • This file is larger than 256 KB, so syntax highlighting is disabled by default.
// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize32,-wavefrontsize64 -show-encoding %s | FileCheck --check-prefixes=GFX10,W32 %s // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize32,-wavefrontsize64 -show-encoding %s | FileCheck --check-prefixes=GFX10,W32 %s
// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=-wavefrontsize32,+wavefrontsize64 -show-encoding %s | FileCheck --check-prefixes=GFX10,W64 %s // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=-wavefrontsize32,+wavefrontsize64 -show-encoding %s | FileCheck --check-prefixes=GFX10,W64 %s
// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize32,-wavefrontsize64 %s 2>&1 | FileCheck --check-prefixes=GFX10-ERR,W32-ERR --implicit-check-not=error: %s // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize32,-wavefrontsize64 %s 2>&1 | FileCheck --check-prefix=W32-ERR --implicit-check-not=error: %s
// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=-wavefrontsize32,+wavefrontsize64 %s 2>&1 | FileCheck --check-prefixes=GFX10-ERR,W64-ERR --implicit-check-not=error: %s // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=-wavefrontsize32,+wavefrontsize64 %s 2>&1 | FileCheck --check-prefix=W64-ERR --implicit-check-not=error: %s
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
// ENC_VOP2. // ENC_VOP2.
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
v_cndmask_b32_e32 v5, v1, v2, vcc v_cndmask_b32_e32 v5, v1, v2, vcc
// W64: encoding: [0x01,0x05,0x0a,0x02] // W64: encoding: [0x01,0x05,0x0a,0x02]
// W32-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: operands are not valid for this GPU or mode // W32-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: operands are not valid for this GPU or mode
▲ Show 20 Lines • Show All 13,080 Lines • Show Last 20 Lines

llvm/test/MC/AMDGPU/gfx10_asm_vopc.s

  • This file is larger than 256 KB, so syntax highlighting is disabled by default.
// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize32,-wavefrontsize64 -show-encoding %s | FileCheck --check-prefixes=GFX10,W32 %s // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize32,-wavefrontsize64 -show-encoding %s | FileCheck --check-prefix=W32 %s
// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=-wavefrontsize32,+wavefrontsize64 -show-encoding %s | FileCheck --check-prefixes=GFX10,W64 %s // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=-wavefrontsize32,+wavefrontsize64 -show-encoding %s | FileCheck --check-prefix=W64 %s
// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize32,-wavefrontsize64 %s 2>&1 | FileCheck --check-prefixes=GFX10-ERR,W32-ERR --implicit-check-not=error: %s // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize32,-wavefrontsize64 %s 2>&1 | FileCheck --check-prefix=W32-ERR --implicit-check-not=error: %s
// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=-wavefrontsize32,+wavefrontsize64 %s 2>&1 | FileCheck --check-prefixes=GFX10-ERR,W64-ERR --implicit-check-not=error: %s // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=-wavefrontsize32,+wavefrontsize64 %s 2>&1 | FileCheck --check-prefix=W64-ERR --implicit-check-not=error: %s
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
// ENC_VOPC - v_cmp_* opcodes. // ENC_VOPC - v_cmp_* opcodes.
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
v_cmp_f_f32 vcc, v1, v2 v_cmp_f_f32 vcc, v1, v2
// W64: encoding: [0x01,0x05,0x00,0x7c] // W64: encoding: [0x01,0x05,0x00,0x7c]
// W32-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: operands are not valid for this GPU or mode // W32-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: operands are not valid for this GPU or mode
▲ Show 20 Lines • Show All 11,564 Lines • Show Last 20 Lines

llvm/test/MC/AMDGPU/gfx10_asm_vopc_e64.s

  • This file is larger than 256 KB, so syntax highlighting is disabled by default.
// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize32,-wavefrontsize64 -show-encoding %s | FileCheck --check-prefixes=GFX10,W32 %s // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize32,-wavefrontsize64 -show-encoding %s | FileCheck --check-prefix=W32 %s
// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=-wavefrontsize32,+wavefrontsize64 -show-encoding %s | FileCheck --check-prefixes=GFX10,W64 %s // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=-wavefrontsize32,+wavefrontsize64 -show-encoding %s | FileCheck --check-prefix=W64 %s
// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize32,-wavefrontsize64 %s 2>&1 | FileCheck --check-prefixes=GFX10-ERR,W32-ERR --implicit-check-not=error: %s // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize32,-wavefrontsize64 %s 2>&1 | FileCheck --check-prefix=W32-ERR --implicit-check-not=error: %s
// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=-wavefrontsize32,+wavefrontsize64 %s 2>&1 | FileCheck --check-prefixes=GFX10-ERR,W64-ERR --implicit-check-not=error: %s // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=-wavefrontsize32,+wavefrontsize64 %s 2>&1 | FileCheck --check-prefix=W64-ERR --implicit-check-not=error: %s
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
// ENC_VOPC, VOP3 variant. // ENC_VOPC, VOP3 variant.
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
v_cmp_f_f32_e64 s[10:11], v1, v2 v_cmp_f_f32_e64 s[10:11], v1, v2
// W64: encoding: [0x0a,0x00,0x00,0xd4,0x01,0x05,0x02,0x00] // W64: encoding: [0x0a,0x00,0x00,0xd4,0x01,0x05,0x02,0x00]
// W32-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: invalid operand for instruction // W32-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: invalid operand for instruction
▲ Show 20 Lines • Show All 21,680 Lines • Show Last 20 Lines

llvm/test/MC/AMDGPU/gfx10_asm_vopc_sdwa.s

  • This file is larger than 256 KB, so syntax highlighting is disabled by default.
// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize32,-wavefrontsize64 -show-encoding %s | FileCheck --check-prefixes=GFX10,W32 %s // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize32,-wavefrontsize64 -show-encoding %s | FileCheck --check-prefix=W32 %s
// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=-wavefrontsize32,+wavefrontsize64 -show-encoding %s | FileCheck --check-prefixes=GFX10,W64 %s // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=-wavefrontsize32,+wavefrontsize64 -show-encoding %s | FileCheck --check-prefix=W64 %s
// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize32,-wavefrontsize64 %s 2>&1 | FileCheck --check-prefixes=GFX10-ERR,W32-ERR --implicit-check-not=error: %s // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize32,-wavefrontsize64 %s 2>&1 | FileCheck --check-prefix=W32-ERR --implicit-check-not=error: %s
// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=-wavefrontsize32,+wavefrontsize64 %s 2>&1 | FileCheck --check-prefixes=GFX10-ERR,W64-ERR --implicit-check-not=error: %s // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=-wavefrontsize32,+wavefrontsize64 %s 2>&1 | FileCheck --check-prefix=W64-ERR --implicit-check-not=error: %s
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
// ENC_VOPC, SDWA variant. // ENC_VOPC, SDWA variant.
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
v_cmp_f_f32_sdwa s[6:7], v1, v2 src0_sel:DWORD src1_sel:DWORD v_cmp_f_f32_sdwa s[6:7], v1, v2 src0_sel:DWORD src1_sel:DWORD
// W64: encoding: [0xf9,0x04,0x00,0x7c,0x01,0x86,0x06,0x06] // W64: encoding: [0xf9,0x04,0x00,0x7c,0x01,0x86,0x06,0x06]
// W32-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: invalid operand for instruction // W32-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: invalid operand for instruction
▲ Show 20 Lines • Show All 13,804 Lines • Show Last 20 Lines

llvm/test/MC/AMDGPU/gfx10_asm_vopcx.s

  • This file is larger than 256 KB, so syntax highlighting is disabled by default.
// RUN: llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize32,-wavefrontsize64 -show-encoding %s | FileCheck --check-prefixes=GFX10,W32 %s // RUN: llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize32,-wavefrontsize64 -show-encoding %s | FileCheck --check-prefix=GFX10 %s
// RUN: llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=-wavefrontsize32,+wavefrontsize64 -show-encoding %s | FileCheck --check-prefixes=GFX10,W64 %s // RUN: llvm-mc -arch=amdgcn -mcpu=gfx1010 -mattr=-wavefrontsize32,+wavefrontsize64 -show-encoding %s | FileCheck --check-prefix=GFX10 %s
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
// ENC_VOPC - v_cmpx_* opcodes. // ENC_VOPC - v_cmpx_* opcodes.
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
v_cmpx_f_f32 v1, v2 v_cmpx_f_f32 v1, v2
// GFX10: encoding: [0x01,0x05,0x20,0x7c] // GFX10: encoding: [0x01,0x05,0x20,0x7c]
▲ Show 20 Lines • Show All 15,965 Lines • Show Last 20 Lines

llvm/test/MC/AMDGPU/hsa-metadata-kernel-args-v3.s

// RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx700 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX700 %s // RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx700 -show-encoding %s | FileCheck --check-prefix=CHECK %s
// RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx800 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX800 %s // RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx800 -show-encoding %s | FileCheck --check-prefix=CHECK %s
// RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx900 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX900 %s // RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx900 -show-encoding %s | FileCheck --check-prefix=CHECK %s
; CHECK: .amdgpu_metadata ; CHECK: .amdgpu_metadata
; CHECK: amdhsa.kernels: ; CHECK: amdhsa.kernels:
; CHECK-NEXT: - .args: ; CHECK-NEXT: - .args:
; CHECK-NEXT: - .offset: 1 ; CHECK-NEXT: - .offset: 1
; CHECK-NEXT: .size: 1 ; CHECK-NEXT: .size: 1
; CHECK-NEXT: .type_name: char ; CHECK-NEXT: .type_name: char
; CHECK-NEXT: .value_kind: by_value ; CHECK-NEXT: .value_kind: by_value
▲ Show 20 Lines • Show All 85 Lines • Show Last 20 Lines

llvm/test/MC/AMDGPU/hsa-metadata-kernel-args.s

// RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx700 --amdhsa-code-object-version=2 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX700 %s // RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx700 --amdhsa-code-object-version=2 -show-encoding %s | FileCheck --check-prefix=CHECK %s
// RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx800 --amdhsa-code-object-version=2 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX800 %s // RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx800 --amdhsa-code-object-version=2 -show-encoding %s | FileCheck --check-prefix=CHECK %s
// RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx900 --amdhsa-code-object-version=2 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX900 %s // RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx900 --amdhsa-code-object-version=2 -show-encoding %s | FileCheck --check-prefix=CHECK %s
// The legacy ValueType field should be parsed without error, but not // The legacy ValueType field should be parsed without error, but not
// re-emitted. // re-emitted.
// CHECK: .amd_amdgpu_hsa_metadata // CHECK: .amd_amdgpu_hsa_metadata
// CHECK: Version: [ 1, 0 ] // CHECK: Version: [ 1, 0 ]
// CHECK: Printf: // CHECK: Printf:
// CHECK: - '1:1:4:%d\n' // CHECK: - '1:1:4:%d\n'
▲ Show 20 Lines • Show All 59 Lines • Show Last 20 Lines

llvm/test/MC/AMDGPU/hsa-metadata-kernel-attrs-v3.s

// RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx700 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX700 %s // RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx700 -show-encoding %s | FileCheck --check-prefix=CHECK %s
// RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx800 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX800 %s // RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx800 -show-encoding %s | FileCheck --check-prefix=CHECK %s
// RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx900 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX900 %s // RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx900 -show-encoding %s | FileCheck --check-prefix=CHECK %s
// CHECK: .amdgpu_metadata // CHECK: .amdgpu_metadata
// CHECK: amdhsa.kernels: // CHECK: amdhsa.kernels:
// CHECK: - .group_segment_fixed_size: 16 // CHECK: - .group_segment_fixed_size: 16
// CHECK: .kernarg_segment_align: 64 // CHECK: .kernarg_segment_align: 64
// CHECK: .kernarg_segment_size: 8 // CHECK: .kernarg_segment_size: 8
// CHECK: .language: OpenCL C // CHECK: .language: OpenCL C
// CHECK: .language_version: // CHECK: .language_version:
▲ Show 20 Lines • Show All 57 Lines • Show Last 20 Lines

llvm/test/MC/AMDGPU/hsa-metadata-kernel-attrs.s

// RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx700 --amdhsa-code-object-version=2 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX700 %s // RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx700 --amdhsa-code-object-version=2 -show-encoding %s | FileCheck --check-prefix=CHECK %s
// RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx800 --amdhsa-code-object-version=2 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX800 %s // RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx800 --amdhsa-code-object-version=2 -show-encoding %s | FileCheck --check-prefix=CHECK %s
// RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx900 --amdhsa-code-object-version=2 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX900 %s // RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx900 --amdhsa-code-object-version=2 -show-encoding %s | FileCheck --check-prefix=CHECK %s
// CHECK: .amd_amdgpu_hsa_metadata // CHECK: .amd_amdgpu_hsa_metadata
// CHECK: Version: [ 1, 0 ] // CHECK: Version: [ 1, 0 ]
// CHECK: Printf: // CHECK: Printf:
// CHECK: - '1:1:4:%d\n' // CHECK: - '1:1:4:%d\n'
// CHECK: - '2:1:8:%g\n' // CHECK: - '2:1:8:%g\n'
// CHECK: Kernels: // CHECK: Kernels:
// CHECK: - Name: test_kernel // CHECK: - Name: test_kernel
Show All 21 Lines

llvm/test/MC/AMDGPU/hsa-metadata-kernel-code-props-v3.s

// RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx700 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX700 %s // RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx700 -show-encoding %s | FileCheck --check-prefix=CHECK %s
// RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx800 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX800 %s // RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx800 -show-encoding %s | FileCheck --check-prefix=CHECK %s
// RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx900 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX900 %s // RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx900 -show-encoding %s | FileCheck --check-prefix=CHECK %s
// CHECK: .amdgpu_metadata // CHECK: .amdgpu_metadata
// CHECK: amdhsa.kernels: // CHECK: amdhsa.kernels:
// CHECK: - .group_segment_fixed_size: 24 // CHECK: - .group_segment_fixed_size: 24
// CHECK: .kernarg_segment_align: 16 // CHECK: .kernarg_segment_align: 16
// CHECK: .kernarg_segment_size: 24 // CHECK: .kernarg_segment_size: 24
// CHECK: .max_flat_workgroup_size: 256 // CHECK: .max_flat_workgroup_size: 256
// CHECK: .name: test_kernel // CHECK: .name: test_kernel
Show All 31 Lines

llvm/test/MC/AMDGPU/hsa-metadata-kernel-code-props.s

// RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx700 --amdhsa-code-object-version=2 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX700 %s // RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx700 --amdhsa-code-object-version=2 -show-encoding %s | FileCheck --check-prefix=CHECK %s
// RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx800 --amdhsa-code-object-version=2 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX800 %s // RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx800 --amdhsa-code-object-version=2 -show-encoding %s | FileCheck --check-prefix=CHECK %s
// RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx900 --amdhsa-code-object-version=2 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX900 %s // RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx900 --amdhsa-code-object-version=2 -show-encoding %s | FileCheck --check-prefix=CHECK %s
// CHECK: .amd_amdgpu_hsa_metadata // CHECK: .amd_amdgpu_hsa_metadata
// CHECK: Version: [ 1, 0 ] // CHECK: Version: [ 1, 0 ]
// CHECK: Kernels: // CHECK: Kernels:
// CHECK: - Name: test_kernel // CHECK: - Name: test_kernel
// CHECK: SymbolName: 'test_kernel@kd' // CHECK: SymbolName: 'test_kernel@kd'
// CHECK: CodeProps: // CHECK: CodeProps:
// CHECK: KernargSegmentSize: 24 // CHECK: KernargSegmentSize: 24
Show All 23 Lines

llvm/test/MC/AMDGPU/hsa-metadata-kernel-debug-props.s

// RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx700 --amdhsa-code-object-version=2 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX700 %s // RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx700 --amdhsa-code-object-version=2 -show-encoding %s | FileCheck --check-prefix=CHECK %s
// RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx800 --amdhsa-code-object-version=2 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX800 %s // RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx800 --amdhsa-code-object-version=2 -show-encoding %s | FileCheck --check-prefix=CHECK %s
// RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx900 --amdhsa-code-object-version=2 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX900 %s // RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx900 --amdhsa-code-object-version=2 -show-encoding %s | FileCheck --check-prefix=CHECK %s
// CHECK: .amd_amdgpu_hsa_metadata // CHECK: .amd_amdgpu_hsa_metadata
// CHECK: Version: [ 1, 0 ] // CHECK: Version: [ 1, 0 ]
// CHECK: Kernels: // CHECK: Kernels:
// CHECK: - Name: test_kernel // CHECK: - Name: test_kernel
// CHECK: SymbolName: 'test_kernel@kd' // CHECK: SymbolName: 'test_kernel@kd'
// CHECK: DebugProps: // CHECK: DebugProps:
// CHECK: DebuggerABIVersion: [ 1, 0 ] // CHECK: DebuggerABIVersion: [ 1, 0 ]
Show All 17 Lines

llvm/test/MC/AMDGPU/hsa-wave-size.s

// RUN: not llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx700 --amdhsa-code-object-version=2 %s | FileCheck --check-prefixes=GCN,GFX7 %s // RUN: not llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx700 --amdhsa-code-object-version=2 %s | FileCheck --check-prefixes=GCN,GFX7 %s
// RUN: not llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx1010 --amdhsa-code-object-version=2 -mattr=+wavefrontsize32,-wavefrontsize64 %s | FileCheck --check-prefixes=GCN,GFX10-W32 %s // RUN: not llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx1010 --amdhsa-code-object-version=2 -mattr=+wavefrontsize32,-wavefrontsize64 %s | FileCheck --check-prefixes=GCN,GFX10-W32 %s
// RUN: not llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx1010 --amdhsa-code-object-version=2 -mattr=-wavefrontsize32,+wavefrontsize64 %s | FileCheck --check-prefixes=GCN,GFX10-W64 %s // RUN: not llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx1010 --amdhsa-code-object-version=2 -mattr=-wavefrontsize32,+wavefrontsize64 %s | FileCheck --check-prefixes=GCN,GFX10-W64 %s
// RUN: not llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx700 --amdhsa-code-object-version=2 %s 2>&1 | FileCheck --check-prefixes=GCN-ERR,GFX7-ERR %s // RUN: not llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx700 --amdhsa-code-object-version=2 %s 2>&1 | FileCheck --check-prefix=GFX7-ERR %s
// RUN: not llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx1010 --amdhsa-code-object-version=2 -mattr=+wavefrontsize32,-wavefrontsize64 %s 2>&1 | FileCheck --check-prefixes=GCN-ERR,GFX10-W32-ERR %s // RUN: not llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx1010 --amdhsa-code-object-version=2 -mattr=+wavefrontsize32,-wavefrontsize64 %s 2>&1 | FileCheck --check-prefix=GFX10-W32-ERR %s
// RUN: not llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx1010 --amdhsa-code-object-version=2 -mattr=-wavefrontsize32,+wavefrontsize64 %s 2>&1 | FileCheck --check-prefixes=GCN-ERR,GFX10-W64-ERR %s // RUN: not llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx1010 --amdhsa-code-object-version=2 -mattr=-wavefrontsize32,+wavefrontsize64 %s 2>&1 | FileCheck --check-prefix=GFX10-W64-ERR %s
// GCN: test0: // GCN: test0:
// GFX7: enable_wavefront_size32 = 0 // GFX7: enable_wavefront_size32 = 0
// GFX7: wavefront_size = 6 // GFX7: wavefront_size = 6
// GFX10-W32: enable_wavefront_size32 = 1 // GFX10-W32: enable_wavefront_size32 = 1
// GFX10-W32: wavefront_size = 5 // GFX10-W32: wavefront_size = 5
// GFX10-W64: enable_wavefront_size32 = 0 // GFX10-W64: enable_wavefront_size32 = 0
// GFX10-W64: wavefront_size = 6 // GFX10-W64: wavefront_size = 6
▲ Show 20 Lines • Show All 50 Lines • Show Last 20 Lines

llvm/test/MC/AMDGPU/regression/bug28165.s

// RUN: llvm-mc -arch=amdgcn -show-encoding %s | FileCheck %s --check-prefix=GCN --check-prefix=SI --check-prefix=SICI // RUN: llvm-mc -arch=amdgcn -show-encoding %s | FileCheck %s --check-prefixes=GCN,SICI
// RUN: llvm-mc -arch=amdgcn -mcpu=tahiti -show-encoding %s | FileCheck %s --check-prefix=GCN --check-prefix=SI --check-prefix=SICI // RUN: llvm-mc -arch=amdgcn -mcpu=tahiti -show-encoding %s | FileCheck %s --check-prefixes=GCN,SICI
// RUN: llvm-mc -arch=amdgcn -mcpu=bonaire -show-encoding %s | FileCheck %s --check-prefix=GCN --check-prefix=SICI --check-prefix=CIVI // RUN: llvm-mc -arch=amdgcn -mcpu=bonaire -show-encoding %s | FileCheck %s --check-prefixes=GCN,SICI
// RUN: llvm-mc -arch=amdgcn -mcpu=tonga -show-encoding %s | FileCheck %s --check-prefix=GCN --check-prefix=CIVI --check-prefix=VI // RUN: llvm-mc -arch=amdgcn -mcpu=tonga -show-encoding %s | FileCheck %s --check-prefixes=GCN,VI
// SICI: v_cmp_eq_f64_e32 vcc, 0.5, v[254:255] ; encoding: [0xf0,0xfc,0x45,0x7c] // SICI: v_cmp_eq_f64_e32 vcc, 0.5, v[254:255] ; encoding: [0xf0,0xfc,0x45,0x7c]
// VI: v_cmp_eq_f64_e32 vcc, 0.5, v[254:255] ; encoding: [0xf0,0xfc,0xc5,0x7c] // VI: v_cmp_eq_f64_e32 vcc, 0.5, v[254:255] ; encoding: [0xf0,0xfc,0xc5,0x7c]
v_cmp_eq_f64 vcc, 0.5, v[254:255] v_cmp_eq_f64 vcc, 0.5, v[254:255]
// GCN: v_cvt_f32_f64_e32 v0, 0.5 ; encoding: [0xf0,0x1e,0x00,0x7e] // GCN: v_cvt_f32_f64_e32 v0, 0.5 ; encoding: [0xf0,0x1e,0x00,0x7e]
v_cvt_f32_f64 v0, 0.5 v_cvt_f32_f64 v0, 0.5

llvm/test/MC/AMDGPU/regression/bug28168.s

// RUN: llvm-mc -arch=amdgcn -mcpu=bonaire -show-encoding %s | FileCheck %s --check-prefix=GCN --check-prefix=SICI --check-prefix=CI // RUN: llvm-mc -arch=amdgcn -mcpu=bonaire -show-encoding %s | FileCheck %s --check-prefix=CI
// RUN: llvm-mc -arch=amdgcn -mcpu=tonga -show-encoding %s | FileCheck %s --check-prefix=GCN --check-prefix=CIVI --check-prefix=VI // RUN: llvm-mc -arch=amdgcn -mcpu=tonga -show-encoding %s | FileCheck %s --check-prefix=VI
v_mqsad_pk_u16_u8 v[2:3], s[0:1], 1, v[254:255] v_mqsad_pk_u16_u8 v[2:3], s[0:1], 1, v[254:255]
// CI: [0x02,0x00,0xe6,0xd2,0x00,0x02,0xf9,0x07] // CI: [0x02,0x00,0xe6,0xd2,0x00,0x02,0xf9,0x07]
// VI: [0x02,0x00,0xe6,0xd1,0x00,0x02,0xf9,0x07] // VI: [0x02,0x00,0xe6,0xd1,0x00,0x02,0xf9,0x07]
v_qsad_pk_u16_u8 v[2:3], v[0:1], 1, s[0:1] v_qsad_pk_u16_u8 v[2:3], v[0:1], 1, s[0:1]
// CI: [0x02,0x00,0xe4,0xd2,0x00,0x03,0x01,0x00] // CI: [0x02,0x00,0xe4,0xd2,0x00,0x03,0x01,0x00]
// VI: [0x02,0x00,0xe5,0xd1,0x00,0x03,0x01,0x00] // VI: [0x02,0x00,0xe5,0xd1,0x00,0x03,0x01,0x00]

llvm/test/MC/AMDGPU/regression/bug28413.s

// RUN: llvm-mc -arch=amdgcn -show-encoding %s | FileCheck %s --check-prefix=GCN --check-prefix=SI --check-prefix=SICI // RUN: llvm-mc -arch=amdgcn -show-encoding %s | FileCheck %s --check-prefixes=GCN,SICI
// RUN: llvm-mc -arch=amdgcn -mcpu=tahiti -show-encoding %s | FileCheck %s --check-prefix=GCN --check-prefix=SI --check-prefix=SICI // RUN: llvm-mc -arch=amdgcn -mcpu=tahiti -show-encoding %s | FileCheck %s --check-prefixes=GCN,SICI
// RUN: llvm-mc -arch=amdgcn -mcpu=bonaire -show-encoding %s | FileCheck %s --check-prefix=GCN --check-prefix=SICI --check-prefix=CIVI // RUN: llvm-mc -arch=amdgcn -mcpu=bonaire -show-encoding %s | FileCheck %s --check-prefixes=GCN,SICI
// RUN: llvm-mc -arch=amdgcn -mcpu=tonga -show-encoding %s | FileCheck %s --check-prefix=GCN --check-prefix=CIVI --check-prefix=VI // RUN: llvm-mc -arch=amdgcn -mcpu=tonga -show-encoding %s | FileCheck %s --check-prefixes=GCN,VI
v_cmp_eq_i32 vcc, 0.5, v0 v_cmp_eq_i32 vcc, 0.5, v0
// SICI: v_cmp_eq_i32_e32 vcc, 0.5, v0 ; encoding: [0xf0,0x00,0x04,0x7d] // SICI: v_cmp_eq_i32_e32 vcc, 0.5, v0 ; encoding: [0xf0,0x00,0x04,0x7d]
// VI: v_cmp_eq_i32_e32 vcc, 0.5, v0 ; encoding: [0xf0,0x00,0x84,0x7d] // VI: v_cmp_eq_i32_e32 vcc, 0.5, v0 ; encoding: [0xf0,0x00,0x84,0x7d]
v_cmpx_eq_u32_e64 s[0:1], -4.0, s0 v_cmpx_eq_u32_e64 s[0:1], -4.0, s0
// SICI: v_cmpx_eq_u32_e64 s[0:1], -4.0, s0 ; encoding: [0x00,0x00,0xa4,0xd1,0xf7,0x00,0x00,0x00] // SICI: v_cmpx_eq_u32_e64 s[0:1], -4.0, s0 ; encoding: [0x00,0x00,0xa4,0xd1,0xf7,0x00,0x00,0x00]
// VI: v_cmpx_eq_u32_e64 s[0:1], -4.0, s0 ; encoding: [0x00,0x00,0xda,0xd0,0xf7,0x00,0x00,0x00] // VI: v_cmpx_eq_u32_e64 s[0:1], -4.0, s0 ; encoding: [0x00,0x00,0xda,0xd0,0xf7,0x00,0x00,0x00]
Show All 14 Lines

llvm/test/MC/AMDGPU/regression/bug28538.s

// RUN: not llvm-mc -arch=amdgcn -mcpu=tonga %s 2>&1 | FileCheck %s --check-prefix=NOCIVI --check-prefix=NOVI --implicit-check-not=error: // RUN: not llvm-mc -arch=amdgcn -mcpu=tonga %s 2>&1 | FileCheck %s --check-prefix=NOVI --implicit-check-not=error:
// RUN: not llvm-mc -arch=amdgcn %s 2>&1 | FileCheck %s --check-prefix=NOSI --check-prefix=NOSICI --implicit-check-not=error: // RUN: not llvm-mc -arch=amdgcn %s 2>&1 | FileCheck %s --check-prefix=NOSICI --implicit-check-not=error:
// RUN: not llvm-mc -arch=amdgcn -mcpu=tahiti %s 2>&1 | FileCheck %s --check-prefix=NOSI --check-prefix=NOSICI --implicit-check-not=error: // RUN: not llvm-mc -arch=amdgcn -mcpu=tahiti %s 2>&1 | FileCheck %s --check-prefix=NOSICI --implicit-check-not=error:
// RUN: not llvm-mc -arch=amdgcn -mcpu=bonaire %s 2>&1 | FileCheck %s --check-prefix=NOSICI --implicit-check-not=error: // RUN: not llvm-mc -arch=amdgcn -mcpu=bonaire %s 2>&1 | FileCheck %s --check-prefix=NOSICI --implicit-check-not=error:
// NOSICI: error: not a valid operand. // NOSICI: error: not a valid operand.
// NOVI: error: failed parsing operand // NOVI: error: failed parsing operand
v_mov_b32 v0, v0 row_bcast:0 v_mov_b32 v0, v0 row_bcast:0
// NOSICI: error: not a valid operand. // NOSICI: error: not a valid operand.
// NOVI: error: failed parsing operand // NOVI: error: failed parsing operand
v_mov_b32 v0, v0 row_bcast:13 v_mov_b32 v0, v0 row_bcast:13