1// RUN: mlir-opt -allow-unregistered-dialect %s -split-input-file -verify-diagnostics 2 3// expected-error@+1{{alignment attribute is not a power of 2}} 4llvm.mlir.global private @invalid_global_alignment(42 : i64) {alignment = 63} : i64 5 6// ----- 7 8// expected-error@+1{{expected llvm.noalias argument attribute to be a unit attribute}} 9func @invalid_noalias(%arg0: i32 {llvm.noalias = 3}) { 10 "llvm.return"() : () -> () 11} 12 13// ----- 14 15// expected-error@+1{{llvm.align argument attribute of non integer type}} 16func @invalid_align(%arg0: i32 {llvm.align = "foo"}) { 17 "llvm.return"() : () -> () 18} 19 20//////////////////////////////////////////////////////////////////////////////// 21 22// Check that parser errors are properly produced and do not crash the compiler. 23 24// ----- 25 26func @icmp_non_string(%arg0 : i32, %arg1 : i16) { 27 // expected-error@+1 {{invalid kind of attribute specified}} 28 llvm.icmp 42 %arg0, %arg0 : i32 29 return 30} 31 32// ----- 33 34func @icmp_wrong_string(%arg0 : i32, %arg1 : i16) { 35 // expected-error@+1 {{'foo' is an incorrect value of the 'predicate' attribute}} 36 llvm.icmp "foo" %arg0, %arg0 : i32 37 return 38} 39 40// ----- 41 42func @alloca_missing_input_result_type(%size : i64) { 43 // expected-error@+1 {{expected trailing function type with one argument and one result}} 44 llvm.alloca %size x i32 : () -> () 45} 46 47// ----- 48 49func @alloca_missing_input_type() { 50 // expected-error@+1 {{expected trailing function type with one argument and one result}} 51 llvm.alloca %size x i32 : () -> (!llvm.ptr<i32>) 52} 53 54// ----- 55 56func @alloca_missing_result_type() { 57 // expected-error@+1 {{expected trailing function type with one argument and one result}} 58 llvm.alloca %size x i32 : (i64) -> () 59} 60 61// ----- 62 63func @alloca_non_function_type() { 64 // expected-error@+1 {{expected trailing function type with one argument and one result}} 65 llvm.alloca %size x i32 : !llvm.ptr<i32> 66} 67 68// ----- 69 70func @alloca_non_integer_alignment() { 71 // expected-error@+1 {{expected integer alignment}} 72 llvm.alloca %size x i32 {alignment = 3.0} : !llvm.ptr<i32> 73} 74 75// ----- 76 77func @gep_missing_input_result_type(%pos : i64, %base : !llvm.ptr<f32>) { 78 // expected-error@+1 {{2 operands present, but expected 0}} 79 llvm.getelementptr %base[%pos] : () -> () 80} 81 82// ----- 83 84func @gep_missing_input_type(%pos : i64, %base : !llvm.ptr<f32>) { 85 // expected-error@+1 {{2 operands present, but expected 0}} 86 llvm.getelementptr %base[%pos] : () -> (!llvm.ptr<f32>) 87} 88 89// ----- 90 91func @gep_missing_result_type(%pos : i64, %base : !llvm.ptr<f32>) { 92 // expected-error@+1 {{op requires one result}} 93 llvm.getelementptr %base[%pos] : (!llvm.ptr<f32>, i64) -> () 94} 95 96// ----- 97 98func @gep_non_function_type(%pos : i64, %base : !llvm.ptr<f32>) { 99 // expected-error@+1 {{invalid kind of type specified}} 100 llvm.getelementptr %base[%pos] : !llvm.ptr<f32> 101} 102 103// ----- 104 105func @load_non_llvm_type(%foo : memref<f32>) { 106 // expected-error@+1 {{expected LLVM pointer type}} 107 llvm.load %foo : memref<f32> 108} 109 110// ----- 111 112func @load_non_ptr_type(%foo : f32) { 113 // expected-error@+1 {{expected LLVM pointer type}} 114 llvm.load %foo : f32 115} 116 117// ----- 118 119func @store_non_llvm_type(%foo : memref<f32>, %bar : f32) { 120 // expected-error@+1 {{expected LLVM pointer type}} 121 llvm.store %bar, %foo : memref<f32> 122} 123 124// ----- 125 126func @store_non_ptr_type(%foo : f32, %bar : f32) { 127 // expected-error@+1 {{expected LLVM pointer type}} 128 llvm.store %bar, %foo : f32 129} 130 131// ----- 132 133func @call_non_function_type(%callee : !llvm.func<i8 (i8)>, %arg : i8) { 134 // expected-error@+1 {{expected function type}} 135 llvm.call %callee(%arg) : !llvm.func<i8 (i8)> 136} 137 138// ----- 139 140func @invalid_call() { 141 // expected-error@+1 {{'llvm.call' op must have either a `callee` attribute or at least an operand}} 142 "llvm.call"() : () -> () 143} 144 145// ----- 146 147func @call_non_function_type(%callee : !llvm.func<i8 (i8)>, %arg : i8) { 148 // expected-error@+1 {{expected function type}} 149 llvm.call %callee(%arg) : !llvm.func<i8 (i8)> 150} 151 152// ----- 153 154func @call_unknown_symbol() { 155 // expected-error@+1 {{'llvm.call' op 'missing_callee' does not reference a symbol in the current scope}} 156 llvm.call @missing_callee() : () -> () 157} 158 159// ----- 160 161func private @standard_func_callee() 162 163func @call_non_llvm() { 164 // expected-error@+1 {{'llvm.call' op 'standard_func_callee' does not reference a valid LLVM function}} 165 llvm.call @standard_func_callee() : () -> () 166} 167 168// ----- 169 170func @call_non_llvm_indirect(%arg0 : tensor<*xi32>) { 171 // expected-error@+1 {{'llvm.call' op operand #0 must be LLVM dialect-compatible type}} 172 "llvm.call"(%arg0) : (tensor<*xi32>) -> () 173} 174 175// ----- 176 177llvm.func @callee_func(i8) -> () 178 179func @callee_arg_mismatch(%arg0 : i32) { 180 // expected-error@+1 {{'llvm.call' op operand type mismatch for operand 0: 'i32' != 'i8'}} 181 llvm.call @callee_func(%arg0) : (i32) -> () 182} 183 184// ----- 185 186func @indirect_callee_arg_mismatch(%arg0 : i32, %callee : !llvm.ptr<func<void(i8)>>) { 187 // expected-error@+1 {{'llvm.call' op operand type mismatch for operand 0: 'i32' != 'i8'}} 188 "llvm.call"(%callee, %arg0) : (!llvm.ptr<func<void(i8)>>, i32) -> () 189} 190 191// ----- 192 193llvm.func @callee_func() -> (i8) 194 195func @callee_return_mismatch() { 196 // expected-error@+1 {{'llvm.call' op result type mismatch: 'i32' != 'i8'}} 197 %res = llvm.call @callee_func() : () -> (i32) 198} 199 200// ----- 201 202func @indirect_callee_return_mismatch(%callee : !llvm.ptr<func<i8()>>) { 203 // expected-error@+1 {{'llvm.call' op result type mismatch: 'i32' != 'i8'}} 204 "llvm.call"(%callee) : (!llvm.ptr<func<i8()>>) -> (i32) 205} 206 207// ----- 208 209func @call_too_many_results(%callee : () -> (i32,i32)) { 210 // expected-error@+1 {{expected function with 0 or 1 result}} 211 llvm.call %callee() : () -> (i32, i32) 212} 213 214// ----- 215 216func @call_non_llvm_result(%callee : () -> (tensor<*xi32>)) { 217 // expected-error@+1 {{expected result to have LLVM type}} 218 llvm.call %callee() : () -> (tensor<*xi32>) 219} 220 221// ----- 222 223func @call_non_llvm_input(%callee : (tensor<*xi32>) -> (), %arg : tensor<*xi32>) { 224 // expected-error@+1 {{expected LLVM types as inputs}} 225 llvm.call %callee(%arg) : (tensor<*xi32>) -> () 226} 227 228// ----- 229 230llvm.func @void_func_result(%arg0: i32) { 231 // expected-error@below {{expected no operands}} 232 // expected-note@above {{when returning from function}} 233 llvm.return %arg0: i32 234} 235 236// ----- 237 238llvm.func @non_void_func_no_result() -> i32 { 239 // expected-error@below {{expected 1 operand}} 240 // expected-note@above {{when returning from function}} 241 llvm.return 242} 243 244// ----- 245 246llvm.func @func_result_mismatch(%arg0: f32) -> i32 { 247 // expected-error@below {{mismatching result types}} 248 // expected-note@above {{when returning from function}} 249 llvm.return %arg0 : f32 250} 251 252// ----- 253 254func @constant_wrong_type() { 255 // expected-error@+1 {{only supports integer, float, string or elements attributes}} 256 llvm.mlir.constant(@constant_wrong_type) : !llvm.ptr<func<void ()>> 257} 258 259// ----- 260 261func @constant_wrong_type_string() { 262 // expected-error@below {{expected array type of 3 i8 elements for the string constant}} 263 llvm.mlir.constant("foo") : !llvm.ptr<i8> 264} 265 266// ----- 267 268llvm.func @array_attribute_one_element() -> !llvm.struct<(f64, f64)> { 269 // expected-error @+1 {{expected array attribute with two elements, representing a complex constant}} 270 %0 = llvm.mlir.constant([1.0 : f64]) : !llvm.struct<(f64, f64)> 271 llvm.return %0 : !llvm.struct<(f64, f64)> 272} 273 274// ----- 275 276llvm.func @array_attribute_two_different_types() -> !llvm.struct<(f64, f64)> { 277 // expected-error @+1 {{expected array attribute with two elements, representing a complex constant}} 278 %0 = llvm.mlir.constant([1.0 : f64, 1.0 : f32]) : !llvm.struct<(f64, f64)> 279 llvm.return %0 : !llvm.struct<(f64, f64)> 280} 281 282// ----- 283 284llvm.func @struct_wrong_attribute_type() -> !llvm.struct<(f64, f64)> { 285 // expected-error @+1 {{expected array attribute with two elements, representing a complex constant}} 286 %0 = llvm.mlir.constant(1.0 : f64) : !llvm.struct<(f64, f64)> 287 llvm.return %0 : !llvm.struct<(f64, f64)> 288} 289 290// ----- 291 292llvm.func @struct_one_element() -> !llvm.struct<(f64)> { 293 // expected-error @+1 {{expected struct type with two elements of the same type, the type of a complex constant}} 294 %0 = llvm.mlir.constant([1.0 : f64, 1.0 : f64]) : !llvm.struct<(f64)> 295 llvm.return %0 : !llvm.struct<(f64)> 296} 297 298// ----- 299 300llvm.func @struct_two_different_elements() -> !llvm.struct<(f64, f32)> { 301 // expected-error @+1 {{expected struct type with two elements of the same type, the type of a complex constant}} 302 %0 = llvm.mlir.constant([1.0 : f64, 1.0 : f64]) : !llvm.struct<(f64, f32)> 303 llvm.return %0 : !llvm.struct<(f64, f32)> 304} 305 306// ----- 307 308llvm.func @struct_wrong_element_types() -> !llvm.struct<(!llvm.array<2 x f64>, !llvm.array<2 x f64>)> { 309 // expected-error @+1 {{expected struct element types to be floating point type or integer type}} 310 %0 = llvm.mlir.constant([dense<[1.0, 1.0]> : tensor<2xf64>, dense<[1.0, 1.0]> : tensor<2xf64>]) : !llvm.struct<(!llvm.array<2 x f64>, !llvm.array<2 x f64>)> 311 llvm.return %0 : !llvm.struct<(!llvm.array<2 x f64>, !llvm.array<2 x f64>)> 312} 313 314// ----- 315 316func @insertvalue_non_llvm_type(%a : i32, %b : i32) { 317 // expected-error@+1 {{expected LLVM IR Dialect type}} 318 llvm.insertvalue %a, %b[0] : tensor<*xi32> 319} 320 321// ----- 322 323func @insertvalue_non_array_position() { 324 // Note the double-type, otherwise attribute parsing consumes the trailing 325 // type of the op as the (wrong) attribute type. 326 // expected-error@+1 {{invalid kind of attribute specified}} 327 llvm.insertvalue %a, %b 0 : i32 : !llvm.struct<(i32)> 328} 329 330// ----- 331 332func @insertvalue_non_integer_position() { 333 // expected-error@+1 {{expected an array of integer literals}} 334 llvm.insertvalue %a, %b[0.0] : !llvm.struct<(i32)> 335} 336 337// ----- 338 339func @insertvalue_struct_out_of_bounds() { 340 // expected-error@+1 {{position out of bounds}} 341 llvm.insertvalue %a, %b[1] : !llvm.struct<(i32)> 342} 343 344// ----- 345 346func @insertvalue_array_out_of_bounds() { 347 // expected-error@+1 {{position out of bounds}} 348 llvm.insertvalue %a, %b[1] : !llvm.array<1 x i32> 349} 350 351// ----- 352 353func @insertvalue_wrong_nesting() { 354 // expected-error@+1 {{expected LLVM IR structure/array type}} 355 llvm.insertvalue %a, %b[0,0] : !llvm.struct<(i32)> 356} 357 358// ----- 359 360func @insertvalue_invalid_type(%a : !llvm.array<1 x i32>) -> !llvm.array<1 x i32> { 361 // expected-error@+1 {{'llvm.insertvalue' op Type mismatch: cannot insert '!llvm.array<1 x i32>' into '!llvm.array<1 x i32>'}} 362 %b = "llvm.insertvalue"(%a, %a) {position = [0]} : (!llvm.array<1 x i32>, !llvm.array<1 x i32>) -> !llvm.array<1 x i32> 363 return %b : !llvm.array<1 x i32> 364} 365 366// ----- 367 368func @extractvalue_invalid_type(%a : !llvm.array<4 x vector<8xf32>>) -> !llvm.array<4 x vector<8xf32>> { 369 // expected-error@+1 {{'llvm.extractvalue' op Type mismatch: extracting from '!llvm.array<4 x vector<8xf32>>' should produce 'vector<8xf32>' but this op returns '!llvm.array<4 x vector<8xf32>>'}} 370 %b = "llvm.extractvalue"(%a) {position = [1]} 371 : (!llvm.array<4 x vector<8xf32>>) -> !llvm.array<4 x vector<8xf32>> 372 return %b : !llvm.array<4 x vector<8xf32>> 373} 374 375 376// ----- 377 378func @extractvalue_non_llvm_type(%a : i32, %b : tensor<*xi32>) { 379 // expected-error@+1 {{expected LLVM IR Dialect type}} 380 llvm.extractvalue %b[0] : tensor<*xi32> 381} 382 383// ----- 384 385func @extractvalue_non_array_position() { 386 // Note the double-type, otherwise attribute parsing consumes the trailing 387 // type of the op as the (wrong) attribute type. 388 // expected-error@+1 {{invalid kind of attribute specified}} 389 llvm.extractvalue %b 0 : i32 : !llvm.struct<(i32)> 390} 391 392// ----- 393 394func @extractvalue_non_integer_position() { 395 // expected-error@+1 {{expected an array of integer literals}} 396 llvm.extractvalue %b[0.0] : !llvm.struct<(i32)> 397} 398 399// ----- 400 401func @extractvalue_struct_out_of_bounds() { 402 // expected-error@+1 {{position out of bounds}} 403 llvm.extractvalue %b[1] : !llvm.struct<(i32)> 404} 405 406// ----- 407 408func @extractvalue_array_out_of_bounds() { 409 // expected-error@+1 {{position out of bounds}} 410 llvm.extractvalue %b[1] : !llvm.array<1 x i32> 411} 412 413// ----- 414 415func @extractvalue_wrong_nesting() { 416 // expected-error@+1 {{expected LLVM IR structure/array type}} 417 llvm.extractvalue %b[0,0] : !llvm.struct<(i32)> 418} 419 420// ----- 421 422func @invalid_vector_type_1(%arg0: vector<4xf32>, %arg1: i32, %arg2: f32) { 423 // expected-error@+1 {{expected LLVM dialect-compatible vector type for operand #1}} 424 %0 = llvm.extractelement %arg2[%arg1 : i32] : f32 425} 426 427// ----- 428 429func @invalid_vector_type_2(%arg0: vector<4xf32>, %arg1: i32, %arg2: f32) { 430 // expected-error@+1 {{expected LLVM dialect-compatible vector type for operand #1}} 431 %0 = llvm.insertelement %arg2, %arg2[%arg1 : i32] : f32 432} 433 434// ----- 435 436func @invalid_vector_type_3(%arg0: vector<4xf32>, %arg1: i32, %arg2: f32) { 437 // expected-error@+1 {{expected LLVM IR dialect vector type for operand #1}} 438 %0 = llvm.shufflevector %arg2, %arg2 [0 : i32, 0 : i32, 0 : i32, 0 : i32, 7 : i32] : f32, f32 439} 440 441// ----- 442 443func @invalid_vector_type_4(%a : vector<4xf32>, %idx : i32) -> vector<4xf32> { 444 // expected-error@+1 {{'llvm.extractelement' op Type mismatch: extracting from 'vector<4xf32>' should produce 'f32' but this op returns 'vector<4xf32>'}} 445 %b = "llvm.extractelement"(%a, %idx) : (vector<4xf32>, i32) -> vector<4xf32> 446 return %b : vector<4xf32> 447} 448 449// ----- 450 451func @invalid_vector_type_5(%a : vector<4xf32>, %idx : i32) -> vector<4xf32> { 452 // expected-error@+1 {{'llvm.insertelement' op Type mismatch: cannot insert 'vector<4xf32>' into 'vector<4xf32>'}} 453 %b = "llvm.insertelement"(%a, %a, %idx) : (vector<4xf32>, vector<4xf32>, i32) -> vector<4xf32> 454 return %b : vector<4xf32> 455} 456 457// ----- 458 459func @null_non_llvm_type() { 460 // expected-error@+1 {{must be LLVM pointer type, but got 'i32'}} 461 llvm.mlir.null : i32 462} 463 464// ----- 465 466func @nvvm_invalid_shfl_pred_1(%arg0 : i32, %arg1 : i32, %arg2 : i32, %arg3 : i32) { 467 // expected-error@+1 {{expected return type to be a two-element struct with i1 as the second element}} 468 %0 = nvvm.shfl.sync.bfly %arg0, %arg3, %arg1, %arg2 {return_value_and_is_valid} : i32 469} 470 471// ----- 472 473func @nvvm_invalid_shfl_pred_2(%arg0 : i32, %arg1 : i32, %arg2 : i32, %arg3 : i32) { 474 // expected-error@+1 {{expected return type to be a two-element struct with i1 as the second element}} 475 %0 = nvvm.shfl.sync.bfly %arg0, %arg3, %arg1, %arg2 {return_value_and_is_valid} : !llvm.struct<(i32)> 476} 477 478// ----- 479 480func @nvvm_invalid_shfl_pred_3(%arg0 : i32, %arg1 : i32, %arg2 : i32, %arg3 : i32) { 481 // expected-error@+1 {{expected return type to be a two-element struct with i1 as the second element}} 482 %0 = nvvm.shfl.sync.bfly %arg0, %arg3, %arg1, %arg2 {return_value_and_is_valid} : !llvm.struct<(i32, i32)> 483} 484 485// ----- 486 487func @nvvm_invalid_mma_0(%a0 : f16, %a1 : vector<2xf16>, 488 %b0 : vector<2xf16>, %b1 : vector<2xf16>, 489 %c0 : f32, %c1 : f32, %c2 : f32, %c3 : f32, 490 %c4 : f32, %c5 : f32, %c6 : f32, %c7 : f32) { 491 // expected-error@+1 {{expected operands to be 4 <halfx2>s followed by either 4 <halfx2>s or 8 floats}} 492 %0 = nvvm.mma.sync %a0, %a1, %b0, %b1, %c0, %c1, %c2, %c3, %c4, %c5, %c6, %c7 {alayout="row", blayout="col"} : (f16, vector<2xf16>, vector<2xf16>, vector<2xf16>, f32, f32, f32, f32, f32, f32, f32, f32) -> !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)> 493 llvm.return %0 : !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)> 494} 495 496// ----- 497 498func @nvvm_invalid_mma_1(%a0 : vector<2xf16>, %a1 : vector<2xf16>, 499 %b0 : vector<2xf16>, %b1 : vector<2xf16>, 500 %c0 : f32, %c1 : f32, %c2 : f32, %c3 : f32, 501 %c4 : f32, %c5 : f32, %c6 : f32, %c7 : f32) { 502 // expected-error@+1 {{expected result type to be a struct of either 4 <halfx2>s or 8 floats}} 503 %0 = nvvm.mma.sync %a0, %a1, %b0, %b1, %c0, %c1, %c2, %c3, %c4, %c5, %c6, %c7 {alayout="row", blayout="col"} : (vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, f32, f32, f32, f32, f32, f32, f32, f32) -> !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f16)> 504 llvm.return %0 : !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f16)> 505} 506 507// ----- 508 509func @nvvm_invalid_mma_2(%a0 : vector<2xf16>, %a1 : vector<2xf16>, 510 %b0 : vector<2xf16>, %b1 : vector<2xf16>, 511 %c0 : f32, %c1 : f32, %c2 : f32, %c3 : f32, 512 %c4 : f32, %c5 : f32, %c6 : f32, %c7 : f32) { 513 // expected-error@+1 {{alayout and blayout attributes must be set to either "row" or "col"}} 514 %0 = nvvm.mma.sync %a0, %a1, %b0, %b1, %c0, %c1, %c2, %c3, %c4, %c5, %c6, %c7 : (vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, f32, f32, f32, f32, f32, f32, f32, f32) -> !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)> 515 llvm.return %0 : !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)> 516} 517 518// ----- 519 520func @nvvm_invalid_mma_3(%a0 : vector<2xf16>, %a1 : vector<2xf16>, 521 %b0 : vector<2xf16>, %b1 : vector<2xf16>, 522 %c0 : vector<2xf16>, %c1 : vector<2xf16>, 523 %c2 : vector<2xf16>, %c3 : vector<2xf16>) { 524 // expected-error@+1 {{unimplemented mma.sync variant}} 525 %0 = nvvm.mma.sync %a0, %a1, %b0, %b1, %c0, %c1, %c2, %c3 {alayout="row", blayout="col"} : (vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>) -> !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)> 526 llvm.return %0 : !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)> 527} 528 529// ----- 530 531func @nvvm_invalid_mma_4(%a0 : vector<2xf16>, %a1 : vector<2xf16>, 532 %b0 : vector<2xf16>, %b1 : vector<2xf16>, 533 %c0 : f32, %c1 : f32, %c2 : f32, %c3 : f32, 534 %c4 : f32, %c5 : f32, %c6 : f32, %c7 : f32) { 535 // expected-error@+1 {{unimplemented mma.sync variant}} 536 %0 = nvvm.mma.sync %a0, %a1, %b0, %b1, %c0, %c1, %c2, %c3, %c4, %c5, %c6, %c7 {alayout="row", blayout="col"} : (vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, f32, f32, f32, f32, f32, f32, f32, f32) -> !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)> 537 llvm.return %0 : !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)> 538} 539 540// ----- 541 542func @nvvm_invalid_mma_5(%a0 : vector<2xf16>, %a1 : vector<2xf16>, 543 %b0 : vector<2xf16>, %b1 : vector<2xf16>, 544 %c0 : f32, %c1 : f32, %c2 : f32, %c3 : f32, 545 %c4 : f32, %c5 : f32, %c6 : f32, %c7 : f32) { 546 // expected-error@+1 {{unimplemented mma.sync variant}} 547 %0 = nvvm.mma.sync %a0, %a1, %b0, %b1, %c0, %c1, %c2, %c3, %c4, %c5, %c6, %c7 {alayout="col", blayout="row"} : (vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, f32, f32, f32, f32, f32, f32, f32, f32) -> !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)> 548 llvm.return %0 : !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)> 549} 550 551// ----- 552 553func @nvvm_invalid_mma_6(%a0 : vector<2xf16>, %a1 : vector<2xf16>, 554 %b0 : vector<2xf16>, %b1 : vector<2xf16>, 555 %c0 : f32, %c1 : f32, %c2 : f32, %c3 : f32, 556 %c4 : f32, %c5 : f32, %c6 : f32, %c7 : f32) { 557 // expected-error@+1 {{invalid kind of type specified}} 558 %0 = nvvm.mma.sync %a0, %a1, %b0, %b1, %c0, %c1, %c2, %c3, %c4, %c5, %c6, %c7 {alayout="col", blayout="row"} : !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)> 559 llvm.return %0 : !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)> 560} 561 562// ----- 563 564func @nvvm_invalid_mma_7(%a0 : vector<2xf16>, %a1 : vector<2xf16>, 565 %b0 : vector<2xf16>, %b1 : vector<2xf16>, 566 %c0 : f32, %c1 : f32, %c2 : f32, %c3 : f32, 567 %c4 : f32, %c5 : f32, %c6 : f32, %c7 : f32) { 568 // expected-error@+1 {{op requires one result}} 569 %0:2 = nvvm.mma.sync %a0, %a1, %b0, %b1, %c0, %c1, %c2, %c3, %c4, %c5, %c6, %c7 {alayout="col", blayout="row"} : (vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, f32, f32, f32, f32, f32, f32, f32, f32) -> (!llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)>, i32) 570 llvm.return %0#0 : !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)> 571} 572 573// ----- 574 575func @atomicrmw_expected_ptr(%f32 : f32) { 576 // expected-error@+1 {{operand #0 must be LLVM pointer to floating point LLVM type or integer}} 577 %0 = "llvm.atomicrmw"(%f32, %f32) {bin_op=11, ordering=1} : (f32, f32) -> f32 578 llvm.return 579} 580 581// ----- 582 583func @atomicrmw_mismatched_operands(%f32_ptr : !llvm.ptr<f32>, %i32 : i32) { 584 // expected-error@+1 {{expected LLVM IR element type for operand #0 to match type for operand #1}} 585 %0 = "llvm.atomicrmw"(%f32_ptr, %i32) {bin_op=11, ordering=1} : (!llvm.ptr<f32>, i32) -> f32 586 llvm.return 587} 588 589// ----- 590 591func @atomicrmw_mismatched_operands(%f32_ptr : !llvm.ptr<f32>, %f32 : f32) { 592 // expected-error@+1 {{expected LLVM IR result type to match type for operand #1}} 593 %0 = "llvm.atomicrmw"(%f32_ptr, %f32) {bin_op=11, ordering=1} : (!llvm.ptr<f32>, f32) -> i32 594 llvm.return 595} 596 597// ----- 598 599func @atomicrmw_expected_float(%i32_ptr : !llvm.ptr<i32>, %i32 : i32) { 600 // expected-error@+1 {{expected LLVM IR floating point type}} 601 %0 = llvm.atomicrmw fadd %i32_ptr, %i32 unordered : i32 602 llvm.return 603} 604 605// ----- 606 607func @atomicrmw_unexpected_xchg_type(%i1_ptr : !llvm.ptr<i1>, %i1 : i1) { 608 // expected-error@+1 {{unexpected LLVM IR type for 'xchg' bin_op}} 609 %0 = llvm.atomicrmw xchg %i1_ptr, %i1 unordered : i1 610 llvm.return 611} 612 613// ----- 614 615func @atomicrmw_expected_int(%f32_ptr : !llvm.ptr<f32>, %f32 : f32) { 616 // expected-error@+1 {{expected LLVM IR integer type}} 617 %0 = llvm.atomicrmw max %f32_ptr, %f32 unordered : f32 618 llvm.return 619} 620 621// ----- 622 623func @cmpxchg_expected_ptr(%f32_ptr : !llvm.ptr<f32>, %f32 : f32) { 624 // expected-error@+1 {{op operand #0 must be LLVM pointer to integer or LLVM pointer type}} 625 %0 = "llvm.cmpxchg"(%f32, %f32, %f32) {success_ordering=2,failure_ordering=2} : (f32, f32, f32) -> !llvm.struct<(f32, i1)> 626 llvm.return 627} 628 629// ----- 630 631func @cmpxchg_mismatched_operands(%i64_ptr : !llvm.ptr<i64>, %i32 : i32) { 632 // expected-error@+1 {{expected LLVM IR element type for operand #0 to match type for all other operands}} 633 %0 = "llvm.cmpxchg"(%i64_ptr, %i32, %i32) {success_ordering=2,failure_ordering=2} : (!llvm.ptr<i64>, i32, i32) -> !llvm.struct<(i32, i1)> 634 llvm.return 635} 636 637// ----- 638 639func @cmpxchg_unexpected_type(%i1_ptr : !llvm.ptr<i1>, %i1 : i1) { 640 // expected-error@+1 {{unexpected LLVM IR type}} 641 %0 = llvm.cmpxchg %i1_ptr, %i1, %i1 monotonic monotonic : i1 642 llvm.return 643} 644 645// ----- 646 647func @cmpxchg_at_least_monotonic_success(%i32_ptr : !llvm.ptr<i32>, %i32 : i32) { 648 // expected-error@+1 {{ordering must be at least 'monotonic'}} 649 %0 = llvm.cmpxchg %i32_ptr, %i32, %i32 unordered monotonic : i32 650 llvm.return 651} 652 653// ----- 654 655func @cmpxchg_at_least_monotonic_failure(%i32_ptr : !llvm.ptr<i32>, %i32 : i32) { 656 // expected-error@+1 {{ordering must be at least 'monotonic'}} 657 %0 = llvm.cmpxchg %i32_ptr, %i32, %i32 monotonic unordered : i32 658 llvm.return 659} 660 661// ----- 662 663func @cmpxchg_failure_release(%i32_ptr : !llvm.ptr<i32>, %i32 : i32) { 664 // expected-error@+1 {{failure ordering cannot be 'release' or 'acq_rel'}} 665 %0 = llvm.cmpxchg %i32_ptr, %i32, %i32 acq_rel release : i32 666 llvm.return 667} 668 669// ----- 670 671func @cmpxchg_failure_acq_rel(%i32_ptr : !llvm.ptr<i32>, %i32 : i32) { 672 // expected-error@+1 {{failure ordering cannot be 'release' or 'acq_rel'}} 673 %0 = llvm.cmpxchg %i32_ptr, %i32, %i32 acq_rel acq_rel : i32 674 llvm.return 675} 676 677// ----- 678 679llvm.func @foo(i32) -> i32 680llvm.func @__gxx_personality_v0(...) -> i32 681 682llvm.func @bad_landingpad(%arg0: !llvm.ptr<ptr<i8>>) -> i32 attributes { personality = @__gxx_personality_v0} { 683 %0 = llvm.mlir.constant(3 : i32) : i32 684 %1 = llvm.mlir.constant(2 : i32) : i32 685 %2 = llvm.invoke @foo(%1) to ^bb1 unwind ^bb2 : (i32) -> i32 686^bb1: // pred: ^bb0 687 llvm.return %1 : i32 688^bb2: // pred: ^bb0 689 // expected-error@+1 {{clause #0 is not a known constant - null, addressof, bitcast}} 690 %3 = llvm.landingpad cleanup (catch %1 : i32) (catch %arg0 : !llvm.ptr<ptr<i8>>) : !llvm.struct<(ptr<i8>, i32)> 691 llvm.return %0 : i32 692} 693 694// ----- 695 696llvm.func @foo(i32) -> i32 697llvm.func @__gxx_personality_v0(...) -> i32 698 699llvm.func @caller(%arg0: i32) -> i32 attributes { personality = @__gxx_personality_v0} { 700 %0 = llvm.mlir.constant(1 : i32) : i32 701 %1 = llvm.alloca %0 x !llvm.ptr<i8> : (i32) -> !llvm.ptr<ptr<i8>> 702 // expected-note@+1 {{global addresses expected as operand to bitcast used in clauses for landingpad}} 703 %2 = llvm.bitcast %1 : !llvm.ptr<ptr<i8>> to !llvm.ptr<i8> 704 %3 = llvm.invoke @foo(%0) to ^bb1 unwind ^bb2 : (i32) -> i32 705^bb1: // pred: ^bb0 706 llvm.return %0 : i32 707^bb2: // pred: ^bb0 708 // expected-error@+1 {{constant clauses expected}} 709 %5 = llvm.landingpad (catch %2 : !llvm.ptr<i8>) : !llvm.struct<(ptr<i8>, i32)> 710 llvm.return %0 : i32 711} 712 713// ----- 714 715llvm.func @foo(i32) -> i32 716llvm.func @__gxx_personality_v0(...) -> i32 717 718llvm.func @caller(%arg0: i32) -> i32 attributes { personality = @__gxx_personality_v0} { 719 %0 = llvm.mlir.constant(1 : i32) : i32 720 %1 = llvm.invoke @foo(%0) to ^bb1 unwind ^bb2 : (i32) -> i32 721^bb1: // pred: ^bb0 722 llvm.return %0 : i32 723^bb2: // pred: ^bb0 724 // expected-error@+1 {{landingpad instruction expects at least one clause or cleanup attribute}} 725 %2 = llvm.landingpad : !llvm.struct<(ptr<i8>, i32)> 726 llvm.return %0 : i32 727} 728 729// ----- 730 731llvm.func @foo(i32) -> i32 732llvm.func @__gxx_personality_v0(...) -> i32 733 734llvm.func @caller(%arg0: i32) -> i32 attributes { personality = @__gxx_personality_v0 } { 735 %0 = llvm.mlir.constant(1 : i32) : i32 736 %1 = llvm.invoke @foo(%0) to ^bb1 unwind ^bb2 : (i32) -> i32 737^bb1: // pred: ^bb0 738 llvm.return %0 : i32 739^bb2: // pred: ^bb0 740 %2 = llvm.landingpad cleanup : !llvm.struct<(ptr<i8>, i32)> 741 // expected-error@+1 {{'llvm.resume' op expects landingpad value as operand}} 742 llvm.resume %0 : i32 743} 744 745// ----- 746 747llvm.func @foo(i32) -> i32 748 749llvm.func @caller(%arg0: i32) -> i32 { 750 %0 = llvm.mlir.constant(1 : i32) : i32 751 %1 = llvm.invoke @foo(%0) to ^bb1 unwind ^bb2 : (i32) -> i32 752^bb1: // pred: ^bb0 753 llvm.return %0 : i32 754^bb2: // pred: ^bb0 755 // expected-error@+1 {{llvm.landingpad needs to be in a function with a personality}} 756 %2 = llvm.landingpad cleanup : !llvm.struct<(ptr<i8>, i32)> 757 llvm.resume %2 : !llvm.struct<(ptr<i8>, i32)> 758} 759 760// ----- 761 762func @invalid_ordering_in_fence() { 763 // expected-error @+1 {{can be given only acquire, release, acq_rel, and seq_cst orderings}} 764 llvm.fence syncscope("agent") monotonic 765} 766 767// ----- 768 769// expected-error @+1 {{invalid data layout descriptor}} 770module attributes {llvm.data_layout = "#vjkr32"} { 771 func @invalid_data_layout() 772} 773 774// ----- 775 776func @switch_wrong_number_of_weights(%arg0 : i32) { 777 // expected-error@+1 {{expects number of branch weights to match number of successors: 3 vs 2}} 778 llvm.switch %arg0, ^bb1 [ 779 42: ^bb2(%arg0, %arg0 : i32, i32) 780 ] {branch_weights = dense<[13, 17, 19]> : vector<3xi32>} 781 782^bb1: // pred: ^bb0 783 llvm.return 784 785^bb2(%1: i32, %2: i32): // pred: ^bb0 786 llvm.return 787} 788 789// ----- 790 791// expected-error@below {{expected zero value for 'common' linkage}} 792llvm.mlir.global common @non_zero_global_common_linkage(42 : i32) : i32 793 794// ----- 795 796// expected-error@below {{expected zero value for 'common' linkage}} 797llvm.mlir.global common @non_zero_compound_global_common_linkage(dense<[0, 0, 0, 1, 0]> : vector<5xi32>) : !llvm.array<5 x i32> 798 799// ----- 800 801// expected-error@below {{expected array type for 'appending' linkage}} 802llvm.mlir.global appending @non_array_type_global_appending_linkage() : i32 803 804// ----- 805 806module { 807 llvm.func @loopOptions() { 808 // expected-error@below {{expected 'llvm.loop' to be a dictionary attribute}} 809 llvm.br ^bb4 {llvm.loop = "test"} 810 ^bb4: 811 llvm.return 812 } 813} 814 815// ----- 816 817module { 818 llvm.func @loopOptions() { 819 // expected-error@below {{expected 'parallel_access' to be an array attribute}} 820 llvm.br ^bb4 {llvm.loop = {parallel_access = "loop"}} 821 ^bb4: 822 llvm.return 823 } 824} 825 826// ----- 827 828module { 829 llvm.func @loopOptions() { 830 // expected-error@below {{expected '"loop"' to be a symbol reference}} 831 llvm.br ^bb4 {llvm.loop = {parallel_access = ["loop"]}} 832 ^bb4: 833 llvm.return 834 } 835} 836 837// ----- 838 839module { 840 llvm.func @loopOptions() { 841 // expected-error@below {{expected '@func1' to reference a metadata op}} 842 llvm.br ^bb4 {llvm.loop = {parallel_access = [@func1]}} 843 ^bb4: 844 llvm.return 845 } 846 llvm.func @func1() { 847 llvm.return 848 } 849} 850 851// ----- 852 853module { 854 llvm.func @loopOptions() { 855 // expected-error@below {{expected '@metadata' to reference an access_group op}} 856 llvm.br ^bb4 {llvm.loop = {parallel_access = [@metadata]}} 857 ^bb4: 858 llvm.return 859 } 860 llvm.metadata @metadata { 861 llvm.return 862 } 863} 864 865// ----- 866 867module { 868 llvm.func @loopOptions() { 869 // expected-error@below {{expected 'options' to be a `loopopts` attribute}} 870 llvm.br ^bb4 {llvm.loop = {options = "name"}} 871 ^bb4: 872 llvm.return 873 } 874} 875 876// ----- 877 878module { 879 llvm.func @loopOptions() { 880 // expected-error@below {{unknown loop option: name}} 881 llvm.br ^bb4 {llvm.loop = {options = #llvm.loopopts<name>}} 882 ^bb4: 883 llvm.return 884 } 885} 886 887// ----- 888 889module { 890 llvm.func @loopOptions() { 891 // expected-error@below {{loop option present twice}} 892 llvm.br ^bb4 {llvm.loop = {options = #llvm.loopopts<disable_licm = true, disable_licm = true>}} 893 ^bb4: 894 llvm.return 895 } 896} 897 898// ----- 899 900module { 901 llvm.func @accessGroups(%arg0 : !llvm.ptr<i32>) { 902 // expected-error@below {{attribute 'access_groups' failed to satisfy constraint: symbol ref array attribute}} 903 %0 = llvm.load %arg0 { "access_groups" = "test" } : !llvm.ptr<i32> 904 llvm.return 905 } 906} 907 908// ----- 909 910module { 911 llvm.func @accessGroups(%arg0 : !llvm.ptr<i32>) { 912 // expected-error@below {{expected '@func1' to specify a fully qualified reference}} 913 %0 = llvm.load %arg0 { "access_groups" = [@func1] } : !llvm.ptr<i32> 914 llvm.return 915 } 916 llvm.func @func1() { 917 llvm.return 918 } 919} 920 921// ----- 922 923module { 924 llvm.func @accessGroups(%arg0 : !llvm.ptr<i32>) { 925 // expected-error@below {{expected '@accessGroups::@group1' to reference a metadata op}} 926 %0 = llvm.load %arg0 { "access_groups" = [@accessGroups::@group1] } : !llvm.ptr<i32> 927 llvm.return 928 } 929 llvm.metadata @metadata { 930 llvm.return 931 } 932} 933 934// ----- 935 936module { 937 llvm.func @accessGroups(%arg0 : !llvm.ptr<i32>) { 938 // expected-error@below {{expected '@metadata::@group1' to be a valid reference}} 939 %0 = llvm.load %arg0 { "access_groups" = [@metadata::@group1] } : !llvm.ptr<i32> 940 llvm.return 941 } 942 llvm.metadata @metadata { 943 llvm.return 944 } 945} 946 947// ----- 948 949module { 950 llvm.func @accessGroups(%arg0 : !llvm.ptr<i32>) { 951 // expected-error@below {{expected '@metadata::@scope' to resolve to a llvm.access_group}} 952 %0 = llvm.load %arg0 { "access_groups" = [@metadata::@scope] } : !llvm.ptr<i32> 953 llvm.return 954 } 955 llvm.metadata @metadata { 956 llvm.alias_scope_domain @domain 957 llvm.alias_scope @scope { domain = @domain } 958 llvm.return 959 } 960} 961 962// ----- 963 964module { 965 llvm.func @accessGroups(%arg0 : !llvm.ptr<i32>) { 966 // expected-error@below {{attribute 'alias_scopes' failed to satisfy constraint: symbol ref array attribute}} 967 %0 = llvm.load %arg0 { "alias_scopes" = "test" } : !llvm.ptr<i32> 968 llvm.return 969 } 970} 971 972// ----- 973 974module { 975 llvm.func @accessGroups(%arg0 : !llvm.ptr<i32>) { 976 // expected-error@below {{attribute 'noalias_scopes' failed to satisfy constraint: symbol ref array attribute}} 977 %0 = llvm.load %arg0 { "noalias_scopes" = "test" } : !llvm.ptr<i32> 978 llvm.return 979 } 980} 981 982// ----- 983 984module { 985 llvm.func @aliasScope(%arg0 : !llvm.ptr<i32>) { 986 // expected-error@below {{expected '@metadata::@group' to resolve to a llvm.alias_scope}} 987 %0 = llvm.load %arg0 { "alias_scopes" = [@metadata::@group] } : !llvm.ptr<i32> 988 llvm.return 989 } 990 llvm.metadata @metadata { 991 llvm.access_group @group 992 llvm.return 993 } 994} 995 996// ----- 997 998module { 999 llvm.func @aliasScope(%arg0 : !llvm.ptr<i32>) { 1000 // expected-error@below {{expected '@metadata::@group' to resolve to a llvm.alias_scope}} 1001 %0 = llvm.load %arg0 { "noalias_scopes" = [@metadata::@group] } : !llvm.ptr<i32> 1002 llvm.return 1003 } 1004 llvm.metadata @metadata { 1005 llvm.access_group @group 1006 llvm.return 1007 } 1008} 1009 1010// ----- 1011 1012llvm.func @wmmaLoadOp_invalid_mem_space(%arg0: !llvm.ptr<i32, 5>, %arg1: i32) { 1013 // expected-error@+1 {{'nvvm.wmma.m16n16k16.load.a.f16.row.stride' op expected operands to be a source pointer in memory space 0, 1, 3 followed by ldm of the source}} 1014 %0 = nvvm.wmma.m16n16k16.load.a.f16.row.stride %arg0, %arg1 : (!llvm.ptr<i32, 5>, i32) -> !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)> 1015 1016 llvm.return 1017} 1018 1019// ----- 1020 1021llvm.func @wmmaLoadOp_invalid_missing_ldm(%arg0: !llvm.ptr<i32, 3>, %arg1: i32) { 1022 // expected-error@+1 {{'nvvm.wmma.m16n16k16.load.a.f16.row.stride' op expected operands to be a source pointer in memory space 0, 1, 3 followed by ldm of the source}} 1023 %0 = nvvm.wmma.m16n16k16.load.a.f16.row.stride %arg0: (!llvm.ptr<i32, 3>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)> 1024 1025 llvm.return 1026} 1027 1028// ----- 1029 1030llvm.func @wmmaLoadOp_invalid_AOp(%arg0: !llvm.ptr<i32, 3>, %arg1: i32) { 1031 // expected-error@+1 {{'nvvm.wmma.m16n16k16.load.a.f16.row.stride' op expected result type of loadAOp and loadBOp to be a struct of 8 <halfx2>s}} 1032 %0 = nvvm.wmma.m16n16k16.load.a.f16.row.stride %arg0, %arg1 : (!llvm.ptr<i32, 3>, i32) -> !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)> 1033 1034 llvm.return 1035} 1036 1037// ----- 1038 1039llvm.func @wmmaLoadOp_invalid_AOp(%arg0: !llvm.ptr<i32, 3>, %arg1: i32) { 1040 // expected-error@+1 {{nvvm.wmma.m16n16k16.load.a.f16.row.stride' op expected result type of loadAOp and loadBOp to be a struct of 8 <halfx2>s}} 1041 %0 = nvvm.wmma.m16n16k16.load.a.f16.row.stride %arg0, %arg1 : (!llvm.ptr<i32, 3>, i32) -> !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)> 1042 1043 llvm.return 1044} 1045 1046// ----- 1047 1048llvm.func @wmmaLoadOp_invalid_BOp(%arg0: !llvm.ptr<i32, 3>, %arg1: i32) { 1049 // expected-error@+1 {{'nvvm.wmma.m16n16k16.load.b.f16.row.stride' op expected result type of loadAOp and loadBOp to be a struct of 8 <halfx2>s}} 1050 %0 = nvvm.wmma.m16n16k16.load.b.f16.row.stride %arg0, %arg1 : (!llvm.ptr<i32, 3>, i32) -> !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)> 1051 1052 llvm.return 1053} 1054 1055// ----- 1056 1057llvm.func @wmmaLoadOp_invalid_COp(%arg0: !llvm.ptr<i32, 3>, %arg1: i32) { 1058 // expected-error@+1 {{'nvvm.wmma.m16n16k16.load.c.f16.row.stride' op expected result type of loadCOp to be a struct of 4 <halfx2>s or 8 f32s}} 1059 %0 = nvvm.wmma.m16n16k16.load.c.f16.row.stride %arg0, %arg1 : (!llvm.ptr<i32, 3>, i32) -> !llvm.struct<(vector<2xf16>, vector<2xf16>)> 1060 1061 llvm.return 1062} 1063 1064// ----- 1065 1066llvm.func @wmmaStoreOp_invalid_mem_space(%arg0: !llvm.ptr<i32, 5>, %arg1: vector<2 x f16>, 1067 %arg2: vector<2 x f16>, %arg3: vector<2 x f16>, 1068 %arg4: vector<2 xf16>, %arg5: i32) { 1069 // expected-error@+1 {{'nvvm.wmma.m16n16k16.store.d.f16.row.stride' op expected operands to be a source pointer in memoryspace 0, 1, 3 followed by ldm of the source}} 1070 nvvm.wmma.m16n16k16.store.d.f16.row.stride %arg0, %arg1, %arg2, %arg3, %arg4, %arg5 : !llvm.ptr<i32, 5>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, i32 1071 llvm.return 1072} 1073 1074// ----- 1075 1076llvm.func @wmmaStoreOp_invalid_missing_ldm(%arg0: !llvm.ptr<i32, 3>, %arg1: vector<2 x f16>, 1077 %arg2: vector<2 x f16>, %arg3: vector<2 x f16>, 1078 %arg4: vector<2 xf16>, %arg5: i32) { 1079 // expected-error@+1 {{'nvvm.wmma.m16n16k16.store.d.f16.row.stride' op expected operands to be a source pointer in memoryspace 0, 1, 3 followed by ldm of the source}} 1080 nvvm.wmma.m16n16k16.store.d.f16.row.stride %arg0, %arg1, %arg2, %arg3, %arg4 : !llvm.ptr<i32, 3>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16> 1081 llvm.return 1082} 1083 1084// ----- 1085 1086llvm.func @gpu_wmma_mma_op_invalid_operands(%arg0: vector<2 x f16>, %arg1: vector<2 x f16>, 1087 %arg2: vector<2 x f16>, %arg3: vector<2 x f16>, 1088 %arg4: vector<2 x f16>, %arg5: vector<2 x f16>, 1089 %arg6: vector<2 x f16>, %arg7: vector<2 x f16>, 1090 %arg8: vector<2 x f16>, %arg9: vector<2 x f16>, 1091 %arg10: vector<2 x f16>, %arg11: vector<2 x f16>, 1092 %arg12: vector<2 x f16>, %arg13: vector<2 x f16>, 1093 %arg14: vector<2 x f16>, %arg15: vector<2 x f16>, 1094 %arg16: vector<2 x f16>, %arg17: vector<2 x f16>, 1095 %arg18: vector<2 x f16>) { 1096 // expected-error@+1 {{'nvvm.wmma.m16n16k16.mma.row.row.f16.f16' op expected 20 <halfx2>s as operands}} 1097 %0 = nvvm.wmma.m16n16k16.mma.row.row.f16.f16 %arg0, %arg1, %arg2, %arg3, %arg4, %arg5, %arg6, %arg7, %arg8, %arg9, %arg10, %arg11, %arg12, %arg13, %arg14, %arg15, %arg16, %arg17, %arg18 : vector<2 x f16> -> !llvm.struct<(vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>)> 1098 llvm.return 1099} 1100 1101// ----- 1102 1103llvm.func @gpu_wmma_mma_op_results(%arg0: vector<2 x f16>, %arg1: vector<2 x f16>, 1104 %arg2: vector<2 x f16>, %arg3: vector<2 x f16>, 1105 %arg4: vector<2 x f16>, %arg5: vector<2 x f16>, 1106 %arg6: vector<2 x f16>, %arg7: vector<2 x f16>, 1107 %arg8: vector<2 x f16>, %arg9: vector<2 x f16>, 1108 %arg10: vector<2 x f16>, %arg11: vector<2 x f16>, 1109 %arg12: vector<2 x f16>, %arg13: vector<2 x f16>, 1110 %arg14: vector<2 x f16>, %arg15: vector<2 x f16>, 1111 %arg16: vector<2 x f16>, %arg17: vector<2 x f16>, 1112 %arg18: vector<2 x f16>, %arg19: vector<2 x f16>) { 1113 // expected-error@+1 {{expected result type to be a struct of 4 <halfx2>s}} 1114 %0 = nvvm.wmma.m16n16k16.mma.row.row.f16.f16 %arg0, %arg1, %arg2, %arg3, %arg4, %arg5, %arg6, %arg7, %arg8, %arg9, %arg10, %arg11, %arg12, %arg13, %arg14, %arg15, %arg16, %arg17, %arg18, %arg19 : vector<2 x f16> -> !llvm.struct<(vector<2 x f16>, vector<2 x f16>, vector<2 x f16>)> 1115 llvm.return 1116} 1117 1118// ----- 1119 1120llvm.func @gpu_wmma_mma_op_invalid_ab_operands(%arg0: vector<2 x f16>, %arg1: vector<2 x f16>, 1121 %arg2: vector<2 x f16>, %arg3: vector<2 x f16>, 1122 %arg4: vector<2 x f16>, %arg5: vector<2 x f16>, 1123 %arg6: vector<2 x f16>, %arg7: vector<2 x f16>, 1124 %arg8: vector<2 x f16>, %arg9: vector<2 x f16>, 1125 %arg10: vector<2 x f16>, %arg11: vector<2 x f16>, 1126 %arg12: vector<2 x f16>, %arg13: vector<2 x f16>, 1127 %arg14: vector<2 x f16>, %arg15: f32, 1128 %arg16: f32, %arg17: f32, %arg18: f32, %arg19: f32, 1129 %arg20: f32, %arg21: f32, %arg22: f32, %arg23: f32) { 1130 // expected-error@+1 {{'nvvm.wmma.m16n16k16.mma.row.row.f32.f32' op expected 16 <halfx2>s for `a` and `b` operand}} 1131 %0 = nvvm.wmma.m16n16k16.mma.row.row.f32.f32 %arg0, %arg1, %arg2, %arg3, %arg4, %arg5, %arg6, %arg7, %arg8, %arg9, %arg10, %arg11, %arg12, %arg13, %arg14, %arg15, %arg16, %arg17, %arg18, %arg19, %arg20, %arg21, %arg22, %arg23 : (vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, f32, f32, f32, f32, f32, f32, f32, f32, f32) -> !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)> 1132 llvm.return 1133} 1134 1135// ----- 1136 1137llvm.func @gpu_wmma_mma_op_invalid_c_operand(%arg0: vector<2 x f16>, %arg1: vector<2 x f16>, 1138 %arg2: vector<2 x f16>, %arg3: vector<2 x f16>, 1139 %arg4: vector<2 x f16>, %arg5: vector<2 x f16>, 1140 %arg6: vector<2 x f16>, %arg7: vector<2 x f16>, 1141 %arg8: vector<2 x f16>, %arg9: vector<2 x f16>, 1142 %arg10: vector<2 x f16>, %arg11: vector<2 x f16>, 1143 %arg12: vector<2 x f16>, %arg13: vector<2 x f16>, 1144 %arg14: vector<2 x f16>, %arg15: vector<2xf16>, 1145 %arg16: f32, %arg17: f32, %arg18: f32, %arg19: f32, 1146 %arg20: f32, %arg21: f32, %arg22: f32, %arg23: vector<2xf16>) { 1147 // expected-error@+1 {{'nvvm.wmma.m16n16k16.mma.row.row.f32.f32' op expected 8 f32s for `c` operand}} 1148 %0 = nvvm.wmma.m16n16k16.mma.row.row.f32.f32 %arg0, %arg1, %arg2, %arg3, %arg4, %arg5, %arg6, %arg7, %arg8, %arg9, %arg10, %arg11, %arg12, %arg13, %arg14, %arg15, %arg16, %arg17, %arg18, %arg19, %arg20, %arg21, %arg22, %arg23 : (vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, f32, f32, f32, f32, f32, f32, f32, vector<2xf16>) -> !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)> 1149 llvm.return 1150} 1151 1152// ----- 1153 1154llvm.func @gpu_wmma_mma_op_invalid_result(%arg0: vector<2 x f16>, %arg1: vector<2 x f16>, 1155 %arg2: vector<2 x f16>, %arg3: vector<2 x f16>, 1156 %arg4: vector<2 x f16>, %arg5: vector<2 x f16>, 1157 %arg6: vector<2 x f16>, %arg7: vector<2 x f16>, 1158 %arg8: vector<2 x f16>, %arg9: vector<2 x f16>, 1159 %arg10: vector<2 x f16>, %arg11: vector<2 x f16>, 1160 %arg12: vector<2 x f16>, %arg13: vector<2 x f16>, 1161 %arg14: vector<2 x f16>, %arg15: vector<2xf16>, 1162 %arg16: f32, %arg17: f32, %arg18: f32, %arg19: f32, 1163 %arg20: f32, %arg21: f32, %arg22: f32, %arg23: f32) { 1164 // expected-error@+1 {{'nvvm.wmma.m16n16k16.mma.row.row.f32.f32' op expected result type to be a struct of 8 f32s}} 1165 %0 = nvvm.wmma.m16n16k16.mma.row.row.f32.f32 %arg0, %arg1, %arg2, %arg3, %arg4, %arg5, %arg6, %arg7, %arg8, %arg9, %arg10, %arg11, %arg12, %arg13, %arg14, %arg15, %arg16, %arg17, %arg18, %arg19, %arg20, %arg21, %arg22, %arg23 : (vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, f32, f32, f32, f32, f32, f32, f32, f32) -> !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, vector<2xf16>)> 1166 llvm.return 1167} 1168 1169// ----- 1170 1171llvm.func @caller() { 1172 // expected-error @below {{expected function call to produce a value}} 1173 llvm.call @callee() : () -> () 1174 llvm.return 1175} 1176 1177llvm.func @callee() -> i32 1178 1179// ----- 1180 1181llvm.func @caller() { 1182 // expected-error @below {{calling function with void result must not produce values}} 1183 %0 = llvm.call @callee() : () -> i32 1184 llvm.return 1185} 1186 1187llvm.func @callee() -> () 1188 1189// ----- 1190 1191llvm.func @caller() { 1192 // expected-error @below {{expected function with 0 or 1 result}} 1193 %0:2 = llvm.call @callee() : () -> (i32, f32) 1194 llvm.return 1195} 1196 1197llvm.func @callee() -> !llvm.struct<(i32, f32)> 1198 1199// ----- 1200 1201func @bitcast(%arg0: vector<2x3xf32>) { 1202 // expected-error @below {{op operand #0 must be LLVM-compatible non-aggregate type}} 1203 llvm.bitcast %arg0 : vector<2x3xf32> to vector<2x3xi32> 1204 return 1205} 1206