2068 lines · plain
1// RUN: mlir-opt -allow-unregistered-dialect %s -split-input-file -verify-diagnostics2 3// expected-error@+1{{alignment attribute is not a power of 2}}4llvm.mlir.global private @invalid_global_alignment(42 : i64) {alignment = 63} : i645 6// -----7 8llvm.func @ctor() {9 llvm.return10}11 12// expected-error@+1{{ctors, priorities, and data must have the same number of elements}}13llvm.mlir.global_ctors ctors = [@ctor], priorities = [], data = [#llvm.zero]14 15// -----16 17llvm.func @dtor() {18 llvm.return19}20 21// expected-error@+1{{dtors, priorities, and data must have the same number of elements}}22llvm.mlir.global_dtors dtors = [@dtor], priorities = [0 : i32, 32767 : i32], data = [#llvm.zero]23 24// -----25 26// expected-error@+1{{'ctor' does not reference a valid LLVM function}}27llvm.mlir.global_ctors ctors = [@ctor], priorities = [0 : i32], data = [#llvm.zero]28 29// -----30 31llvm.func @dtor()32 33// expected-error@+1{{'dtor' does not have a definition}}34llvm.mlir.global_dtors dtors = [@dtor], priorities = [0 : i32], data = [#llvm.zero]35 36// -----37 38llvm.func @dtor() {39 llvm.return40}41 42// expected-error@+1{{data element must be symbol or #llvm.zero}}43llvm.mlir.global_dtors dtors = [@dtor], priorities = [0 : i32], data = [0 : i32]44 45////////////////////////////////////////////////////////////////////////////////46 47// Check that parser errors are properly produced and do not crash the compiler.48 49// -----50 51func.func @icmp_non_string(%arg0 : i32, %arg1 : i16) {52 // expected-error@+1 {{invalid kind of attribute specified}}53 llvm.icmp 42 %arg0, %arg0 : i3254 return55}56 57// -----58 59func.func @icmp_wrong_string(%arg0 : i32, %arg1 : i16) {60 // expected-error@+1 {{'foo' is an incorrect value of the 'predicate' attribute}}61 llvm.icmp "foo" %arg0, %arg0 : i3262 return63}64 65// -----66 67func.func @alloca_missing_input_result_type(%size : i64) {68 // expected-error@+1 {{expected trailing function type with one argument and one result}}69 llvm.alloca %size x i32 : () -> ()70}71 72// -----73 74func.func @alloca_missing_input_type() {75 // expected-error@+1 {{expected trailing function type with one argument and one result}}76 llvm.alloca %size x i32 : () -> (!llvm.ptr)77}78 79// -----80 81func.func @alloca_missing_result_type() {82 // expected-error@+1 {{expected trailing function type with one argument and one result}}83 llvm.alloca %size x i32 : (i64) -> ()84}85 86// -----87 88func.func @alloca_non_function_type() {89 // expected-error@+1 {{expected trailing function type with one argument and one result}}90 llvm.alloca %size x i32 : !llvm.ptr91}92 93// -----94 95func.func @alloca_non_integer_alignment() {96 // expected-error@+1 {{expected integer alignment}}97 llvm.alloca %size x i32 {alignment = 3.0} : !llvm.ptr98}99 100// -----101 102func.func @gep_missing_input_result_type(%pos : i64, %base : !llvm.ptr) {103 // expected-error@+1 {{number of operands and types do not match: got 2 operands and 0 types}}104 llvm.getelementptr %base[%pos] : () -> (), i64105}106 107// -----108 109func.func @gep_missing_input_type(%pos : i64, %base : !llvm.ptr) {110 // expected-error@+1 {{number of operands and types do not match: got 2 operands and 0 types}}111 llvm.getelementptr %base[%pos] : () -> (!llvm.ptr), i64112}113 114// -----115 116func.func @gep_missing_result_type(%pos : i64, %base : !llvm.ptr) {117 // expected-error@+1 {{op requires one result}}118 llvm.getelementptr %base[%pos] : (!llvm.ptr, i64) -> (), i64119}120 121// -----122 123func.func @gep_non_function_type(%pos : i64, %base : !llvm.ptr) {124 // expected-error@+1 {{invalid kind of type specified: expected builtin.function, but found '!llvm.ptr'}}125 llvm.getelementptr %base[%pos] : !llvm.ptr126}127 128// -----129 130func.func @gep_too_few_dynamic(%base : !llvm.ptr) {131 // expected-error@+1 {{expected as many dynamic indices as specified in 'rawConstantIndices'}}132 %1 = "llvm.getelementptr"(%base) {elem_type = f32, rawConstantIndices = array<i32: -2147483648>} : (!llvm.ptr) -> !llvm.ptr133}134 135// -----136 137func.func @load_non_llvm_type(%foo : memref<f32>) {138 // expected-error@+1 {{op operand #0 must be LLVM pointer type}}139 llvm.load %foo : memref<f32> -> f32140}141 142// -----143 144func.func @load_syncscope(%ptr : !llvm.ptr) {145 // expected-error@below {{expected syncscope to be null for non-atomic access}}146 %1 = "llvm.load"(%ptr) {syncscope = "singlethread"} : (!llvm.ptr) -> (f32)147}148 149// -----150 151func.func @load_unsupported_ordering(%ptr : !llvm.ptr) {152 // expected-error@below {{unsupported ordering 'release'}}153 %1 = llvm.load %ptr atomic release {alignment = 4 : i64} : !llvm.ptr -> f32154}155 156// -----157 158func.func @load_unsupported_type(%ptr : !llvm.ptr) {159 // expected-error@below {{unsupported type 'f80' for atomic access}}160 %1 = llvm.load %ptr atomic monotonic {alignment = 16 : i64} : !llvm.ptr -> f80161}162 163// -----164 165func.func @load_unsupported_type(%ptr : !llvm.ptr) {166 // expected-error@below {{unsupported type 'i1' for atomic access}}167 %1 = llvm.load %ptr atomic monotonic {alignment = 16 : i64} : !llvm.ptr -> i1168}169 170// -----171 172func.func @load_unsupported_type(%ptr : !llvm.ptr) {173 // expected-error@below {{unsupported type 'i33' for atomic access}}174 %1 = llvm.load %ptr atomic monotonic {alignment = 16 : i64} : !llvm.ptr -> i33175}176 177// -----178 179func.func @load_unaligned_atomic(%ptr : !llvm.ptr) {180 // expected-error@below {{expected alignment for atomic access}}181 %1 = llvm.load %ptr atomic monotonic : !llvm.ptr -> f32182}183 184// -----185 186func.func @store_syncscope(%val : f32, %ptr : !llvm.ptr) {187 // expected-error@below {{expected syncscope to be null for non-atomic access}}188 "llvm.store"(%val, %ptr) {syncscope = "singlethread"} : (f32, !llvm.ptr) -> ()189}190 191// -----192 193func.func @store_unsupported_ordering(%val : f32, %ptr : !llvm.ptr) {194 // expected-error@below {{unsupported ordering 'acquire'}}195 llvm.store %val, %ptr atomic acquire {alignment = 4 : i64} : f32, !llvm.ptr196}197 198// -----199 200func.func @store_unsupported_type(%val : f80, %ptr : !llvm.ptr) {201 // expected-error@below {{unsupported type 'f80' for atomic access}}202 llvm.store %val, %ptr atomic monotonic {alignment = 16 : i64} : f80, !llvm.ptr203}204 205// -----206 207func.func @store_unsupported_type(%val : i1, %ptr : !llvm.ptr) {208 // expected-error@below {{unsupported type 'i1' for atomic access}}209 llvm.store %val, %ptr atomic monotonic {alignment = 16 : i64} : i1, !llvm.ptr210}211 212// -----213 214func.func @store_unsupported_type(%val : i48, %ptr : !llvm.ptr) {215 // expected-error@below {{unsupported type 'i48' for atomic access}}216 llvm.store %val, %ptr atomic monotonic {alignment = 16 : i64} : i48, !llvm.ptr217}218 219// -----220 221func.func @store_unaligned_atomic(%val : f32, %ptr : !llvm.ptr) {222 // expected-error@below {{expected alignment for atomic access}}223 llvm.store %val, %ptr atomic monotonic : f32, !llvm.ptr224}225 226// -----227 228func.func @invalid_call() {229 // expected-error@+1 {{'llvm.call' op must have either a `callee` attribute or at least an operand}}230 "llvm.call"() {op_bundle_sizes = array<i32>} : () -> ()231 llvm.return232}233 234// -----235 236func.func @call_missing_ptr_type(%callee : !llvm.func<i8 (i8)>, %arg : i8) {237 // expected-error@+1 {{expected indirect call to have 2 trailing types}}238 llvm.call %callee(%arg) : (i8) -> (i8)239 llvm.return240}241 242// -----243 244func.func private @standard_func_callee()245 246func.func @call_missing_ptr_type(%arg : i8) {247 // expected-error@+2 {{expected '('}}248 // expected-error@+1 {{expected direct call to have 1 trailing type}}249 llvm.call @standard_func_callee(%arg) : !llvm.ptr, (i8) -> (i8)250 llvm.return251}252 253// -----254 255func.func @call_non_pointer_type(%callee : !llvm.func<i8 (i8)>, %arg : i8) {256 // expected-error@+1 {{indirect call expects a pointer as callee: '!llvm.func<i8 (i8)>'}}257 llvm.call %callee(%arg) : !llvm.func<i8 (i8)>, (i8) -> (i8)258 llvm.return259}260 261// -----262 263func.func @call_non_function_type(%callee : !llvm.ptr, %arg : i8) {264 // expected-error@+2 {{expected '('}}265 // expected-error@+1 {{expected trailing function type}}266 llvm.call %callee(%arg) : !llvm.ptr, !llvm.func<i8 (i8)>267 llvm.return268}269 270// -----271 272func.func @call_void_result_type(%callee : !llvm.ptr, %arg : i8) {273 // expected-error@+1 {{expected a non-void result type}}274 llvm.call %callee(%arg) : !llvm.ptr, (i8) -> (!llvm.void)275 llvm.return276}277 278// -----279 280func.func @call_unknown_symbol() {281 // expected-error@+1 {{'llvm.call' op 'missing_callee' does not reference a symbol in the current scope}}282 llvm.call @missing_callee() : () -> ()283 llvm.return284}285 286// -----287 288func.func private @standard_func_callee()289 290func.func @call_non_llvm() {291 // expected-error@+1 {{'llvm.call' op 'standard_func_callee' does not reference a valid LLVM function}}292 llvm.call @standard_func_callee() : () -> ()293 llvm.return294}295 296// -----297 298func.func @call_non_llvm_arg(%arg0 : tensor<*xi32>) {299 // expected-error@+1 {{'llvm.call' op operand #0 must be variadic of LLVM dialect-compatible type}}300 "llvm.call"(%arg0) {operandSegmentSizes = array<i32: 1, 0>, op_bundle_sizes = array<i32>} : (tensor<*xi32>) -> ()301 llvm.return302}303 304// -----305 306func.func @call_non_llvm_res(%callee : !llvm.ptr) {307 // expected-error@+1 {{'llvm.call' op result #0 must be LLVM dialect-compatible type}}308 llvm.call %callee() : !llvm.ptr, () -> (tensor<*xi32>)309 llvm.return310}311 312// -----313 314llvm.func @callee_func(i8) -> ()315 316func.func @callee_arg_mismatch(%arg0 : i32) {317 // expected-error@+1 {{'llvm.call' op operand type mismatch for operand 0: 'i32' != 'i8'}}318 llvm.call @callee_func(%arg0) : (i32) -> ()319 llvm.return320}321 322// -----323 324llvm.func @callee_func() -> (i8)325 326func.func @callee_return_mismatch() {327 // expected-error@+1 {{'llvm.call' op result type mismatch: 'i32' != 'i8'}}328 %res = llvm.call @callee_func() : () -> (i32)329 llvm.return330}331 332// -----333 334func.func @call_too_many_results(%callee : !llvm.ptr) {335 // expected-error@+1 {{expected function with 0 or 1 result}}336 llvm.call %callee() : !llvm.ptr, () -> (i32, i32)337 llvm.return338}339 340// -----341 342llvm.func @void_func_result(%arg0: i32) {343 // expected-error@below {{expected no operands}}344 // expected-note@above {{when returning from function}}345 llvm.return %arg0: i32346}347 348// -----349 350llvm.func @non_void_func_no_result() -> i32 {351 // expected-error@below {{expected 1 operand}}352 // expected-note@above {{when returning from function}}353 llvm.return354}355 356// -----357 358llvm.func @func_result_mismatch(%arg0: f32) -> i32 {359 // expected-error@below {{mismatching result types}}360 // expected-note@above {{when returning from function}}361 llvm.return %arg0 : f32362}363 364// -----365 366func.func @constant_wrong_type() {367 // expected-error@+1 {{only supports integer, float, string or elements attributes}}368 llvm.mlir.constant(@constant_wrong_type) : !llvm.ptr369}370 371// -----372 373func.func @constant_wrong_type_string() {374 // expected-error@below {{expected array type of 3 i8 elements for the string constant}}375 llvm.mlir.constant("foo") : !llvm.ptr376}377 378// -----379 380llvm.func @array_attribute_one_element() -> !llvm.struct<(f64, f64)> {381 // expected-error @+1 {{expected array attribute of size 2}}382 %0 = llvm.mlir.constant([1.0 : f64]) : !llvm.struct<(f64, f64)>383 llvm.return %0 : !llvm.struct<(f64, f64)>384}385 386// -----387 388llvm.func @array_attribute_two_different_types() -> !llvm.struct<(f64, f64)> {389 // expected-error @+1 {{struct element at index 1 is of wrong type}}390 %0 = llvm.mlir.constant([1.0 : f64, 1.0 : f32]) : !llvm.struct<(f64, f64)>391 llvm.return %0 : !llvm.struct<(f64, f64)>392}393 394// -----395 396llvm.func @struct_wrong_attribute_type() -> !llvm.struct<(f64, f64)> {397 // expected-error @+1 {{expected array attribute for struct type}}398 %0 = llvm.mlir.constant(1.0 : f64) : !llvm.struct<(f64, f64)>399 llvm.return %0 : !llvm.struct<(f64, f64)>400}401 402// -----403 404llvm.func @struct_one_element() -> !llvm.struct<(f64)> {405 // expected-error @+1 {{expected array attribute of size 1}}406 %0 = llvm.mlir.constant([1.0 : f64, 1.0 : f64]) : !llvm.struct<(f64)>407 llvm.return %0 : !llvm.struct<(f64)>408}409 410// -----411 412llvm.func @struct_two_different_elements() -> !llvm.struct<(f64, f32)> {413 // expected-error @+1 {{struct element at index 1 is of wrong type}}414 %0 = llvm.mlir.constant([1.0 : f64, 1.0 : f64]) : !llvm.struct<(f64, f32)>415 llvm.return %0 : !llvm.struct<(f64, f32)>416}417 418// -----419 420llvm.func @struct_wrong_element_types() -> !llvm.struct<(!llvm.array<2 x f64>, !llvm.array<2 x f64>)> {421 // expected-error @+1 {{expected struct element types to be floating point type or integer type}}422 %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>)>423 llvm.return %0 : !llvm.struct<(!llvm.array<2 x f64>, !llvm.array<2 x f64>)>424}425 426// -----427 428llvm.func @const_wrong_number_of_elements() -> vector<5xf64> {429 // expected-error @+1{{type and attribute have a different number of elements: 5 vs. 2}}430 %0 = llvm.mlir.constant(dense<[1.0, 1.0]> : tensor<2xf64>) : vector<5xf64>431 llvm.return %0 : vector<5xf64>432}433 434// -----435 436llvm.func @scalable_vec_requires_splat() -> vector<[4]xf64> {437 // expected-error @+1{{scalable vector type requires a splat attribute}}438 %0 = llvm.mlir.constant(dense<[1.0, 1.0, 2.0, 2.0]> : tensor<4xf64>) : vector<[4]xf64>439 llvm.return %0 : vector<[4]xf64>440}441 442 443// -----444 445llvm.func @int_attr_requires_int_type() -> f32 {446 // expected-error @below{{expected integer type}}447 %0 = llvm.mlir.constant(1 : index) : f32448 llvm.return %0 : f32449}450 451// -----452 453llvm.func @vector_int_attr_requires_int_type() -> vector<2xf32> {454 // expected-error @below{{expected integer element type}}455 %0 = llvm.mlir.constant(dense<[1, 2]> : vector<2xi32>) : vector<2xf32>456 llvm.return %0 : vector<2xf32>457}458 459// -----460 461llvm.func @float_attr_and_type_required_same() -> f16 {462 // expected-error @below{{attribute and type have different float semantics}}463 %cst = llvm.mlir.constant(1.0 : bf16) : f16464 llvm.return %cst : f16465}466 467// -----468 469llvm.func @vector_float_attr_and_type_required_same() -> vector<2xf16> {470 // expected-error @below{{attribute and type have different float semantics}}471 %cst = llvm.mlir.constant(dense<[1.0, 2.0]> : vector<2xbf16>) : vector<2xf16>472 llvm.return %cst : vector<2xf16>473}474 475// -----476 477llvm.func @incompatible_integer_type_for_float_attr() -> i32 {478 // expected-error @below{{expected integer type of width 16}}479 %cst = llvm.mlir.constant(1.0 : f16) : i32480 llvm.return %cst : i32481}482 483// -----484 485llvm.func @vector_incompatible_integer_type_for_float_attr() -> vector<2xi8> {486 // expected-error @below{{expected integer type of width 16}}487 %cst = llvm.mlir.constant(dense<[1.0, 2.0]> : vector<2xf16>) : vector<2xi8>488 llvm.return %cst : vector<2xi8>489}490 491// -----492 493llvm.func @vector_with_non_vector_type() -> f32 {494 // expected-error @below{{expected vector or array type}}495 %cst = llvm.mlir.constant(dense<100.0> : vector<1xf64>) : f32496 llvm.return %cst : f32497}498 499// -----500 501llvm.func @array_attr_with_invalid_type() -> i32 {502 // expected-error @below{{expected array or struct type for array attribute}}503 %0 = llvm.mlir.constant([1 : i32]) : i32504 llvm.return %0 : i32505}506 507// -----508 509llvm.func @elements_attribute_incompatible_nested_array_struct1_type() -> !llvm.array<2 x array<2 x array<2 x struct<(i32)>>>> {510 // expected-error @below{{expected integer element type for integer elements attribute}}511 %0 = llvm.mlir.constant(dense<[[[1, 2], [3, 4]], [[42, 43], [44, 45]]]> : tensor<2x2x2xi32>) : !llvm.array<2 x array<2 x array<2 x struct<(i32)>>>>512 llvm.return %0 : !llvm.array<2 x array<2 x array<2 x struct<(i32)>>>>513}514 515// -----516 517llvm.func @elements_attribute_incompatible_nested_array_struct3_type() -> !llvm.array<2 x array<2 x array<2 x struct<(i32, i32, i32)>>>> {518 // expected-error @below{{expected integer element type for integer elements attribute}}519 %0 = llvm.mlir.constant(dense<[[[1, 2], [3, 4]], [[42, 43], [44, 45]]]> : tensor<2x2x2xi32>) : !llvm.array<2 x array<2 x array<2 x struct<(i32, i32, i32)>>>>520 llvm.return %0 : !llvm.array<2 x array<2 x array<2 x struct<(i32, i32, i32)>>>>521}522 523// -----524 525llvm.func @invalid_struct_element_type() -> !llvm.struct<(f64, array<2 x i32>)> {526 // expected-error @below{{expected struct element types to be floating point type or integer type}}527 %0 = llvm.mlir.constant([1.0 : f64, dense<[1, 2]> : tensor<2xi32>]) : !llvm.struct<(f64, array<2 x i32>)>528 llvm.return %0 : !llvm.struct<(f64, array<2 x i32>)>529}530 531// -----532 533llvm.func @wrong_struct_element_attr_type() -> !llvm.struct<(f64, f64)> {534 // expected-error @below{{expected element of array attribute to be floating point or integer}}535 %0 = llvm.mlir.constant([dense<[1, 2]> : tensor<2xi32>, 2.0 : f64]) : !llvm.struct<(f64, f64)>536 llvm.return %0 : !llvm.struct<(f64, f64)>537}538 539// -----540 541llvm.func @struct_wrong_attribute_element_type() -> !llvm.struct<(f64, f64)> {542 // expected-error @below{{struct element at index 0 is of wrong type}}543 %0 = llvm.mlir.constant([1.0 : f32, 1.0 : f32]) : !llvm.struct<(f64, f64)>544 llvm.return %0 : !llvm.struct<(f64, f64)>545}546 547// -----548 549func.func @insertvalue_non_llvm_type(%a : i32, %b : i32) {550 // expected-error@+2 {{expected LLVM IR Dialect type}}551 llvm.insertvalue %a, %b[0] : tensor<*xi32>552}553 554// -----555 556func.func @insertvalue_struct_out_of_bounds() {557 // expected-error@+2 {{position out of bounds}}558 llvm.insertvalue %a, %b[1] : !llvm.struct<(i32)>559}560 561// -----562 563func.func @insertvalue_array_out_of_bounds() {564 // expected-error@+2 {{position out of bounds}}565 llvm.insertvalue %a, %b[1] : !llvm.array<1 x i32>566}567 568// -----569 570func.func @insertvalue_wrong_nesting() {571 // expected-error@+2 {{expected LLVM IR structure/array type}}572 llvm.insertvalue %a, %b[0,0] : !llvm.struct<(i32)>573}574 575// -----576 577func.func @insertvalue_invalid_type(%a : !llvm.array<1 x i32>) -> !llvm.array<1 x i32> {578 // expected-error@+1 {{'llvm.insertvalue' op Type mismatch: cannot insert '!llvm.array<1 x i32>' into '!llvm.array<1 x i32>'}}579 %b = "llvm.insertvalue"(%a, %a) {position = array<i64: 0>} : (!llvm.array<1 x i32>, !llvm.array<1 x i32>) -> !llvm.array<1 x i32>580 return %b : !llvm.array<1 x i32>581}582 583// -----584 585func.func @extractvalue_invalid_type(%a : !llvm.array<4 x vector<8xf32>>) -> !llvm.array<4 x vector<8xf32>> {586 // 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>>'}}587 %b = "llvm.extractvalue"(%a) {position = array<i64: 1>}588 : (!llvm.array<4 x vector<8xf32>>) -> !llvm.array<4 x vector<8xf32>>589 return %b : !llvm.array<4 x vector<8xf32>>590}591 592// -----593 594func.func @extractvalue_non_llvm_type(%a : i32, %b : tensor<*xi32>) {595 // expected-error@+2 {{expected LLVM IR Dialect type}}596 llvm.extractvalue %b[0] : tensor<*xi32>597}598 599// -----600 601func.func @extractvalue_struct_out_of_bounds() {602 // expected-error@+2 {{position out of bounds}}603 llvm.extractvalue %b[1] : !llvm.struct<(i32)>604}605 606// -----607 608func.func @extractvalue_array_out_of_bounds() {609 // expected-error@+2 {{position out of bounds}}610 llvm.extractvalue %b[1] : !llvm.array<1 x i32>611}612 613// -----614 615func.func @extractvalue_wrong_nesting() {616 // expected-error@+2 {{expected LLVM IR structure/array type}}617 llvm.extractvalue %b[0,0] : !llvm.struct<(i32)>618}619 620// -----621 622func.func @invalid_vector_type_1(%arg0: vector<4xf32>, %arg1: i32, %arg2: f32) {623 // expected-error@+1 {{invalid kind of type specified: expected builtin.vector, but found 'f32'}}624 %0 = llvm.extractelement %arg2[%arg1 : i32] : f32625}626 627// -----628 629func.func @invalid_vector_type_2(%arg0: vector<4xf32>, %arg1: i32, %arg2: f32) {630 // expected-error@+1 {{invalid kind of type specified: expected builtin.vector, but found 'f32'}}631 %0 = llvm.insertelement %arg2, %arg2[%arg1 : i32] : f32632}633 634// -----635 636func.func @invalid_vector_type_3(%arg0: vector<4xf32>, %arg1: i32, %arg2: f32) {637 // expected-error@+1 {{invalid kind of type specified: expected builtin.vector, but found 'f32'}}638 %0 = llvm.shufflevector %arg2, %arg2 [0, 0, 0, 0, 7] : f32639}640 641// -----642 643func.func @invalid_vector_type_4(%a : vector<4xf32>, %idx : i32) -> vector<4xf32> {644 // expected-error@+1 {{failed to verify that result type matches vector element type}}645 %b = "llvm.extractelement"(%a, %idx) : (vector<4xf32>, i32) -> vector<4xf32>646 return %b : vector<4xf32>647}648 649// -----650 651func.func @invalid_vector_type_5(%a : vector<4xf32>, %idx : i32) -> vector<4xf32> {652 // expected-error@+1 {{failed to verify that argument type matches vector element type}}653 %b = "llvm.insertelement"(%a, %a, %idx) : (vector<4xf32>, vector<4xf32>, i32) -> vector<4xf32>654 return %b : vector<4xf32>655}656 657// -----658 659func.func @zero_non_llvm_type() {660 // expected-error@+1 {{'llvm.mlir.zero' op result #0 must be LLVM dialect-compatible type, but got 'tensor<4xi32>'}}661 llvm.mlir.zero : tensor<4xi32>662}663 664// -----665 666func.func @nvvm_invalid_shfl_pred_1(%arg0 : i32, %arg1 : i32, %arg2 : i32, %arg3 : i32) {667 // expected-error@+1 {{expected return type to be a two-element struct}}668 %0 = nvvm.shfl.sync bfly %arg0, %arg3, %arg1, %arg2 {return_value_and_is_valid} : i32 -> i32669}670 671// -----672 673func.func @nvvm_invalid_shfl_pred_2(%arg0 : i32, %arg1 : i32, %arg2 : i32, %arg3 : i32) {674 // expected-error@+1 {{expected return type to be a two-element struct}}675 %0 = nvvm.shfl.sync bfly %arg0, %arg3, %arg1, %arg2 {return_value_and_is_valid} : i32 -> !llvm.struct<(i32)>676}677 678// -----679 680func.func @nvvm_invalid_shfl_pred_3(%arg0 : i32, %arg1 : i32, %arg2 : i32, %arg3 : i32) {681 // expected-error@+1 {{expected second element in the returned struct to be of type 'i1' but got 'i32' instead}}682 %0 = nvvm.shfl.sync bfly %arg0, %arg3, %arg1, %arg2 {return_value_and_is_valid} : i32 -> !llvm.struct<(i32, i32)>683}684 685// -----686 687func.func @nvvm_invalid_mma_0(%a0 : f16, %a1 : f16,688 %b0 : vector<2xf16>, %b1 : vector<2xf16>,689 %c0 : f32, %c1 : f32, %c2 : f32, %c3 : f32,690 %c4 : f32, %c5 : f32, %c6 : f32, %c7 : f32) {691 // expected-error@+1 {{Could not match types for the A operands; expected one of 2xvector<2xf16> but got f16, f16}}692 %0 = nvvm.mma.sync A[%a0, %a1] B[%b0, %b1] C[%c0, %c1, %c2, %c3, %c4, %c5, %c6, %c7]693 {layoutA=#nvvm.mma_layout<row>, layoutB=#nvvm.mma_layout<col>, shape = #nvvm.shape<m = 8, n = 8, k = 4>} : (f16, vector<2xf16>, f32) -> !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)>694 llvm.return %0 : !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)>695}696 697// -----698 699func.func @nvvm_invalid_mma_1(%a0 : vector<2xf16>, %a1 : vector<2xf16>,700 %b0 : vector<2xf16>, %b1 : vector<2xf16>,701 %c0 : f32, %c1 : f32, %c2 : f32, %c3 : f32,702 %c4 : f32, %c5 : f32, %c6 : f32, %c7 : f32) {703 // expected-error@+1 {{Could not match allowed types for the result; expected one of !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>, !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)> but got !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f16)>}}704 %0 = nvvm.mma.sync A[%a0, %a1] B[%b0, %b1] C[%c0, %c1, %c2, %c3, %c4, %c5, %c6, %c7]705 {layoutA=#nvvm.mma_layout<row>, layoutB=#nvvm.mma_layout<col>, shape = #nvvm.shape<m = 8, n = 8, k = 4>} : (vector<2xf16>, vector<2xf16>, f32) -> !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f16)>706 llvm.return %0 : !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f16)>707}708 709// -----710 711func.func @nvvm_invalid_mma_2(%a0 : vector<2xf16>, %a1 : vector<2xf16>,712 %b0 : vector<2xf16>, %b1 : vector<2xf16>,713 %c0 : f32, %c1 : f32, %c2 : f32, %c3 : f32,714 %c4 : f32, %c5 : f32, %c6 : f32, %c7 : f32) {715 // expected-error@+1 {{op requires attribute 'layoutA'}}716 %0 = nvvm.mma.sync A[%a0, %a1] B[%b0, %b1] C[%c0, %c1, %c2, %c3, %c4, %c5, %c6, %c7]717 {shape = #nvvm.shape<m = 8, n = 8, k = 4>}: (vector<2xf16>, vector<2xf16>, f32) -> !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)>718 llvm.return %0 : !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)>719}720 721// -----722 723func.func @nvvm_invalid_mma_3(%a0 : vector<2xf16>, %a1 : vector<2xf16>,724 %b0 : vector<2xf16>, %b1 : vector<2xf16>,725 %c0 : vector<2xf16>, %c1 : vector<2xf16>) {726 // expected-error@+1 {{unimplemented variant for MMA shape <8, 8, 16>}}727 %0 = nvvm.mma.sync A[%a0, %a1] B[%b0, %b1] C[%c0, %c1] {layoutA=#nvvm.mma_layout<row>, layoutB=#nvvm.mma_layout<col>, shape = #nvvm.shape<m = 8, n = 8, k = 16>} : (vector<2xf16>, vector<2xf16>, vector<2xf16>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>)>728 llvm.return %0 : !llvm.struct<(vector<2xf16>, vector<2xf16>)>729}730 731// -----732 733func.func @nvvm_invalid_mma_8(%a0 : i32, %a1 : i32,734 %b0 : i32,735 %c0 : i32, %c1 : i32, %c2 : i32, %c3 : i32) {736 // expected-error@+1 {{op requires b1Op attribute}}737 %0 = nvvm.mma.sync A[%a0, %a1] B[%b0] C[%c0, %c1, %c2, %c3]738 {layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>,739 multiplicandAPtxType = #nvvm.mma_type<b1>, multiplicandBPtxType = #nvvm.mma_type<b1>,740 shape = #nvvm.shape<m = 16, n = 8, k = 128>} : (i32, i32, i32) -> !llvm.struct<(i32,i32,i32,i32)>741 llvm.return %0 : !llvm.struct<(i32,i32,i32,i32)>742}743 744// -----745 746// f32 return type, f16 accumulate type747llvm.func @nvvm_mma_m16n8k16_f32_f16(%a0 : vector<2xf16>, %a1 : vector<2xf16>,748 %a2 : vector<2xf16>, %a3 : vector<2xf16>,749 %b0 : vector<2xf16>, %b1 : vector<2xf16>,750 %c0 : vector<2xf16>, %c1 : vector<2xf16>) -> !llvm.struct<(f32, f32, f32, f32)> {751 // C and D should have the same type according to PTX ISA752 // expected-error@+1 {{'nvvm.mma.sync' op ctype does not match dtype}}753 %0 = nvvm.mma.sync A[%a0, %a1, %a2, %a3] B[%b0, %b1] C[%c0, %c1]754 {layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>,755 shape = #nvvm.shape<m = 16, n = 8, k = 16>} : (vector<2xf16>, vector<2xf16>, vector<2xf16>) -> !llvm.struct<(f32, f32, f32, f32)>756 llvm.return %0 : !llvm.struct<(f32, f32, f32, f32)>757}758 759// -----760 761// f16 return type, f32 accumulate type762llvm.func @nvvm_mma_m16n8k16_f16_f32(%a0 : vector<2xf16>, %a1 : vector<2xf16>,763 %a2 : vector<2xf16>, %a3 : vector<2xf16>,764 %b0 : vector<2xf16>, %b1 : vector<2xf16>,765 %c0 : f32, %c1 : f32, %c2 : f32, %c3 : f32) -> !llvm.struct<(vector<2xf16>, vector<2xf16>)> {766 // C and D should have the same type according to PTX ISA767 // expected-error@+1 {{'nvvm.mma.sync' op ctype does not match dtype}}768 %0 = nvvm.mma.sync A[%a0, %a1, %a2, %a3] B[%b0, %b1] C[%c0, %c1, %c2, %c3]769 {layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>,770 shape = #nvvm.shape<m = 16, n = 8, k = 16>} : (vector<2xf16>, vector<2xf16>, f32) -> !llvm.struct<(vector<2xf16>, vector<2xf16>)>771 llvm.return %0 : !llvm.struct<(vector<2xf16>, vector<2xf16>)>772}773 774// -----775 776func.func @atomicrmw_mismatched_operands(%f32_ptr : !llvm.ptr, %f32 : f32) {777 // expected-error@+1 {{op failed to verify that result #0 and operand #1 have the same type}}778 %0 = "llvm.atomicrmw"(%f32_ptr, %f32) {bin_op=11, ordering=1} : (!llvm.ptr, f32) -> i32779 llvm.return780}781 782// -----783 784func.func @atomicrmw_expected_float(%i32_ptr : !llvm.ptr, %i32 : i32) {785 // expected-error@+1 {{expected LLVM IR floating point type}}786 %0 = llvm.atomicrmw fadd %i32_ptr, %i32 unordered : !llvm.ptr, i32787 llvm.return788}789 790// -----791 792func.func @atomicrmw_scalable_vector(%ptr : !llvm.ptr, %f32_vec : vector<[2]xf32>) {793 // expected-error@+1 {{'val' must be floating point LLVM type or LLVM pointer type or signless integer or LLVM dialect-compatible fixed-length vector type}}794 %0 = llvm.atomicrmw fadd %ptr, %f32_vec unordered : !llvm.ptr, vector<[2]xf32>795 llvm.return796}797 798// -----799 800func.func @atomicrmw_vector_expected_float(%ptr : !llvm.ptr, %i32_vec : vector<3xi32>) {801 // expected-error@+1 {{expected LLVM IR floating point type for vector element}}802 %0 = llvm.atomicrmw fadd %ptr, %i32_vec unordered : !llvm.ptr, vector<3xi32>803 llvm.return804}805 806// -----807 808func.func @atomicrmw_unexpected_xchg_type(%i1_ptr : !llvm.ptr, %i1 : i1) {809 // expected-error@+1 {{unexpected LLVM IR type for 'xchg' bin_op}}810 %0 = llvm.atomicrmw xchg %i1_ptr, %i1 unordered : !llvm.ptr, i1811 llvm.return812}813 814// -----815 816func.func @atomicrmw_expected_int(%f32_ptr : !llvm.ptr, %f32 : f32) {817 // expected-error@+1 {{expected LLVM IR integer type}}818 %0 = llvm.atomicrmw max %f32_ptr, %f32 unordered : !llvm.ptr, f32819 llvm.return820}821 822// -----823 824func.func @cmpxchg_mismatched_value_operands(%ptr : !llvm.ptr, %i32 : i32, %i64 : i64) {825 // expected-error@+1 {{op failed to verify that operand #1 and operand #2 have the same type}}826 %0 = "llvm.cmpxchg"(%ptr, %i32, %i64) {success_ordering=2,failure_ordering=2} : (!llvm.ptr, i32, i64) -> !llvm.struct<(i32, i1)>827 llvm.return828}829 830// -----831 832func.func @cmpxchg_mismatched_result(%ptr : !llvm.ptr, %i64 : i64) {833 // expected-error@+1 {{op failed to verify that result #0 has an LLVM struct type consisting of the type of operand #2 and a bool}}834 %0 = "llvm.cmpxchg"(%ptr, %i64, %i64) {success_ordering=2,failure_ordering=2} : (!llvm.ptr, i64, i64) -> !llvm.struct<(i64, i64)>835 llvm.return836}837 838// -----839 840func.func @cmpxchg_unexpected_type(%i1_ptr : !llvm.ptr, %i1 : i1) {841 // expected-error@+1 {{unexpected LLVM IR type}}842 %0 = llvm.cmpxchg %i1_ptr, %i1, %i1 monotonic monotonic : !llvm.ptr, i1843 llvm.return844}845 846// -----847 848func.func @cmpxchg_at_least_monotonic_success(%i32_ptr : !llvm.ptr, %i32 : i32) {849 // expected-error@+1 {{ordering must be at least 'monotonic'}}850 %0 = llvm.cmpxchg %i32_ptr, %i32, %i32 unordered monotonic : !llvm.ptr, i32851 llvm.return852}853 854// -----855 856func.func @cmpxchg_at_least_monotonic_failure(%i32_ptr : !llvm.ptr, %i32 : i32) {857 // expected-error@+1 {{ordering must be at least 'monotonic'}}858 %0 = llvm.cmpxchg %i32_ptr, %i32, %i32 monotonic unordered : !llvm.ptr, i32859 llvm.return860}861 862// -----863 864func.func @cmpxchg_failure_release(%i32_ptr : !llvm.ptr, %i32 : i32) {865 // expected-error@+1 {{failure ordering cannot be 'release' or 'acq_rel'}}866 %0 = llvm.cmpxchg %i32_ptr, %i32, %i32 acq_rel release : !llvm.ptr, i32867 llvm.return868}869 870// -----871 872func.func @cmpxchg_failure_acq_rel(%i32_ptr : !llvm.ptr, %i32 : i32) {873 // expected-error@+1 {{failure ordering cannot be 'release' or 'acq_rel'}}874 %0 = llvm.cmpxchg %i32_ptr, %i32, %i32 acq_rel acq_rel : !llvm.ptr, i32875 llvm.return876}877 878// -----879 880llvm.func @foo(i32) -> i32881llvm.func @__gxx_personality_v0(...) -> i32882 883llvm.func @bad_landingpad(%arg0: !llvm.ptr) -> i32 attributes { personality = @__gxx_personality_v0} {884 %0 = llvm.mlir.constant(3 : i32) : i32885 %1 = llvm.mlir.constant(2 : i32) : i32886 %2 = llvm.invoke @foo(%1) to ^bb1 unwind ^bb2 : (i32) -> i32887^bb1: // pred: ^bb0888 llvm.return %1 : i32889^bb2: // pred: ^bb0890 // expected-error@+1 {{clause #0 is not a known constant - null, addressof, bitcast}}891 %3 = llvm.landingpad cleanup (catch %1 : i32) (catch %arg0 : !llvm.ptr) : !llvm.struct<(ptr, i32)>892 llvm.return %0 : i32893}894 895// -----896 897llvm.func @foo(i32) -> i32898llvm.func @__gxx_personality_v0(...) -> i32899 900llvm.func @caller(%arg0: i32) -> i32 attributes { personality = @__gxx_personality_v0} {901 %0 = llvm.mlir.constant(1 : i32) : i32902 %1 = llvm.alloca %0 x !llvm.ptr : (i32) -> !llvm.ptr903 // expected-note@+1 {{global addresses expected as operand to bitcast used in clauses for landingpad}}904 %2 = llvm.bitcast %1 : !llvm.ptr to !llvm.ptr905 %3 = llvm.invoke @foo(%0) to ^bb1 unwind ^bb2 : (i32) -> i32906^bb1: // pred: ^bb0907 llvm.return %0 : i32908^bb2: // pred: ^bb0909 // expected-error@+1 {{constant clauses expected}}910 %5 = llvm.landingpad (catch %2 : !llvm.ptr) : !llvm.struct<(ptr, i32)>911 llvm.return %0 : i32912}913 914// -----915 916llvm.func @foo(i32) -> i32917llvm.func @__gxx_personality_v0(...) -> i32918 919llvm.func @caller(%arg0: i32) -> i32 attributes { personality = @__gxx_personality_v0} {920 %0 = llvm.mlir.constant(1 : i32) : i32921 %1 = llvm.invoke @foo(%0) to ^bb1 unwind ^bb2 : (i32) -> i32922^bb1: // pred: ^bb0923 llvm.return %0 : i32924^bb2: // pred: ^bb0925 // expected-error@+1 {{landingpad instruction expects at least one clause or cleanup attribute}}926 %2 = llvm.landingpad : !llvm.struct<(ptr, i32)>927 llvm.return %0 : i32928}929 930// -----931 932llvm.func @foo(i32) -> i32933llvm.func @__gxx_personality_v0(...) -> i32934 935// expected-error@below {{'llvm.resume' should have a consistent input type inside a function}}936llvm.func @caller(%arg0: i32) -> i32 attributes { personality = @__gxx_personality_v0 } {937 %0 = llvm.invoke @foo(%arg0) to ^bb1 unwind ^bb2 : (i32) -> i32938^bb1:939 %1 = llvm.invoke @foo(%0) to ^bb3 unwind ^bb4 : (i32) -> i32940^bb2:941 %2 = llvm.landingpad cleanup : !llvm.struct<(ptr, i32)>942 llvm.resume %0 : i32943^bb3:944 llvm.return %1 : i32945^bb4:946 %3 = llvm.landingpad cleanup : !llvm.struct<(ptr, i32)>947 llvm.resume %3 : !llvm.struct<(ptr, i32)>948}949 950// -----951 952llvm.func @foo(i32) -> i32953llvm.func @__gxx_personality_v0(...) -> i32954 955// expected-error@below {{'llvm.landingpad' should have a consistent result type inside a function}}956llvm.func @caller(%arg0: i32) -> i32 attributes { personality = @__gxx_personality_v0 } {957 %0 = llvm.invoke @foo(%arg0) to ^bb1 unwind ^bb2 : (i32) -> i32958^bb1:959 %1 = llvm.invoke @foo(%0) to ^bb3 unwind ^bb4 : (i32) -> i32960^bb2:961 %2 = llvm.landingpad cleanup : !llvm.struct<(ptr, i32)>962 llvm.resume %2 : !llvm.struct<(ptr, i32)>963^bb3:964 llvm.return %1 : i32965^bb4:966 %3 = llvm.landingpad cleanup : i32967 llvm.resume %3 : i32968}969 970// -----971 972llvm.func @foo(i32) -> i32973 974llvm.func @caller(%arg0: i32) -> i32 {975 %0 = llvm.mlir.constant(1 : i32) : i32976 %1 = llvm.invoke @foo(%0) to ^bb1 unwind ^bb2 : (i32) -> i32977^bb1: // pred: ^bb0978 llvm.return %0 : i32979^bb2: // pred: ^bb0980 // expected-error@+1 {{llvm.landingpad needs to be in a function with a personality}}981 %2 = llvm.landingpad cleanup : !llvm.struct<(ptr, i32)>982 llvm.resume %2 : !llvm.struct<(ptr, i32)>983}984 985// -----986 987func.func @invalid_ordering_in_fence() {988 // expected-error @+1 {{can be given only acquire, release, acq_rel, and seq_cst orderings}}989 llvm.fence syncscope("agent") monotonic990}991 992// -----993 994// expected-error @+1 {{invalid data layout descriptor}}995module attributes {llvm.data_layout = "#vjkr32"} {996 func.func @invalid_data_layout()997}998 999// -----1000 1001func.func @switch_superfluous_comma(%arg0 : i64) {1002 // expected-error@+3 {{custom op 'llvm.switch' expected integer value}}1003 llvm.switch %arg0 : i32, ^bb1 [1004 42: ^bb2,1005 ]1006^bb1:1007 llvm.return1008^bb2:1009 llvm.return1010}1011 1012// -----1013 1014func.func @switch_wrong_number_of_weights(%arg0 : i32) {1015 // expected-error@+1 {{expects number of branch weights to match number of successors: 3 vs 2}}1016 llvm.switch %arg0 : i32, ^bb1 [1017 42: ^bb2(%arg0, %arg0 : i32, i32)1018 ] {branch_weights = array<i32: 13, 17, 19>}1019 1020^bb1: // pred: ^bb01021 llvm.return1022 1023^bb2(%1: i32, %2: i32): // pred: ^bb01024 llvm.return1025}1026 1027// -----1028 1029func.func @switch_case_type_mismatch(%arg0 : i64) {1030 // expected-error@below {{expects case value type to match condition value type}}1031 "llvm.switch"(%arg0)[^bb1, ^bb2] <{case_operand_segments = array<i32: 0>, case_values = dense<42> : vector<1xi32>, operandSegmentSizes = array<i32: 1, 0, 0>}> : (i64) -> ()1032^bb1: // pred: ^bb01033 llvm.return1034^bb2: // pred: ^bb01035 llvm.return1036}1037 1038// -----1039 1040// expected-error@below {{expected zero value for 'common' linkage}}1041llvm.mlir.global common @non_zero_global_common_linkage(42 : i32) : i321042 1043// -----1044 1045// expected-error@below {{expected zero value for 'common' linkage}}1046llvm.mlir.global common @non_zero_compound_global_common_linkage(dense<[0, 0, 0, 1, 0]> : vector<5xi32>) : !llvm.array<5 x i32>1047 1048// -----1049 1050// expected-error@below {{expected array type for 'appending' linkage}}1051llvm.mlir.global appending @non_array_type_global_appending_linkage() : i321052 1053// -----1054 1055module {1056 llvm.func @accessGroups(%arg0 : !llvm.ptr) {1057 // expected-error@below {{attribute 'access_groups' failed to satisfy constraint: LLVM dialect access group metadata array}}1058 %0 = llvm.load %arg0 { "access_groups" = [@func1] } : !llvm.ptr -> i321059 llvm.return1060 }1061 llvm.func @func1() {1062 llvm.return1063 }1064}1065 1066// -----1067 1068module {1069 llvm.func @accessGroups(%arg0 : !llvm.ptr, %arg1 : i32, %arg2 : i32) {1070 // expected-error@below {{attribute 'access_groups' failed to satisfy constraint: LLVM dialect access group metadata array}}1071 %0 = llvm.cmpxchg %arg0, %arg1, %arg2 acq_rel monotonic { "access_groups" = [@metadata::@scope] } : !llvm.ptr, i321072 llvm.return1073 }1074 llvm.metadata @metadata {1075 llvm.func @scope()1076 }1077}1078 1079// -----1080 1081module {1082 llvm.func @aliasScope(%arg0 : !llvm.ptr, %arg1 : i32, %arg2 : i32) {1083 // expected-error@below {{attribute 'alias_scopes' failed to satisfy constraint: LLVM dialect alias scope array}}1084 %0 = llvm.cmpxchg %arg0, %arg1, %arg2 acq_rel monotonic { "alias_scopes" = "test" } : !llvm.ptr, i321085 llvm.return1086 }1087}1088 1089// -----1090 1091module {1092 llvm.func @noAliasScopes(%arg0 : !llvm.ptr) {1093 // expected-error@below {{attribute 'noalias_scopes' failed to satisfy constraint: LLVM dialect alias scope array}}1094 %0 = llvm.load %arg0 { "noalias_scopes" = "test" } : !llvm.ptr -> i321095 llvm.return1096 }1097}1098 1099// -----1100 1101llvm.func @wmmaLoadOp_invalid_mem_space(%arg0: !llvm.ptr<5>, %arg1: i32) {1102 // expected-error@+1 {{'nvvm.wmma.load' op expected source pointer in memory space 0, 1, 3}}1103 %0 = nvvm.wmma.load %arg0, %arg11104 {eltype = #nvvm.mma_type<f16>, frag = #nvvm.mma_frag<a>, k = 16 : i32, layout = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}1105 : (!llvm.ptr<5>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>1106 llvm.return1107}1108 1109// -----1110 1111llvm.func @wmmaLoadOp_invalid_AOp(%arg0: !llvm.ptr<3>, %arg1: i32) {1112 // expected-error@+1 {{'nvvm.wmma.load' op expected destination type is a structure of 8 elements of type 'vector<2xf16>'}}1113 %0 = nvvm.wmma.load %arg0, %arg11114 {eltype = #nvvm.mma_type<f16>, frag = #nvvm.mma_frag<a>, k = 16 : i32, layout = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}1115 : (!llvm.ptr<3>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>1116 llvm.return1117}1118 1119// -----1120 1121llvm.func @wmmaLoadOp_invalid_BOp(%arg0: !llvm.ptr<3>, %arg1: i32) {1122 // expected-error@+1 {{'nvvm.wmma.load' op expected destination type is a structure of 8 elements of type 'vector<2xf16>'}}1123 %0 = nvvm.wmma.load %arg0, %arg11124 {eltype = #nvvm.mma_type<f16>, frag = #nvvm.mma_frag<b>, k = 16 : i32, layout = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}1125 : (!llvm.ptr<3>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>1126 1127 llvm.return1128}1129 1130// -----1131 1132llvm.func @wmmaLoadOp_invalid_COp(%arg0: !llvm.ptr<3>, %arg1: i32) {1133 // expected-error@+1 {{'nvvm.wmma.load' op expected destination type is a structure of 4 elements of type 'vector<2xf16>'}}1134 %0 = nvvm.wmma.load %arg0, %arg11135 {eltype = #nvvm.mma_type<f16>, frag = #nvvm.mma_frag<c>, k = 16 : i32, layout = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}1136 : (!llvm.ptr<3>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>)>1137 1138 llvm.return1139}1140 1141// -----1142 1143llvm.func @wmmaStoreOp_invalid_mem_space(%arg0: !llvm.ptr<5>, %arg1: i32,1144 %arg2: vector<2 x f16>, %arg3: vector<2 x f16>,1145 %arg4: vector<2 x f16>, %arg5: vector<2 xf16>) {1146 // expected-error@+1 {{'nvvm.wmma.store' op expected operands to be a source pointer in memory space 0, 1, 3}}1147 nvvm.wmma.store %arg0, %arg1, %arg2, %arg3, %arg4, %arg51148 {eltype = #nvvm.mma_type<f16>, k = 16 : i32, layout = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}1149 : !llvm.ptr<5>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>1150 llvm.return1151}1152 1153// -----1154 1155llvm.func @gpu_wmma_mma_op_invalid_operands(%arg0: vector<2 x f16>, %arg1: vector<2 x f16>,1156 %arg2: vector<2 x f16>, %arg3: vector<2 x f16>,1157 %arg4: vector<2 x f16>, %arg5: vector<2 x f16>,1158 %arg6: vector<2 x f16>, %arg7: vector<2 x f16>,1159 %arg8: vector<2 x f16>, %arg9: vector<2 x f16>,1160 %arg10: vector<2 x f16>, %arg11: vector<2 x f16>,1161 %arg12: vector<2 x f16>, %arg13: vector<2 x f16>,1162 %arg14: vector<2 x f16>, %arg15: vector<2 x f16>,1163 %arg16: vector<2 x f16>, %arg17: vector<2 x f16>,1164 %arg18: vector<2 x f16>) {1165 // expected-error@+1 {{'nvvm.wmma.mma' op expected 20 arguments}}1166 %0 = nvvm.wmma.mma %arg0, %arg1, %arg2, %arg3, %arg4, %arg5, %arg6, %arg7, %arg8, %arg9, %arg10, %arg11, %arg12, %arg13, %arg14, %arg15, %arg16, %arg17, %arg181167 {eltypeA = #nvvm.mma_type<f16>, eltypeB = #nvvm.mma_type<f16>, k = 16 : i32, layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}1168 : (vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>,1169 vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>)1170 -> !llvm.struct<(vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>)>1171 llvm.return1172}1173 1174// -----1175 1176llvm.func @gpu_wmma_mma_op_results(%arg0: vector<2 x f16>, %arg1: vector<2 x f16>,1177 %arg2: vector<2 x f16>, %arg3: vector<2 x f16>,1178 %arg4: vector<2 x f16>, %arg5: vector<2 x f16>,1179 %arg6: vector<2 x f16>, %arg7: vector<2 x f16>,1180 %arg8: vector<2 x f16>, %arg9: vector<2 x f16>,1181 %arg10: vector<2 x f16>, %arg11: vector<2 x f16>,1182 %arg12: vector<2 x f16>, %arg13: vector<2 x f16>,1183 %arg14: vector<2 x f16>, %arg15: vector<2 x f16>,1184 %arg16: vector<2 x f16>, %arg17: vector<2 x f16>,1185 %arg18: vector<2 x f16>, %arg19: vector<2 x f16>) {1186 // expected-error@+1 {{'nvvm.wmma.mma' op expected destination type is a structure of 4 elements of type 'vector<2xf16>'}}1187 %0 = nvvm.wmma.mma %arg0, %arg1, %arg2, %arg3, %arg4, %arg5, %arg6, %arg7, %arg8, %arg9, %arg10, %arg11, %arg12, %arg13, %arg14, %arg15, %arg16, %arg17, %arg18, %arg191188 {eltypeA = #nvvm.mma_type<f16>, eltypeB = #nvvm.mma_type<f16>, k = 16 : i32, layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}1189 : (vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>,1190 vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>)1191 -> !llvm.struct<(vector<2 x f16>, vector<2 x f16>, vector<2 x f16>)> llvm.return1192}1193 1194// -----1195 1196llvm.func @gpu_wmma_mma_op_invalid_ab_operands(%arg0: vector<2 x f16>, %arg1: vector<2 x f16>,1197 %arg2: vector<2 x f16>, %arg3: vector<2 x f16>,1198 %arg4: vector<2 x f16>, %arg5: vector<2 x f16>,1199 %arg6: vector<2 x f16>, %arg7: vector<2 x f16>,1200 %arg8: vector<2 x f16>, %arg9: vector<2 x f16>,1201 %arg10: vector<2 x f16>, %arg11: vector<2 x f16>,1202 %arg12: vector<2 x f16>, %arg13: vector<2 x f16>,1203 %arg14: vector<2 x f16>, %arg15: f32,1204 %arg16: f32, %arg17: f32, %arg18: f32, %arg19: f32,1205 %arg20: f32, %arg21: f32, %arg22: f32, %arg23: f32) {1206 // expected-error@+1 {{'nvvm.wmma.mma' op expected argument 15 to be of type 'vector<2xf16>'}}1207 %0 = nvvm.wmma.mma %arg0, %arg1, %arg2, %arg3, %arg4, %arg5, %arg6, %arg7, %arg8, %arg9, %arg10, %arg11, %arg12, %arg13, %arg14, %arg15, %arg16, %arg17, %arg18, %arg19, %arg20, %arg21, %arg22, %arg231208 {eltypeA = #nvvm.mma_type<f16>, eltypeB = #nvvm.mma_type<f32>, k = 16 : i32, layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}1209 : (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)>1210 llvm.return1211}1212 1213// -----1214 1215llvm.func @gpu_wmma_mma_op_invalid_c_operand(%arg0: vector<2 x f16>, %arg1: vector<2 x f16>,1216 %arg2: vector<2 x f16>, %arg3: vector<2 x f16>,1217 %arg4: vector<2 x f16>, %arg5: vector<2 x f16>,1218 %arg6: vector<2 x f16>, %arg7: vector<2 x f16>,1219 %arg8: vector<2 x f16>, %arg9: vector<2 x f16>,1220 %arg10: vector<2 x f16>, %arg11: vector<2 x f16>,1221 %arg12: vector<2 x f16>, %arg13: vector<2 x f16>,1222 %arg14: vector<2 x f16>, %arg15: vector<2xf16>,1223 %arg16: f32, %arg17: f32, %arg18: f32, %arg19: f32,1224 %arg20: f32, %arg21: f32, %arg22: f32, %arg23: vector<2xf16>) {1225 // expected-error@+1 {{'nvvm.wmma.mma' op expected argument 23 to be of type 'f32'}}1226 %0 = nvvm.wmma.mma %arg0, %arg1, %arg2, %arg3, %arg4, %arg5, %arg6, %arg7, %arg8, %arg9, %arg10, %arg11, %arg12, %arg13, %arg14, %arg15, %arg16, %arg17, %arg18, %arg19, %arg20, %arg21, %arg22, %arg231227 {eltypeA = #nvvm.mma_type<f16>, eltypeB = #nvvm.mma_type<f32>, k = 16 : i32, layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}1228 : (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)>1229 llvm.return1230}1231 1232// -----1233 1234llvm.func @gpu_wmma_mma_op_invalid_result(%arg0: vector<2 x f16>, %arg1: vector<2 x f16>,1235 %arg2: vector<2 x f16>, %arg3: vector<2 x f16>,1236 %arg4: vector<2 x f16>, %arg5: vector<2 x f16>,1237 %arg6: vector<2 x f16>, %arg7: vector<2 x f16>,1238 %arg8: vector<2 x f16>, %arg9: vector<2 x f16>,1239 %arg10: vector<2 x f16>, %arg11: vector<2 x f16>,1240 %arg12: vector<2 x f16>, %arg13: vector<2 x f16>,1241 %arg14: vector<2 x f16>, %arg15: vector<2xf16>,1242 %arg16: f32, %arg17: f32, %arg18: f32, %arg19: f32,1243 %arg20: f32, %arg21: f32, %arg22: f32, %arg23: f32) {1244 // expected-error@+1 {{'nvvm.wmma.mma' op expected destination type is a structure of 8 elements of type 'f32'}}1245 %0 = nvvm.wmma.mma %arg0, %arg1, %arg2, %arg3, %arg4, %arg5, %arg6, %arg7, %arg8, %arg9, %arg10, %arg11, %arg12, %arg13, %arg14, %arg15, %arg16, %arg17, %arg18, %arg19, %arg20, %arg21, %arg22, %arg231246 {eltypeA = #nvvm.mma_type<f16>, eltypeB = #nvvm.mma_type<f32>, k = 16 : i32, layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}1247 : (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>)>1248 llvm.return1249}1250 1251// -----1252 1253llvm.func @caller() {1254 // expected-error @below {{expected function call to produce a value}}1255 llvm.call @callee() : () -> ()1256 llvm.return1257}1258 1259llvm.func @callee() -> i321260 1261// -----1262 1263llvm.func @caller() {1264 // expected-error @below {{calling function with void result must not produce values}}1265 %0 = llvm.call @callee() : () -> i321266 llvm.return1267}1268 1269llvm.func @callee() -> ()1270 1271// -----1272 1273llvm.func @caller() {1274 // expected-error @below {{expected function with 0 or 1 result}}1275 %0:2 = llvm.call @callee() : () -> (i32, f32)1276 llvm.return1277}1278 1279llvm.func @callee() -> !llvm.struct<(i32, f32)>1280 1281// -----1282 1283func.func @bitcast(%arg0: vector<2x3xf32>) {1284 // expected-error @below {{op operand #0 must be LLVM-compatible non-aggregate type}}1285 llvm.bitcast %arg0 : vector<2x3xf32> to vector<2x3xi32>1286 return1287}1288 1289// -----1290 1291func.func @cp_async(%arg0: !llvm.ptr<3>, %arg1: !llvm.ptr<1>) {1292 // expected-error @below {{expected byte size to be either 4, 8 or 16.}}1293 nvvm.cp.async.shared.global %arg0, %arg1, 32, cache = cg : !llvm.ptr<3>, !llvm.ptr<1>1294 return1295}1296 1297// -----1298 1299func.func @cp_async(%arg0: !llvm.ptr<3>, %arg1: !llvm.ptr<1>) {1300 // expected-error @below {{CG cache modifier is only support for 16 bytes copy.}}1301 nvvm.cp.async.shared.global %arg0, %arg1, 8, cache = cg : !llvm.ptr<3>, !llvm.ptr<1>1302 return1303}1304 1305// -----1306 1307func.func @mapa(%a: !llvm.ptr, %b : i32) {1308 // expected-error @below {{'nvvm.mapa' op failed to verify that Valid address-space check(or mapping) for mapa Op}}1309 %0 = nvvm.mapa %a, %b: !llvm.ptr -> !llvm.ptr<7>1310 return1311}1312 1313// -----1314 1315func.func @gep_struct_variable(%arg0: !llvm.ptr, %arg1: i32, %arg2: i32) {1316 // expected-error @below {{op expected index 1 indexing a struct to be constant}}1317 llvm.getelementptr %arg0[%arg1, %arg1] : (!llvm.ptr, i32, i32) -> !llvm.ptr, !llvm.struct<(i32)>1318 return1319}1320 1321// -----1322 1323func.func @gep_out_of_bounds(%ptr: !llvm.ptr, %idx: i64) {1324 // expected-error @below {{index 2 indexing a struct is out of bounds}}1325 llvm.getelementptr %ptr[%idx, 1, 3] : (!llvm.ptr, i64) -> !llvm.ptr, !llvm.struct<(i32, struct<(i32, f32)>)>1326 return1327}1328 1329// -----1330 1331func.func @non_splat_shuffle_on_scalable_vector(%arg0: vector<[4]xf32>) {1332 // expected-error@below {{expected a splat operation for scalable vectors}}1333 %0 = llvm.shufflevector %arg0, %arg0 [0, 0, 0, 1] : vector<[4]xf32>1334 return1335}1336 1337// -----1338 1339llvm.mlir.global internal @side_effecting_global() : !llvm.struct<(i8)> {1340 %0 = llvm.mlir.constant(1 : i64) : i641341 // expected-error@below {{ops with side effects not allowed in global initializers}}1342 %1 = llvm.alloca %0 x !llvm.struct<(i8)> : (i64) -> !llvm.ptr1343 %2 = llvm.load %1 : !llvm.ptr -> !llvm.struct<(i8)>1344 llvm.return %2 : !llvm.struct<(i8)>1345}1346 1347// -----1348 1349func.func @insert_vector_invalid_source_vector_size(%arg0 : vector<16385xi8>, %arg1 : vector<[16]xi8>) {1350 // expected-error@+1 {{op failed to verify that vectors are not bigger than 2^17 bits.}}1351 %0 = llvm.intr.vector.insert %arg0, %arg1[0] : vector<16385xi8> into vector<[16]xi8>1352}1353 1354// -----1355 1356func.func @insert_vector_invalid_dest_vector_size(%arg0 : vector<16xi8>, %arg1 : vector<[16385]xi8>) {1357 // expected-error@+1 {{op failed to verify that vectors are not bigger than 2^17 bits.}}1358 %0 = llvm.intr.vector.insert %arg0, %arg1[0] : vector<16xi8> into vector<[16385]xi8>1359}1360 1361// -----1362 1363func.func @insert_scalable_into_fixed_length_vector(%arg0 : vector<[8]xf32>, %arg1 : vector<16xf32>) {1364 // expected-error@+1 {{op failed to verify that it is not inserting scalable into fixed-length vectors.}}1365 %0 = llvm.intr.vector.insert %arg0, %arg1[0] : vector<[8]xf32> into vector<16xf32>1366}1367 1368// -----1369 1370func.func @extract_vector_invalid_source_vector_size(%arg0 : vector<[16385]xi8>) {1371 // expected-error@+1 {{op failed to verify that vectors are not bigger than 2^17 bits.}}1372 %0 = llvm.intr.vector.extract %arg0[0] : vector<16xi8> from vector<[16385]xi8>1373}1374 1375// -----1376 1377func.func @extract_vector_invalid_result_vector_size(%arg0 : vector<[16]xi8>) {1378 // expected-error@+1 {{op failed to verify that vectors are not bigger than 2^17 bits.}}1379 %0 = llvm.intr.vector.extract %arg0[0] : vector<16385xi8> from vector<[16]xi8>1380}1381 1382// -----1383 1384func.func @extract_scalable_from_fixed_length_vector(%arg0 : vector<16xf32>) {1385 // expected-error@+1 {{op failed to verify that it is not extracting scalable from fixed-length vectors.}}1386 %0 = llvm.intr.vector.extract %arg0[0] : vector<[8]xf32> from vector<16xf32>1387}1388 1389 1390// -----1391 1392func.func @vector_interleave2_bad_type0(%vec1: vector<[2]xf16>, %vec2 : vector<[4]xf16>) {1393 // expected-error@+1 {{op failed to verify that all of {vec1, vec2} have same type}}1394 %0 = "llvm.intr.vector.interleave2"(%vec1, %vec2) : (vector<[2]xf16>, vector<[4]xf16>) -> vector<[8]xf16>1395 return1396}1397 1398// -----1399 1400func.func @vector_interleave2_bad_type1(%vec1: vector<[2]xf16>, %vec2 : vector<[2]xf16>) {1401 // expected-error@+1 {{op failed to verify that result has twice as many elements as 'vec1'}}1402 %0 = "llvm.intr.vector.interleave2"(%vec1, %vec2) : (vector<[2]xf16>, vector<[2]xf16>) -> vector<[8]xf16>1403 return1404}1405 1406// -----1407 1408/// result vector type is not scalable.1409 1410func.func @vector_interleave2_bad_type2(%vec1: vector<[2]xf16>, %vec2 : vector<[2]xf16>) {1411 // expected-error@+1 {{op failed to verify that result has twice as many elements as 'vec1'}}1412 %0 = "llvm.intr.vector.interleave2"(%vec1, %vec2) : (vector<[2]xf16>, vector<[2]xf16>) -> vector<4xf16>1413 return1414}1415 1416// -----1417 1418 1419/// element type doesn't match.1420 1421func.func @vector_interleave2_bad_type3(%vec1: vector<[2]xf16>, %vec2 : vector<[2]xf16>) {1422 // expected-error@+1 {{op failed to verify that result has twice as many elements as 'vec1'}}1423 %0 = "llvm.intr.vector.interleave2"(%vec1, %vec2) : (vector<[2]xf16>, vector<[2]xf16>) -> vector<[4]xf32>1424 return1425}1426 1427// -----1428 1429func.func @invalid_bitcast_ptr_to_i64(%arg : !llvm.ptr) {1430 // expected-error@+1 {{can only cast pointers from and to pointers}}1431 %1 = llvm.bitcast %arg : !llvm.ptr to i641432}1433 1434// -----1435 1436func.func @invalid_bitcast_i64_to_ptr() {1437 %0 = llvm.mlir.constant(2 : i64) : i641438 // expected-error@+1 {{can only cast pointers from and to pointers}}1439 %1 = llvm.bitcast %0 : i64 to !llvm.ptr1440}1441 1442// -----1443 1444func.func @invalid_bitcast_vec_to_ptr(%arg : vector<4x!llvm.ptr>) {1445 // expected-error@+1 {{cannot cast vector of pointers to pointer}}1446 %0 = llvm.bitcast %arg : vector<4x!llvm.ptr> to !llvm.ptr1447}1448 1449// -----1450 1451func.func @invalid_bitcast_ptr_to_vec(%arg : !llvm.ptr) {1452 // expected-error@+1 {{cannot cast pointer to vector of pointers}}1453 %0 = llvm.bitcast %arg : !llvm.ptr to vector<4x!llvm.ptr>1454}1455 1456// -----1457 1458func.func @invalid_bitcast_addr_cast(%arg : !llvm.ptr<1>) {1459 // expected-error@+1 {{cannot cast pointers of different address spaces, use 'llvm.addrspacecast' instead}}1460 %0 = llvm.bitcast %arg : !llvm.ptr<1> to !llvm.ptr1461}1462 1463// -----1464 1465func.func @invalid_bitcast_addr_cast_vec(%arg : vector<4x!llvm.ptr<1>>) {1466 // expected-error@+1 {{cannot cast pointers of different address spaces, use 'llvm.addrspacecast' instead}}1467 %0 = llvm.bitcast %arg : vector<4x!llvm.ptr<1>> to vector<4x!llvm.ptr>1468}1469 1470// -----1471 1472func.func @invalid_target_ext_alloca() {1473 %0 = llvm.mlir.constant(1 : i64) : i641474 // expected-error@+1 {{this target extension type cannot be used in alloca}}1475 %1 = llvm.alloca %0 x !llvm.target<"no_alloca"> : (i64) -> !llvm.ptr1476}1477 1478// -----1479 1480func.func @invalid_target_ext_load(%arg0 : !llvm.ptr) {1481 // expected-error@+1 {{result #0 must be LLVM type with size, but got '!llvm.target<"no_load">'}}1482 %0 = llvm.load %arg0 {alignment = 8 : i64} : !llvm.ptr -> !llvm.target<"no_load">1483}1484 1485// -----1486 1487func.func @invalid_target_ext_atomic(%arg0 : !llvm.ptr) {1488 // expected-error@+1 {{unsupported type '!llvm.target<"spirv.Event">' for atomic access}}1489 %0 = llvm.load %arg0 atomic monotonic {alignment = 8 : i64} : !llvm.ptr -> !llvm.target<"spirv.Event">1490}1491 1492// -----1493 1494func.func @invalid_target_ext_constant_unsupported() {1495 // expected-error@+1 {{target extension type does not support zero-initializer}}1496 %0 = llvm.mlir.zero : !llvm.target<"invalid_constant">1497 llvm.return1498}1499 1500// -----1501 1502func.func @invalid_target_ext_constant() {1503 // expected-error@+1 {{does not support target extension type.}}1504 %0 = llvm.mlir.constant(0 : index) : !llvm.target<"spirv.Event">1505 llvm.return1506}1507 1508// -----1509 1510llvm.comdat @__llvm_comdat {1511 // expected-error@below {{only comdat selector symbols can appear in a comdat region}}1512 llvm.return1513}1514 1515// -----1516 1517llvm.mlir.global @not_comdat(0 : i32) : i321518// expected-error@below {{expected comdat symbol}}1519llvm.mlir.global @invalid_global_comdat(0 : i32) comdat(@not_comdat) : i321520 1521// -----1522 1523// expected-error@below {{expected comdat symbol}}1524llvm.func @invalid_func_comdat() comdat(@foo) {1525 llvm.return1526}1527 1528// -----1529 1530func.func @invalid_zext_target_size_equal(%arg: i32) {1531 // expected-error@+1 {{integer width of the output type is smaller or equal to the integer width of the input type}}1532 %0 = llvm.zext %arg : i32 to i321533}1534 1535// -----1536 1537func.func @invalid_zext_target_size(%arg: i32) {1538 // expected-error@+1 {{integer width of the output type is smaller or equal to the integer width of the input type}}1539 %0 = llvm.zext %arg : i32 to i161540}1541 1542// -----1543 1544func.func @invalid_zext_target_size_vector(%arg: vector<1xi32>) {1545 // expected-error@+1 {{integer width of the output type is smaller or equal to the integer width of the input type}}1546 %0 = llvm.zext %arg : vector<1xi32> to vector<1xi16>1547}1548 1549// -----1550 1551func.func @invalid_zext_target_shape(%arg: vector<1xi32>) {1552 // expected-error@+1 {{input and output vectors are of incompatible shape}}1553 %0 = llvm.zext %arg : vector<1xi32> to vector<2xi64>1554}1555 1556// -----1557 1558func.func @invalid_zext_target_type(%arg: i32) {1559 // expected-error@+1 {{input type is an integer but output type is a vector}}1560 %0 = llvm.zext %arg : i32 to vector<1xi64>1561}1562 1563// -----1564 1565func.func @invalid_zext_target_type_two(%arg: vector<1xi32>) {1566 // expected-error@+1 {{input type is a vector but output type is an integer}}1567 %0 = llvm.zext %arg : vector<1xi32> to i641568}1569 1570// -----1571 1572llvm.func @non_variadic(%arg: i32)1573 1574llvm.func @invalid_var_callee_type(%arg: i32) {1575 // expected-error@below {{expected var_callee_type to be a variadic function type}}1576 llvm.call @non_variadic(%arg) vararg(!llvm.func<void (i32)>) : (i32) -> ()1577 llvm.return1578}1579 1580// -----1581 1582llvm.func @variadic(%arg: i32, ...)1583 1584llvm.func @invalid_var_callee_type_num_parameters(%arg: i32) {1585 // expected-error@below {{expected var_callee_type to have at most 1 parameters}}1586 llvm.call @variadic(%arg) vararg(!llvm.func<void (i32, i64, ...)>) : (i32) -> ()1587 llvm.return1588}1589 1590// -----1591 1592llvm.func @invalid_var_callee_type_num_parameters_indirect(%callee : !llvm.ptr, %arg: i32) {1593 // expected-error@below {{expected var_callee_type to have at most 1 parameters}}1594 llvm.call %callee(%arg) vararg(!llvm.func<void (i32, i64, ...)>) : !llvm.ptr, (i32) -> ()1595 llvm.return1596}1597 1598// -----1599 1600llvm.func @variadic(%arg: i32, ...)1601 1602llvm.func @invalid_var_callee_type_parameter_type_mismatch(%arg: i32) {1603 // expected-error@below {{var_callee_type parameter type mismatch: 'i64' != 'i32'}}1604 llvm.call @variadic(%arg) vararg(!llvm.func<void (i64, ...)>) : (i32) -> ()1605 llvm.return1606}1607 1608// -----1609 1610llvm.func @invalid_var_callee_type_parameter_type_mismatch_indirect(%callee : !llvm.ptr, %arg: i32) {1611 // expected-error@below {{var_callee_type parameter type mismatch: 'i64' != 'i32'}}1612 llvm.call %callee(%arg) vararg(!llvm.func<void (i64, ...)>) : !llvm.ptr, (i32) -> ()1613 llvm.return1614}1615 1616// -----1617 1618llvm.func @variadic(%arg: i32, ...)1619 1620llvm.func @invalid_var_callee_type_non_void(%arg: i32) {1621 // expected-error@below {{expected var_callee_type to return void}}1622 llvm.call @variadic(%arg) vararg(!llvm.func<i8 (i32, ...)>) : (i32) -> ()1623 llvm.return1624}1625 1626// -----1627 1628llvm.func @variadic(%arg: i32, ...) -> i321629 1630llvm.func @invalid_var_callee_type_return_type_mismatch(%arg: i32) {1631 // expected-error@below {{var_callee_type return type mismatch: 'i8' != 'i32'}}1632 %0 = llvm.call @variadic(%arg) vararg(!llvm.func<i8 (i32, ...)>) : (i32) -> (i32)1633 llvm.return1634}1635 1636// -----1637 1638llvm.func @non_variadic(%arg: i32)1639 1640llvm.func @invalid_var_callee_type(%arg: i32) {1641 // expected-error@below {{expected var_callee_type to be a variadic function type}}1642 llvm.invoke @non_variadic(%arg) to ^bb2 unwind ^bb1 vararg(!llvm.func<void (i32)>) : (i32) -> ()1643^bb1:1644 llvm.return1645^bb2:1646 llvm.return1647}1648 1649// -----1650 1651llvm.func @variadic(%arg: i32, ...)1652 1653llvm.func @invalid_var_callee_type_num_parameters(%arg: i32) {1654 // expected-error@below {{expected var_callee_type to have at most 1 parameters}}1655 llvm.invoke @variadic(%arg) to ^bb2 unwind ^bb1 vararg(!llvm.func<void (i32, i64, ...)>) : (i32) -> ()1656^bb1:1657 llvm.return1658^bb2:1659 llvm.return1660}1661 1662// -----1663 1664llvm.func @invalid_var_callee_type_num_parameters_indirect(%callee : !llvm.ptr, %arg: i32) {1665 // expected-error@below {{expected var_callee_type to have at most 1 parameters}}1666 llvm.invoke %callee(%arg) to ^bb2 unwind ^bb1 vararg(!llvm.func<void (i32, i64, ...)>) : !llvm.ptr, (i32) -> ()1667^bb1:1668 llvm.return1669^bb2:1670 llvm.return1671}1672 1673// -----1674 1675llvm.func @variadic(%arg: i32, ...)1676 1677llvm.func @invalid_var_callee_type_parameter_type_mismatch(%arg: i32) {1678 // expected-error@below {{var_callee_type parameter type mismatch: 'i64' != 'i32'}}1679 llvm.invoke @variadic(%arg) to ^bb2 unwind ^bb1 vararg(!llvm.func<void (i64, ...)>) : (i32) -> ()1680^bb1:1681 llvm.return1682^bb2:1683 llvm.return1684}1685 1686// -----1687 1688llvm.func @invalid_var_callee_type_parameter_type_mismatch_indirect(%callee : !llvm.ptr, %arg: i32) {1689 // expected-error@below {{var_callee_type parameter type mismatch: 'i64' != 'i32'}}1690 llvm.invoke %callee(%arg) to ^bb2 unwind ^bb1 vararg(!llvm.func<void (i64, ...)>) : !llvm.ptr, (i32) -> ()1691^bb1:1692 llvm.return1693^bb2:1694 llvm.return1695}1696 1697// -----1698 1699llvm.func @variadic(%arg: i32, ...)1700 1701llvm.func @invalid_var_callee_type_non_void(%arg: i32) {1702 // expected-error@below {{expected var_callee_type to return void}}1703 llvm.invoke @variadic(%arg) to ^bb2 unwind ^bb1 vararg(!llvm.func<i8 (i32, ...)>) : (i32) -> ()1704^bb1:1705 llvm.return1706^bb2:1707 llvm.return1708}1709 1710// -----1711 1712llvm.func @variadic(%arg: i32, ...) -> i321713 1714llvm.func @invalid_var_callee_type_return_type_mismatch(%arg: i32) {1715 // expected-error@below {{var_callee_type return type mismatch: 'i8' != 'i32'}}1716 %0 = llvm.invoke @variadic(%arg) to ^bb2 unwind ^bb1 vararg(!llvm.func<i8 (i32, ...)>) : (i32) -> (i32)1717^bb1:1718 llvm.return1719^bb2:1720 llvm.return1721}1722 1723// -----1724 1725llvm.func @variadic(...)1726 1727llvm.func @invalid_variadic_call(%arg: i32) {1728 // expected-error@+1 {{missing var_callee_type attribute for vararg call}}1729 "llvm.call"(%arg) <{callee = @variadic}> {operandSegmentSizes = array<i32: 1, 0>, op_bundle_sizes = array<i32>} : (i32) -> ()1730 llvm.return1731}1732 1733// -----1734 1735llvm.func @variadic(...)1736 1737llvm.func @invalid_variadic_call(%arg: i32) {1738 // expected-error@+1 {{missing var_callee_type attribute for vararg call}}1739 "llvm.call"(%arg) <{callee = @variadic}> {operandSegmentSizes = array<i32: 1, 0>, op_bundle_sizes = array<i32>} : (i32) -> ()1740 llvm.return1741}1742 1743// -----1744 1745llvm.func @foo(%arg: !llvm.ptr) {1746 // expected-error@+1 {{type '!llvm.ptr' cannot be indexed (index #1)}}1747 %0 = llvm.getelementptr %arg[0, 4] : (!llvm.ptr) -> !llvm.ptr, !llvm.ptr1748 llvm.return1749}1750 1751// -----1752 1753// expected-error @below {{no_inline and always_inline attributes are incompatible}}1754llvm.func @alwaysinline_noinline() attributes { always_inline, no_inline } {1755 llvm.return1756}1757 1758// -----1759 1760// expected-error @below {{'llvm.func' op with optimize_none must also be no_inline}}1761llvm.func @optnone_requires_noinline() attributes { optimize_none } {1762 llvm.return1763}1764 1765// -----1766 1767llvm.func @foo()1768llvm.func @wrong_number_of_bundle_types() {1769 %0 = llvm.mlir.constant(0 : i32) : i321770 // expected-error@+1 {{expected 1 types for operand bundle operands for operand bundle #0, but actually got 2}}1771 llvm.call @foo() ["tag"(%0 : i32, i32)] : () -> ()1772 llvm.return1773}1774 1775// -----1776 1777llvm.func @wrong_number_of_bundle_types_intrin(%arg0: i32) -> i32 {1778 %0 = llvm.mlir.constant(0 : i32) : i321779 // expected-error@+1 {{expected 1 types for operand bundle operands for operand bundle #0, but actually got 2}}1780 %1 = llvm.call_intrinsic "llvm.riscv.sha256sig0"(%arg0) ["tag"(%0 : i32, i32)] : (i32 {llvm.signext}) -> (i32)1781 llvm.return %1 : i321782}1783 1784// -----1785 1786llvm.func @foo()1787llvm.func @wrong_number_of_bundle_tags() {1788 %0 = llvm.mlir.constant(0 : i32) : i321789 %1 = llvm.mlir.constant(1 : i32) : i321790 // expected-error@+1 {{expected 2 operand bundle tags, but actually got 1}}1791 "llvm.call"(%0, %1) <{ op_bundle_tags = ["tag"] }> {1792 callee = @foo,1793 operandSegmentSizes = array<i32: 0, 2>,1794 op_bundle_sizes = array<i32: 1, 1>1795 } : (i32, i32) -> ()1796 llvm.return1797}1798 1799// -----1800 1801llvm.mlir.global external @x(42 : i32) : i321802 1803// expected-error@+1 {{expects type to be a valid element type for an LLVM global alias}}1804llvm.mlir.alias external @y : !llvm.label {1805 %0 = llvm.mlir.addressof @x : !llvm.ptr1806 llvm.return %0 : !llvm.ptr1807}1808 1809// -----1810 1811llvm.mlir.global external @x(42 : i32) : i321812 1813// expected-error@+1 {{linkage not supported in aliases, available options}}1814llvm.mlir.alias appending @y2 : i32 {1815 %0 = llvm.mlir.addressof @x : !llvm.ptr1816 llvm.return %0 : !llvm.ptr1817}1818 1819// -----1820 1821// expected-error@+1 {{initializer region must always return a pointer}}1822llvm.mlir.alias external @y3 : i32 {1823 %c = llvm.mlir.constant(42 : i64) : i641824 llvm.return %c : i641825}1826 1827// -----1828 1829llvm.mlir.global external @x(42 : i32) : i321830 1831llvm.mlir.alias external @y4 : i32 {1832 %0 = llvm.mlir.addressof @x : !llvm.ptr1833 // expected-error@+1 {{ops with side effects are not allowed in alias initializers}}1834 %2 = llvm.load %0 : !llvm.ptr -> i321835 llvm.return %0 : !llvm.ptr1836}1837 1838// -----1839 1840llvm.mlir.global external @x(42 : i32) : i321841 1842llvm.mlir.alias external @y5 : i32 {1843 // expected-error@+1 {{pointer address space must match address space}}1844 %0 = llvm.mlir.addressof @x : !llvm.ptr<4>1845 llvm.return %0 : !llvm.ptr<4>1846}1847 1848// -----1849 1850module {1851 llvm.func @foo()1852 1853 // expected-error@below {{only integer and string values are currently supported}}1854 llvm.module_flags [#llvm.mlir.module_flag<error, "yolo", @foo>]1855}1856 1857// -----1858 1859module {1860 // expected-error@below {{'CG Profile' key expects an array of '#llvm.cgprofile_entry'}}1861 llvm.module_flags [#llvm.mlir.module_flag<append, "CG Profile", [1862 "yo"1863 ]>]1864}1865 1866// -----1867 1868module {1869 // expected-error@below {{'CG Profile' key expects an array of '#llvm.cgprofile_entry'}}1870 llvm.module_flags [#llvm.mlir.module_flag<append, "CG Profile", 3 : i64>]1871}1872 1873// -----1874 1875module {1876 // expected-error@below {{'ProfileSummary' key expects a '#llvm.profile_summary' attribute}}1877 llvm.module_flags [#llvm.mlir.module_flag<append, "ProfileSummary", 3 : i64>]1878}1879 1880// -----1881 1882llvm.module_flags [#llvm.mlir.module_flag<error, "ProfileSummary",1883 // expected-error@below {{expected one of [SampleProfile, InstrProf, CSInstrProf] for LLVM ProfileSummary format kinds, got: YoloFmt}}1884 #llvm.profile_summary<format = "YoloFmt", total_count = 263646, max_count = 86427,1885 // expected-error@above {{failed to parse ModuleFlagProfileSummaryAttr parameter 'format' which is to be a `ProfileSummaryFormatKind`}}1886 max_internal_count = 86427, max_function_count = 4691,1887 num_counts = 3712, num_functions = 796,1888 is_partial_profile = 0,1889 partial_profile_ratio = 0.000000e+00 : f64,1890 detailed_summary =1891 <cut_off = 10000, min_count = 86427, num_counts = 1>,1892 <cut_off = 100000, min_count = 86427, num_counts = 1>1893 // expected-error@below {{failed to parse ModuleFlagAttr parameter}}1894>>]1895 1896// -----1897 1898llvm.func @t0() -> !llvm.ptr {1899 %0 = llvm.blockaddress <function = @t0, tag = <id = 1>> : !llvm.ptr1900 llvm.blocktag <id = 1>1901 llvm.br ^bb11902^bb1:1903 // expected-error@+1 {{duplicate block tag '1' in the same function}}1904 llvm.blocktag <id = 1>1905 llvm.return %0 : !llvm.ptr1906}1907 1908// -----1909 1910llvm.func @t1() -> !llvm.ptr {1911 // expected-error@+1 {{expects an existing block label target in the referenced function}}1912 %0 = llvm.blockaddress <function = @t1, tag = <id = 1>> : !llvm.ptr1913 llvm.br ^bb11914^bb1:1915 llvm.return %0 : !llvm.ptr1916}1917 1918// -----1919 1920llvm.func @gep_inbounds_flag_usage(%ptr: !llvm.ptr, %idx: i64) {1921 // expected-error@+1 {{'inbounds_flag' cannot be used directly}}1922 llvm.getelementptr inbounds_flag %ptr[%idx, 0, %idx] : (!llvm.ptr, i64, i64) -> !llvm.ptr, !llvm.struct<(array<10 x f32>)>1923 llvm.return1924}1925 1926// -----1927 1928llvm.mlir.global @bad_struct_array_init_size() : !llvm.array<2x!llvm.struct<(i32, f32)>> {1929 // expected-error@below {{'llvm.mlir.constant' op array attribute size does not match array type size in dimension 0: 1 vs. 2}}1930 %0 = llvm.mlir.constant([[42 : i32, 1.000000e+00 : f32]]) : !llvm.array<2x!llvm.struct<(i32, f32)>>1931 llvm.return %0 : !llvm.array<2x!llvm.struct<(i32, f32)>>1932}1933 1934// -----1935 1936llvm.mlir.global @bad_struct_array_init_nesting() : !llvm.array<1x!llvm.array<1x!llvm.array<1x!llvm.struct<(i32)>>>> {1937 // expected-error@below {{'llvm.mlir.constant' op nested attribute for sub-array in dimension 1 at index 0 must be a zero, or undef, or array attribute}}1938 %0 = llvm.mlir.constant([[1 : i32]]) : !llvm.array<1x!llvm.array<1x!llvm.array<1x!llvm.struct<(i32)>>>>1939 llvm.return %0 : !llvm.array<1x!llvm.array<1x!llvm.array<1x!llvm.struct<(i32)>>>>1940}1941 1942// -----1943 1944llvm.mlir.global @bad_struct_array_init_elements() : !llvm.array<1x!llvm.struct<(i32, f32)>> {1945 // expected-error@below {{'llvm.mlir.constant' op nested array attribute size for struct element at index 0 must match struct size: 1 vs. 2}}1946 %0 = llvm.mlir.constant([[1 : i32]]) : !llvm.array<1x!llvm.struct<(i32, f32)>>1947 llvm.return %0 : !llvm.array<1x!llvm.struct<(i32, f32)>>1948}1949 1950// -----1951 1952llvm.mlir.global internal constant @bad_array_attr_simple_type() : !llvm.array<2 x f64> {1953 // expected-error@below {{'llvm.mlir.constant' op for array with an array attribute must have a struct element type}}1954 %0 = llvm.mlir.constant([2.5, 7.4]) : !llvm.array<2 x f64>1955 llvm.return %0 : !llvm.array<2 x f64>1956}1957 1958// -----1959 1960llvm.func @inlineAsmMustTail(%arg0: i32, %arg1 : !llvm.ptr) {1961 // expected-error@+1 {{op tail call kind 'musttail' is not supported}}1962 %8 = llvm.inline_asm tail_call_kind = <musttail> "foo", "=r,=r,r" %arg0 : (i32) -> !llvm.struct<(i8, i8)>1963 llvm.return1964}1965 1966// -----1967 1968llvm.func @invalid_xevm_prefetch(%arg0: !llvm.ptr) {1969 // expected-error@+1 {{op operand #0 must be LLVM pointer in address space 1 or LLVM pointer in address space 4}}1970 xevm.prefetch %arg0 <{cache_control = #xevm.load_cache_control<L1uc_L2uc_L3uc>}> : (!llvm.ptr)1971 llvm.return1972}1973 1974// -----1975llvm.func @invalid_xevm_blockload(%arg0: !llvm.ptr<1>) {1976 // expected-error@+1 {{op vector size must be 2, 4 or 8 for element type > 8 bits}}1977 %0 = xevm.blockload %arg0 : (!llvm.ptr<1>) -> vector<3xi16>1978 llvm.return1979}1980 1981// -----1982llvm.func @invalid_xevm_blockstore(%arg0: !llvm.ptr<1>, %arg1: vector<5xi8>) {1983 // expected-error@+1 {{op vector size must be 2, 4, 8 or 16 for 8-bit element type}}1984 xevm.blockstore %arg0, %arg1 : (!llvm.ptr<1>, vector<5xi8>)1985 llvm.return1986}1987 1988// -----1989 1990llvm.func @invalid_xevm_mma(%loaded_c_casted: vector<4xf32>, %loaded_a: vector<8xi16>, %loaded_b_casted: vector<8xi32>) -> vector<8xf32> {1991 // expected-error@+1 {{op type of C operand must match result type}}1992 %c_result = xevm.mma %loaded_a, %loaded_b_casted, %loaded_c_casted {shape = <m = 8, n = 16, k = 16>, types = <d = f32, a = f16, b = f16, c = f32>} : (vector<8xi16>, vector<8xi32>, vector<4xf32>) -> vector<8xf32>1993 llvm.return %c_result : vector<8xf32>1994}1995 1996// -----1997 1998llvm.func @invalid_xevm_matrix_1(%c: !llvm.ptr<1>, %base_width_c: i32, %base_height_c: i32, %base_pitch_c: i32, %x: i32, %y: i32, %c_result_casted: vector<8xi32>) {1999 // expected-error@+1 {{op expecting tile_width to be between 1 and 8}}2000 xevm.blockstore2d %c, %base_width_c, %base_height_c, %base_pitch_c, %x, %y, %c_result_casted <{elem_size_in_bits=64 : i32, tile_width=16 : i32, tile_height=8 : i32}> : (!llvm.ptr<1>, i32, i32, i32, i32, i32, vector<8xi32>)2001 llvm.return2002}2003 2004// -----2005 2006llvm.func @invalid_xevm_matrix_2(%c: !llvm.ptr<1>, %base_width_c: i32, %base_height_c: i32, %base_pitch_c: i32, %x: i32, %y: i32, %c_result_casted: vector<8xi32>) {2007 // expected-error@+1 {{op expecting elem_size_in_bits to be 8, 16, 32, or 64}}2008 xevm.blockstore2d %c, %base_width_c, %base_height_c, %base_pitch_c, %x, %y, %c_result_casted <{elem_size_in_bits=18 : i32, tile_width=16 : i32, tile_height=8 : i32}> : (!llvm.ptr<1>, i32, i32, i32, i32, i32, vector<8xi32>)2009 llvm.return2010}2011 2012// -----2013 2014llvm.func @invalid_xevm_matrix_3(%a: !llvm.ptr<1>, %base_width_a: i32, %base_height_a: i32, %base_pitch_a: i32, %x: i32, %y: i32) -> vector<8xi16> {2015 // expected-error@+1 {{op result size of 128 bits does not match the expected size of 208 bits}}2016 %loaded_a = xevm.blockload2d %a, %base_width_a, %base_height_a, %base_pitch_a, %x, %y <{elem_size_in_bits=16 : i32, tile_width=26 : i32, tile_height=8 : i32, v_blocks=1 : i32, transpose=false, pack_register=false, cache_control=#xevm.load_cache_control<L1uc_L2uc_L3uc>}> : (!llvm.ptr<1>, i32, i32, i32, i32, i32) -> vector<8xi16>2017 llvm.return %loaded_a : vector<8xi16>2018}2019 2020// -----2021 2022llvm.func external @resolve_foo() -> !llvm.ptr attributes {dso_local}2023// expected-error@+1 {{'llvm.mlir.ifunc' op resolver must be a definition}}2024llvm.mlir.ifunc external @foo : !llvm.func<void (ptr, i32)>, !llvm.ptr @resolve_foo {dso_local}2025 2026// -----2027 2028llvm.mlir.global external @resolve_foo() : !llvm.ptr2029// expected-error@+1 {{'llvm.mlir.ifunc' op must have a function resolver}}2030llvm.mlir.ifunc external @foo : !llvm.func<void (ptr, i32)>, !llvm.ptr @resolve_foo {dso_local}2031 2032// -----2033 2034llvm.func external @resolve_foo() -> !llvm.ptr2035// expected-error@+1 {{'llvm.mlir.ifunc' op 'common' linkage not supported in ifuncs}}2036llvm.mlir.ifunc common @foo : !llvm.func<void (ptr, i32)>, !llvm.ptr @resolve_foo {dso_local}2037 2038// -----2039 2040llvm.mlir.global external @resolve_foo() : !llvm.ptr2041llvm.mlir.alias external @alias_resolver : !llvm.ptr {2042 %0 = llvm.mlir.addressof @resolve_foo : !llvm.ptr2043 llvm.return %0 : !llvm.ptr2044}2045// expected-error@+1 {{'llvm.mlir.ifunc' op must have a function resolver}}2046llvm.mlir.ifunc external @foo : !llvm.func<void (ptr, i32)>, !llvm.ptr @alias_resolver {dso_local}2047 2048// -----2049 2050llvm.func @invalid_sincos_nonhomogeneous_return_type(%f: f32) -> () {2051 // expected-error@+1 {{op expected result type to be an homogeneous struct with two elements matching the operand type}}2052 llvm.intr.sincos(%f) : (f32) -> !llvm.struct<(f32, f64)>2053}2054 2055// -----2056 2057llvm.func @invalid_sincos_non_struct_return_type(%f: f32) -> () {2058 // expected-error@+1 {{op expected result type to be an homogeneous struct with two elements matching the operand type}}2059 llvm.intr.sincos(%f) : (f32) -> f322060}2061 2062// -----2063 2064llvm.func @invalid_sincos_gt_2_element_struct_return_type(%f: f32) -> () {2065 // expected-error@+1 {{op expected result type to be an homogeneous struct with two elements matching the operand type}}2066 llvm.intr.sincos(%f) : (f32) -> !llvm.struct<(f32, f32, f32)>2067}2068