1a6fc489bSMatt Arsenault// RUN: not llvm-mc -triple=amdgcn -mcpu=gfx950 %s 2>&1 | FileCheck -check-prefix=ERR %s 2*e97fb220SMatt Arsenault// RUN: not llvm-mc -triple=amdgcn -mcpu=gfx950 -mattr=+wavefrontsize32,-wavefrontsize64 %s 2>&1 | FileCheck -check-prefix=W32-ERR %s 3a6fc489bSMatt Arsenault 4a6fc489bSMatt Arsenault//===----------------------------------------------------------------------===// 5a6fc489bSMatt Arsenault// v_mfma_f32_32x32x4_xf32 6a6fc489bSMatt Arsenault//===----------------------------------------------------------------------===// 7a6fc489bSMatt Arsenault 8a6fc489bSMatt Arsenaultv_mfma_f32_32x32x4_xf32 a[0:3], v[2:3], v[4:5], a[2:5] 9a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 10a6fc489bSMatt Arsenault 11a6fc489bSMatt Arsenaultv_mfma_f32_32x32x4_xf32 v[0:3], v[0:3], v[0:3], v[0:3] 12a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 13a6fc489bSMatt Arsenault 14a6fc489bSMatt Arsenaultv_mfma_f32_32x32x4_xf32 v[0:3], v[0:3], v[0:3], v[0:3] 15a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 16a6fc489bSMatt Arsenault 17a6fc489bSMatt Arsenaultv_mfma_f32_32x32x4_xf32 a[0:3], a[0:3], a[0:3], a[0:3] 18a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 19a6fc489bSMatt Arsenault 20a6fc489bSMatt Arsenaultv_mfma_f32_32x32x4_xf32 a[0:3], a[0:3], a[0:3], a[0:3] 21a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 22a6fc489bSMatt Arsenault 23a6fc489bSMatt Arsenaultv_mfma_f32_32x32x4_xf32 v[0:3], a[0:3], v[0:3], 1.0 24a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 25a6fc489bSMatt Arsenault 26a6fc489bSMatt Arsenaultv_mfma_f32_32x32x4_xf32 a[0:3], v[0:3], a[0:3], 1.0 27a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 28a6fc489bSMatt Arsenault 29a6fc489bSMatt Arsenaultv_mfma_f32_32x32x4_xf32 v[0:3], v[0:3], v[0:3], v[0:3] blgp:5 30a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 31a6fc489bSMatt Arsenault 32a6fc489bSMatt Arsenaultv_mfma_f32_32x32x4_xf32 a[0:3], a[0:3], a[0:3], a[0:3] blgp:1 33a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 34a6fc489bSMatt Arsenault 35a6fc489bSMatt Arsenaultv_mfma_f32_32x32x4_xf32 a[0:3], a[0:3], a[0:3], a[0:3] cbsz:3 36a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 37a6fc489bSMatt Arsenault 38a6fc489bSMatt Arsenaultv_mfma_f32_32x32x4_xf32 a[0:3], a[0:3], a[0:3], a[0:3] abid:1 39a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 40a6fc489bSMatt Arsenault 41a6fc489bSMatt Arsenaultv_mfma_f32_32x32x4_xf32 a[0:3], a[0:3], a[0:3], a[0:3] cbsz:3 abid:1 42a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 43a6fc489bSMatt Arsenault 44a6fc489bSMatt Arsenaultv_mfma_f32_32x32x4_xf32 a[0:3], v[0:3], v[0:3], a[4:7] 45a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 46a6fc489bSMatt Arsenault 47a6fc489bSMatt Arsenaultv_mfma_f32_32x32x4_xf32 v[0:3], a[0:3], a[0:3], v[4:7] 48a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 49a6fc489bSMatt Arsenault 50a6fc489bSMatt Arsenaultv_mfma_f32_32x32x4_xf32 a[0:3], v[2:3], v[4:5], a[2:5] 51a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 52a6fc489bSMatt Arsenault 53a6fc489bSMatt Arsenaultv_mfma_f32_32x32x4_xf32 v[0:3], v[0:3], v[0:3], v[0:3] 54a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 55a6fc489bSMatt Arsenault 56a6fc489bSMatt Arsenaultv_mfma_f32_32x32x4_xf32 v[0:3], v[0:3], v[0:3], v[0:3] 57a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 58a6fc489bSMatt Arsenault 59a6fc489bSMatt Arsenaultv_mfma_f32_32x32x4_xf32 a[0:3], a[0:3], a[0:3], a[0:3] 60a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 61a6fc489bSMatt Arsenault 62a6fc489bSMatt Arsenaultv_mfma_f32_32x32x4_xf32 a[0:3], a[0:3], a[0:3], a[0:3] 63a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 64a6fc489bSMatt Arsenault 65a6fc489bSMatt Arsenaultv_mfma_f32_32x32x4_xf32 v[0:3], a[0:3], v[0:3], 1.0 66a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 67a6fc489bSMatt Arsenault 68a6fc489bSMatt Arsenaultv_mfma_f32_32x32x4_xf32 a[0:3], v[0:3], a[0:3], 1.0 69a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 70a6fc489bSMatt Arsenault 71a6fc489bSMatt Arsenaultv_mfma_f32_32x32x4_xf32 v[0:3], v[0:3], v[0:3], v[0:3] blgp:5 72a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 73a6fc489bSMatt Arsenault 74a6fc489bSMatt Arsenaultv_mfma_f32_32x32x4_xf32 a[0:3], a[0:3], a[0:3], a[0:3] blgp:1 75a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 76a6fc489bSMatt Arsenault 77a6fc489bSMatt Arsenaultv_mfma_f32_32x32x4_xf32 a[0:3], a[0:3], a[0:3], a[0:3] cbsz:3 78a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 79a6fc489bSMatt Arsenault 80a6fc489bSMatt Arsenaultv_mfma_f32_32x32x4_xf32 a[0:3], a[0:3], a[0:3], a[0:3] abid:1 81a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 82a6fc489bSMatt Arsenault 83a6fc489bSMatt Arsenaultv_mfma_f32_32x32x4_xf32 a[0:3], a[0:3], a[0:3], a[0:3] cbsz:3 abid:1 84a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 85a6fc489bSMatt Arsenault 86a6fc489bSMatt Arsenaultv_mfma_f32_32x32x4_xf32 a[0:3], v[0:3], v[0:3], a[4:7] 87a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 88a6fc489bSMatt Arsenault 89a6fc489bSMatt Arsenaultv_mfma_f32_32x32x4_xf32 v[0:3], a[0:3], a[0:3], v[4:7] 90a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 91a6fc489bSMatt Arsenault 92a6fc489bSMatt Arsenault 93a6fc489bSMatt Arsenault//===----------------------------------------------------------------------===// 94a6fc489bSMatt Arsenault// v_mfma_f32_16x16x8_xf32 95a6fc489bSMatt Arsenault//===----------------------------------------------------------------------===// 96a6fc489bSMatt Arsenault 97a6fc489bSMatt Arsenaultv_mfma_f32_16x16x8_xf32 a[0:3], v[2:3], v[4:5], a[2:5] 98a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 99a6fc489bSMatt Arsenault 100a6fc489bSMatt Arsenaultv_mfma_f32_16x16x8_xf32 v[0:3], v[0:3], v[0:3], v[0:3] 101a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 102a6fc489bSMatt Arsenault 103a6fc489bSMatt Arsenaultv_mfma_f32_16x16x8_xf32 v[0:3], v[0:3], v[0:3], v[0:3] 104a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 105a6fc489bSMatt Arsenault 106a6fc489bSMatt Arsenaultv_mfma_f32_16x16x8_xf32 a[0:3], a[0:3], a[0:3], a[0:3] 107a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 108a6fc489bSMatt Arsenault 109a6fc489bSMatt Arsenaultv_mfma_f32_16x16x8_xf32 a[0:3], a[0:3], a[0:3], a[0:3] 110a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 111a6fc489bSMatt Arsenault 112a6fc489bSMatt Arsenaultv_mfma_f32_16x16x8_xf32 v[0:3], a[0:3], v[0:3], 1.0 113a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 114a6fc489bSMatt Arsenault 115a6fc489bSMatt Arsenaultv_mfma_f32_16x16x8_xf32 a[0:3], v[0:3], a[0:3], 1.0 116a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 117a6fc489bSMatt Arsenault 118a6fc489bSMatt Arsenaultv_mfma_f32_16x16x8_xf32 v[0:3], v[0:3], v[0:3], v[0:3] blgp:5 119a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 120a6fc489bSMatt Arsenault 121a6fc489bSMatt Arsenaultv_mfma_f32_16x16x8_xf32 a[0:3], a[0:3], a[0:3], a[0:3] blgp:1 122a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 123a6fc489bSMatt Arsenault 124a6fc489bSMatt Arsenaultv_mfma_f32_16x16x8_xf32 a[0:3], a[0:3], a[0:3], a[0:3] cbsz:3 125a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 126a6fc489bSMatt Arsenault 127a6fc489bSMatt Arsenaultv_mfma_f32_16x16x8_xf32 a[0:3], a[0:3], a[0:3], a[0:3] abid:1 128a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 129a6fc489bSMatt Arsenault 130a6fc489bSMatt Arsenaultv_mfma_f32_16x16x8_xf32 a[0:3], a[0:3], a[0:3], a[0:3] cbsz:3 abid:1 131a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 132a6fc489bSMatt Arsenault 133a6fc489bSMatt Arsenaultv_mfma_f32_16x16x8_xf32 a[0:3], v[0:3], v[0:3], a[4:7] 134a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 135a6fc489bSMatt Arsenault 136a6fc489bSMatt Arsenaultv_mfma_f32_16x16x8_xf32 v[0:3], a[0:3], a[0:3], v[4:7] 137a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 138a6fc489bSMatt Arsenault 139a6fc489bSMatt Arsenault 140a6fc489bSMatt Arsenaultv_mfma_f32_16x16x8_xf32 a[0:3], v[2:3], v[4:5], a[2:5] 141a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 142a6fc489bSMatt Arsenault 143a6fc489bSMatt Arsenaultv_mfma_f32_16x16x8_xf32 v[0:3], v[0:3], v[0:3], v[0:3] 144a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 145a6fc489bSMatt Arsenault 146a6fc489bSMatt Arsenaultv_mfma_f32_16x16x8_xf32 v[0:3], v[0:3], v[0:3], v[0:3] 147a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 148a6fc489bSMatt Arsenault 149a6fc489bSMatt Arsenaultv_mfma_f32_16x16x8_xf32 a[0:3], a[0:3], a[0:3], a[0:3] 150a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 151a6fc489bSMatt Arsenault 152a6fc489bSMatt Arsenaultv_mfma_f32_16x16x8_xf32 a[0:3], a[0:3], a[0:3], a[0:3] 153a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 154a6fc489bSMatt Arsenault 155a6fc489bSMatt Arsenaultv_mfma_f32_16x16x8_xf32 v[0:3], a[0:3], v[0:3], 1.0 156a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 157a6fc489bSMatt Arsenault 158a6fc489bSMatt Arsenaultv_mfma_f32_16x16x8_xf32 a[0:3], v[0:3], a[0:3], 1.0 159a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 160a6fc489bSMatt Arsenault 161a6fc489bSMatt Arsenaultv_mfma_f32_16x16x8_xf32 v[0:3], v[0:3], v[0:3], v[0:3] blgp:5 162a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 163a6fc489bSMatt Arsenault 164a6fc489bSMatt Arsenaultv_mfma_f32_16x16x8_xf32 a[0:3], a[0:3], a[0:3], a[0:3] blgp:1 165a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 166a6fc489bSMatt Arsenault 167a6fc489bSMatt Arsenaultv_mfma_f32_16x16x8_xf32 a[0:3], a[0:3], a[0:3], a[0:3] cbsz:3 168a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 169a6fc489bSMatt Arsenault 170a6fc489bSMatt Arsenaultv_mfma_f32_16x16x8_xf32 a[0:3], a[0:3], a[0:3], a[0:3] abid:1 171a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 172a6fc489bSMatt Arsenault 173a6fc489bSMatt Arsenaultv_mfma_f32_16x16x8_xf32 a[0:3], a[0:3], a[0:3], a[0:3] cbsz:3 abid:1 174a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 175a6fc489bSMatt Arsenault 176a6fc489bSMatt Arsenaultv_mfma_f32_16x16x8_xf32 a[0:3], v[0:3], v[0:3], a[4:7] 177a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 178a6fc489bSMatt Arsenault 179a6fc489bSMatt Arsenaultv_mfma_f32_16x16x8_xf32 v[0:3], a[0:3], a[0:3], v[4:7] 180a6fc489bSMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU 181*e97fb220SMatt Arsenault 182*e97fb220SMatt Arsenault//===----------------------------------------------------------------------===// 183*e97fb220SMatt Arsenault// ds_read_b64_tr_b4 184*e97fb220SMatt Arsenault//===----------------------------------------------------------------------===// 185*e97fb220SMatt Arsenaultds_read_b64_tr_b4 v[1:2], v0 186*e97fb220SMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned 187*e97fb220SMatt Arsenault// W32-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU 188*e97fb220SMatt Arsenault 189*e97fb220SMatt Arsenaultds_read_b64_tr_b4 v1, v0 190*e97fb220SMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: invalid operand for instruction 191*e97fb220SMatt Arsenault// W32-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU 192*e97fb220SMatt Arsenault 193*e97fb220SMatt Arsenaultds_read_b64_tr_b4 v[0:1], s0 194*e97fb220SMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: invalid operand for instruction 195*e97fb220SMatt Arsenault// W32-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU 196*e97fb220SMatt Arsenault 197*e97fb220SMatt Arsenaultds_read_b64_tr_b4 v[2:3], v2 offset:-64 198*e97fb220SMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: expected a 16-bit unsigned offset 199*e97fb220SMatt Arsenault// W32-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU 200*e97fb220SMatt Arsenault 201*e97fb220SMatt Arsenault//===----------------------------------------------------------------------===// 202*e97fb220SMatt Arsenault//ds_read_b64_tr_b8 203*e97fb220SMatt Arsenault//===----------------------------------------------------------------------===// 204*e97fb220SMatt Arsenaultds_read_b64_tr_b8 v[1:2], v0 205*e97fb220SMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned 206*e97fb220SMatt Arsenault// W32-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU 207*e97fb220SMatt Arsenault 208*e97fb220SMatt Arsenaultds_read_b64_tr_b8 v1, v0 209*e97fb220SMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: invalid operand for instruction 210*e97fb220SMatt Arsenault// W32-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU 211*e97fb220SMatt Arsenault 212*e97fb220SMatt Arsenaultds_read_b64_tr_b8 v[0:1], s0 213*e97fb220SMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: invalid operand for instruction 214*e97fb220SMatt Arsenault// W32-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU 215*e97fb220SMatt Arsenault 216*e97fb220SMatt Arsenaultds_read_b64_tr_b8 v[2:3], v2 offset:-64 217*e97fb220SMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: expected a 16-bit unsigned offset 218*e97fb220SMatt Arsenault// W32-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU 219*e97fb220SMatt Arsenault 220*e97fb220SMatt Arsenault//===----------------------------------------------------------------------===// 221*e97fb220SMatt Arsenault// ds_read_b64_tr_b16 222*e97fb220SMatt Arsenault//===----------------------------------------------------------------------===// 223*e97fb220SMatt Arsenaultds_read_b64_tr_b16 v[1:2], v0 224*e97fb220SMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned 225*e97fb220SMatt Arsenault// W32-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU 226*e97fb220SMatt Arsenault 227*e97fb220SMatt Arsenaultds_read_b64_tr_b16 v1, v0 228*e97fb220SMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: invalid operand for instruction 229*e97fb220SMatt Arsenault// W32-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU 230*e97fb220SMatt Arsenault 231*e97fb220SMatt Arsenaultds_read_b64_tr_b16 v[0:1], s0 232*e97fb220SMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: invalid operand for instruction 233*e97fb220SMatt Arsenault// W32-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU 234*e97fb220SMatt Arsenault 235*e97fb220SMatt Arsenaultds_read_b64_tr_b16 v[2:3], v2 offset:-64 236*e97fb220SMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: expected a 16-bit unsigned offset 237*e97fb220SMatt Arsenault// W32-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU 238*e97fb220SMatt Arsenault 239*e97fb220SMatt Arsenault//===----------------------------------------------------------------------===// 240*e97fb220SMatt Arsenault// ds_read_b96_tr_b6 241*e97fb220SMatt Arsenault//===----------------------------------------------------------------------===// 242*e97fb220SMatt Arsenaultds_read_b96_tr_b6 v[1:3], v0 243*e97fb220SMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: invalid register class: vgpr tuples must be 64 bit aligned 244*e97fb220SMatt Arsenault// W32-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU 245*e97fb220SMatt Arsenault 246*e97fb220SMatt Arsenaultds_read_b96_tr_b6 v1, v0 247*e97fb220SMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: invalid operand for instruction 248*e97fb220SMatt Arsenault// W32-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU 249*e97fb220SMatt Arsenault 250*e97fb220SMatt Arsenaultds_read_b96_tr_b6 v[0:3], s0 251*e97fb220SMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: invalid operand for instruction 252*e97fb220SMatt Arsenault// W32-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU 253*e97fb220SMatt Arsenault 254*e97fb220SMatt Arsenaultds_read_b96_tr_b6 v[2:4], v2 offset:-64 255*e97fb220SMatt Arsenault// ERR: :[[@LINE-1]]:{{[0-9]+}}: error: expected a 16-bit unsigned offset 256*e97fb220SMatt Arsenault// W32-ERR: :[[@LINE-2]]:{{[0-9]+}}: error: instruction not supported on this GPU 257