1// RUN: mlir-opt -allow-unregistered-dialect -split-input-file -verify-diagnostics %s | FileCheck %s 2 3//===----------------------------------------------------------------------===// 4// spirv.mlir.addressof 5//===----------------------------------------------------------------------===// 6 7spirv.module Logical GLSL450 { 8 spirv.GlobalVariable @var1 : !spirv.ptr<!spirv.struct<(f32, !spirv.array<4xf32>)>, Input> 9 spirv.func @access_chain() -> () "None" { 10 %0 = spirv.Constant 1: i32 11 // CHECK: [[VAR1:%.*]] = spirv.mlir.addressof @var1 : !spirv.ptr<!spirv.struct<(f32, !spirv.array<4 x f32>)>, Input> 12 // CHECK-NEXT: spirv.AccessChain [[VAR1]][{{.*}}, {{.*}}] : !spirv.ptr<!spirv.struct<(f32, !spirv.array<4 x f32>)>, Input> 13 %1 = spirv.mlir.addressof @var1 : !spirv.ptr<!spirv.struct<(f32, !spirv.array<4xf32>)>, Input> 14 %2 = spirv.AccessChain %1[%0, %0] : !spirv.ptr<!spirv.struct<(f32, !spirv.array<4xf32>)>, Input>, i32, i32 -> !spirv.ptr<f32, Input> 15 spirv.Return 16 } 17} 18 19// ----- 20 21// Allow taking address of global variables in other module-like ops 22spirv.GlobalVariable @var : !spirv.ptr<!spirv.struct<(f32, !spirv.array<4xf32>)>, Input> 23func.func @addressof() -> () { 24 // CHECK: spirv.mlir.addressof @var 25 %1 = spirv.mlir.addressof @var : !spirv.ptr<!spirv.struct<(f32, !spirv.array<4xf32>)>, Input> 26 return 27} 28 29// ----- 30 31spirv.module Logical GLSL450 { 32 spirv.GlobalVariable @var1 : !spirv.ptr<!spirv.struct<(f32, !spirv.array<4xf32>)>, Input> 33 spirv.func @foo() -> () "None" { 34 // expected-error @+1 {{expected spirv.GlobalVariable symbol}} 35 %0 = spirv.mlir.addressof @var2 : !spirv.ptr<!spirv.struct<(f32, !spirv.array<4xf32>)>, Input> 36 } 37} 38 39// ----- 40 41spirv.module Logical GLSL450 { 42 spirv.GlobalVariable @var1 : !spirv.ptr<!spirv.struct<(f32, !spirv.array<4xf32>)>, Input> 43 spirv.func @foo() -> () "None" { 44 // expected-error @+1 {{result type mismatch with the referenced global variable's type}} 45 %0 = spirv.mlir.addressof @var1 : !spirv.ptr<f32, Input> 46 } 47} 48 49// ----- 50 51//===----------------------------------------------------------------------===// 52// spirv.Constant 53//===----------------------------------------------------------------------===// 54 55func.func @const() -> () { 56 // CHECK: spirv.Constant true 57 // CHECK: spirv.Constant 42 : i32 58 // CHECK: spirv.Constant 5.000000e-01 : f32 59 // CHECK: spirv.Constant dense<[2, 3]> : vector<2xi32> 60 // CHECK: spirv.Constant [dense<3.000000e+00> : vector<2xf32>] : !spirv.array<1 x vector<2xf32>> 61 // CHECK: spirv.Constant dense<1> : tensor<2x3xi32> : !spirv.array<2 x !spirv.array<3 x i32>> 62 // CHECK: spirv.Constant dense<1.000000e+00> : tensor<2x3xf32> : !spirv.array<2 x !spirv.array<3 x f32>> 63 // CHECK: spirv.Constant dense<{{\[}}[1, 2, 3], [4, 5, 6]]> : tensor<2x3xi32> : !spirv.array<2 x !spirv.array<3 x i32>> 64 // CHECK: spirv.Constant dense<{{\[}}[1.000000e+00, 2.000000e+00, 3.000000e+00], [4.000000e+00, 5.000000e+00, 6.000000e+00]]> : tensor<2x3xf32> : !spirv.array<2 x !spirv.array<3 x f32>> 65 66 %0 = spirv.Constant true 67 %1 = spirv.Constant 42 : i32 68 %2 = spirv.Constant 0.5 : f32 69 %3 = spirv.Constant dense<[2, 3]> : vector<2xi32> 70 %4 = spirv.Constant [dense<3.0> : vector<2xf32>] : !spirv.array<1xvector<2xf32>> 71 %5 = spirv.Constant dense<1> : tensor<2x3xi32> : !spirv.array<2 x !spirv.array<3 x i32>> 72 %6 = spirv.Constant dense<1.0> : tensor<2x3xf32> : !spirv.array<2 x !spirv.array<3 x f32>> 73 %7 = spirv.Constant dense<[[1, 2, 3], [4, 5, 6]]> : tensor<2x3xi32> : !spirv.array<2 x !spirv.array<3 x i32>> 74 %8 = spirv.Constant dense<[[1.0, 2.0, 3.0], [4.0, 5.0, 6.0]]> : tensor<2x3xf32> : !spirv.array<2 x !spirv.array<3 x f32>> 75 %9 = spirv.Constant [[dense<3.0> : vector<2xf32>]] : !spirv.array<1 x !spirv.array<1xvector<2xf32>>> 76 return 77} 78 79// ----- 80 81func.func @unaccepted_std_attr() -> () { 82 // expected-error @+1 {{cannot have attribute: unit}} 83 %0 = spirv.Constant unit : none 84 return 85} 86 87// ----- 88 89func.func @array_constant() -> () { 90 // expected-error @+1 {{result or element type ('vector<2xf32>') does not match value type ('vector<2xi32>')}} 91 %0 = spirv.Constant [dense<3.0> : vector<2xf32>, dense<4> : vector<2xi32>] : !spirv.array<2xvector<2xf32>> 92 return 93} 94 95// ----- 96 97func.func @array_constant() -> () { 98 // expected-error @+1 {{must have spirv.array result type for array value}} 99 %0 = spirv.Constant [dense<3.0> : vector<2xf32>] : !spirv.rtarray<vector<2xf32>> 100 return 101} 102 103// ----- 104 105func.func @non_nested_array_constant() -> () { 106 // expected-error @+1 {{only support nested array result type}} 107 %0 = spirv.Constant dense<3.0> : tensor<2x2xf32> : !spirv.array<2xvector<2xf32>> 108 return 109} 110 111// ----- 112 113func.func @value_result_type_mismatch() -> () { 114 // expected-error @+1 {{result or element type ('vector<4xi32>') does not match value type ('tensor<4xi32>')}} 115 %0 = "spirv.Constant"() {value = dense<0> : tensor<4xi32>} : () -> (vector<4xi32>) 116} 117 118// ----- 119 120func.func @value_result_type_mismatch() -> () { 121 // expected-error @+1 {{result element type ('i32') does not match value element type ('f32')}} 122 %0 = spirv.Constant dense<1.0> : tensor<2x3xf32> : !spirv.array<2 x !spirv.array<3 x i32>> 123} 124 125// ----- 126 127func.func @value_result_num_elements_mismatch() -> () { 128 // expected-error @+1 {{result number of elements (6) does not match value number of elements (4)}} 129 %0 = spirv.Constant dense<1.0> : tensor<2x2xf32> : !spirv.array<2 x !spirv.array<3 x f32>> 130 return 131} 132 133// ----- 134 135//===----------------------------------------------------------------------===// 136// spirv.EntryPoint 137//===----------------------------------------------------------------------===// 138 139spirv.module Logical GLSL450 { 140 spirv.func @do_nothing() -> () "None" { 141 spirv.Return 142 } 143 // CHECK: spirv.EntryPoint "GLCompute" @do_nothing 144 spirv.EntryPoint "GLCompute" @do_nothing 145} 146 147spirv.module Logical GLSL450 { 148 spirv.GlobalVariable @var2 : !spirv.ptr<f32, Input> 149 spirv.GlobalVariable @var3 : !spirv.ptr<f32, Output> 150 spirv.func @do_something(%arg0 : !spirv.ptr<f32, Input>, %arg1 : !spirv.ptr<f32, Output>) -> () "None" { 151 %1 = spirv.Load "Input" %arg0 : f32 152 spirv.Store "Output" %arg1, %1 : f32 153 spirv.Return 154 } 155 // CHECK: spirv.EntryPoint "GLCompute" @do_something, @var2, @var3 156 spirv.EntryPoint "GLCompute" @do_something, @var2, @var3 157} 158 159// ----- 160 161spirv.module Logical GLSL450 { 162 spirv.func @do_nothing() -> () "None" { 163 spirv.Return 164 } 165 // expected-error @+1 {{invalid kind of attribute specified}} 166 spirv.EntryPoint "GLCompute" "do_nothing" 167} 168 169// ----- 170 171spirv.module Logical GLSL450 { 172 spirv.func @do_nothing() -> () "None" { 173 spirv.Return 174 } 175 // expected-error @+1 {{function 'do_something' not found in 'spirv.module'}} 176 spirv.EntryPoint "GLCompute" @do_something 177} 178 179/// TODO: Add a test that verifies an error is thrown 180/// when interface entries of EntryPointOp are not 181/// spirv.Variables. There is currently no other op that has a spirv.ptr 182/// return type 183 184// ----- 185 186spirv.module Logical GLSL450 { 187 spirv.func @do_nothing() -> () "None" { 188 // expected-error @+1 {{op must appear in a module-like op's block}} 189 spirv.EntryPoint "GLCompute" @do_something 190 } 191} 192 193// ----- 194 195spirv.module Logical GLSL450 { 196 spirv.func @do_nothing() -> () "None" { 197 spirv.Return 198 } 199 spirv.EntryPoint "GLCompute" @do_nothing 200 // expected-error @+1 {{duplicate of a previous EntryPointOp}} 201 spirv.EntryPoint "GLCompute" @do_nothing 202} 203 204// ----- 205 206spirv.module Logical GLSL450 { 207 spirv.func @do_nothing() -> () "None" { 208 spirv.Return 209 } 210 spirv.EntryPoint "GLCompute" @do_nothing 211 // expected-error @+1 {{'spirv.EntryPoint' invalid execution_model attribute specification: "ContractionOff"}} 212 spirv.EntryPoint "ContractionOff" @do_nothing 213} 214 215// ----- 216 217//===----------------------------------------------------------------------===// 218// spirv.ExecutionMode 219//===----------------------------------------------------------------------===// 220 221spirv.module Logical GLSL450 { 222 spirv.func @do_nothing() -> () "None" { 223 spirv.Return 224 } 225 spirv.EntryPoint "GLCompute" @do_nothing 226 // CHECK: spirv.ExecutionMode {{@.*}} "ContractionOff" 227 spirv.ExecutionMode @do_nothing "ContractionOff" 228} 229 230spirv.module Logical GLSL450 { 231 spirv.func @do_nothing() -> () "None" { 232 spirv.Return 233 } 234 spirv.EntryPoint "GLCompute" @do_nothing 235 // CHECK: spirv.ExecutionMode {{@.*}} "LocalSizeHint", 3, 4, 5 236 spirv.ExecutionMode @do_nothing "LocalSizeHint", 3, 4, 5 237} 238 239// ----- 240 241spirv.module Logical GLSL450 { 242 spirv.func @do_nothing() -> () "None" { 243 spirv.Return 244 } 245 spirv.EntryPoint "GLCompute" @do_nothing 246 // expected-error @+1 {{custom op 'spirv.ExecutionMode' invalid execution_mode attribute specification: "GLCompute"}} 247 spirv.ExecutionMode @do_nothing "GLCompute", 3, 4, 5 248} 249 250// ----- 251 252//===----------------------------------------------------------------------===// 253// spirv.func 254//===----------------------------------------------------------------------===// 255 256// CHECK: spirv.func @foo() "None" 257spirv.func @foo() "None" 258 259// CHECK: spirv.func @bar(%{{.+}}: i32) -> i32 "Inline|Pure" { 260spirv.func @bar(%arg: i32) -> (i32) "Inline|Pure" { 261 // CHECK-NEXT: spirv. 262 spirv.ReturnValue %arg: i32 263// CHECK-NEXT: } 264} 265 266// CHECK: spirv.func @baz(%{{.+}}: i32) "DontInline" attributes {additional_stuff = 64 : i64} 267spirv.func @baz(%arg: i32) "DontInline" attributes { 268 additional_stuff = 64 269} { spirv.Return } 270 271// ----- 272 273spirv.module Logical GLSL450 requires #spirv.vce<v1.0, [Shader, Linkage], []> { 274 // CHECK: linkage_attributes = #spirv.linkage_attributes<linkage_name = "outside.func", linkage_type = <Import>> 275 spirv.func @outside.func.with.linkage(%arg0 : i8) -> () "Pure" attributes { 276 linkage_attributes=#spirv.linkage_attributes< 277 linkage_name="outside.func", 278 linkage_type=<Import> 279 > 280 } 281 spirv.func @inside.func() -> () "Pure" attributes {} {spirv.Return} 282} 283// ----- 284 285spirv.module Logical GLSL450 requires #spirv.vce<v1.0, [Shader, Linkage], []> { 286 // expected-error @+1 {{'spirv.module' cannot contain external functions without 'Import' linkage_attributes (LinkageAttributes)}} 287 spirv.func @outside.func.without.linkage(%arg0 : i8) -> () "Pure" 288 spirv.func @inside.func() -> () "Pure" attributes {} {spirv.Return} 289} 290 291// ----- 292 293// expected-error @+1 {{expected function_control attribute specified as string}} 294spirv.func @missing_function_control() { spirv.Return } 295 296// ----- 297 298// expected-error @+1 {{cannot have more than one result}} 299spirv.func @cannot_have_more_than_one_result(%arg: i32) -> (i32, i32) "None" 300 301// ----- 302 303// expected-error @+1 {{expected SSA identifier}} 304spirv.func @cannot_have_variadic_arguments(%arg: i32, ...) "None" 305 306// ----- 307 308// Nested function 309spirv.module Logical GLSL450 { 310 spirv.func @outer_func() -> () "None" { 311 // expected-error @+1 {{must appear in a module-like op's block}} 312 spirv.func @inner_func() -> () "None" { 313 spirv.Return 314 } 315 spirv.Return 316 } 317} 318 319// ----- 320 321//===----------------------------------------------------------------------===// 322// spirv.GlobalVariable 323//===----------------------------------------------------------------------===// 324 325spirv.module Logical GLSL450 { 326 // CHECK: spirv.GlobalVariable @var0 : !spirv.ptr<f32, Input> 327 spirv.GlobalVariable @var0 : !spirv.ptr<f32, Input> 328} 329 330// TODO: Fix test case after initialization with normal constant is addressed 331// spirv.module Logical GLSL450 { 332// %0 = spirv.Constant 4.0 : f32 333// COM: CHECK: spirv.Variable init(%0) : !spirv.ptr<f32, Private> 334// spirv.GlobalVariable @var1 init(%0) : !spirv.ptr<f32, Private> 335// } 336 337// ----- 338 339spirv.module Logical GLSL450 { 340 spirv.SpecConstant @sc = 4.0 : f32 341 // CHECK: spirv.GlobalVariable @var initializer(@sc) : !spirv.ptr<f32, Private> 342 spirv.GlobalVariable @var initializer(@sc) : !spirv.ptr<f32, Private> 343} 344 345// ----- 346 347// Allow initializers coming from other module-like ops 348spirv.SpecConstant @sc = 4.0 : f32 349// CHECK: spirv.GlobalVariable @var initializer(@sc) 350spirv.GlobalVariable @var initializer(@sc) : !spirv.ptr<f32, Private> 351 352 353// ----- 354// Allow SpecConstantComposite as initializer 355 spirv.module Logical GLSL450 { 356 spirv.SpecConstant @sc1 = 1 : i8 357 spirv.SpecConstant @sc2 = 2 : i8 358 spirv.SpecConstant @sc3 = 3 : i8 359 spirv.SpecConstantComposite @scc (@sc1, @sc2, @sc3) : !spirv.array<3 x i8> 360 361 // CHECK: spirv.GlobalVariable @var initializer(@scc) : !spirv.ptr<!spirv.array<3 x i8>, Private> 362 spirv.GlobalVariable @var initializer(@scc) : !spirv.ptr<!spirv.array<3 x i8>, Private> 363} 364 365// ----- 366 367spirv.module Logical GLSL450 { 368 // CHECK: spirv.GlobalVariable @var0 bind(1, 2) : !spirv.ptr<f32, Uniform> 369 spirv.GlobalVariable @var0 bind(1, 2) : !spirv.ptr<f32, Uniform> 370} 371 372// TODO: Fix test case after initialization with constant is addressed 373// spirv.module Logical GLSL450 { 374// %0 = spirv.Constant 4.0 : f32 375// COM: CHECK: spirv.GlobalVariable @var1 initializer(%0) {binding = 5 : i32} : !spirv.ptr<f32, Private> 376// spirv.GlobalVariable @var1 initializer(%0) {binding = 5 : i32} : !spirv.ptr<f32, Private> 377// } 378 379// ----- 380 381spirv.module Logical GLSL450 { 382 // CHECK: spirv.GlobalVariable @var1 built_in("GlobalInvocationID") : !spirv.ptr<vector<3xi32>, Input> 383 spirv.GlobalVariable @var1 built_in("GlobalInvocationID") : !spirv.ptr<vector<3xi32>, Input> 384 // CHECK: spirv.GlobalVariable @var2 built_in("GlobalInvocationID") : !spirv.ptr<vector<3xi32>, Input> 385 spirv.GlobalVariable @var2 {built_in = "GlobalInvocationID"} : !spirv.ptr<vector<3xi32>, Input> 386} 387 388// ----- 389 390// Allow in other module-like ops 391module { 392 // CHECK: spirv.GlobalVariable 393 spirv.GlobalVariable @var0 : !spirv.ptr<f32, Input> 394} 395 396// ----- 397 398spirv.module Logical GLSL450 requires #spirv.vce<v1.0, [Shader, Linkage], []> { 399 // CHECK: linkage_attributes = #spirv.linkage_attributes<linkage_name = "outSideGlobalVar1", linkage_type = <Import>> 400 spirv.GlobalVariable @var1 { 401 linkage_attributes=#spirv.linkage_attributes< 402 linkage_name="outSideGlobalVar1", 403 linkage_type=<Import> 404 > 405 } : !spirv.ptr<f32, Private> 406} 407 408 409// ----- 410 411spirv.module Logical GLSL450 { 412 // expected-error @+1 {{expected spirv.ptr type}} 413 spirv.GlobalVariable @var0 : f32 414} 415 416// ----- 417 418spirv.module Logical GLSL450 { 419 // expected-error @+1 {{result must be of a !spv.ptr type}} 420 "spirv.GlobalVariable"() {sym_name = "var0", type = none} : () -> () 421} 422 423// ----- 424 425spirv.module Logical GLSL450 { 426 // expected-error @+1 {{op initializer must be result of a spirv.SpecConstant or spirv.GlobalVariable or spirv.SpecConstantCompositeOp op}} 427 spirv.GlobalVariable @var0 initializer(@var1) : !spirv.ptr<f32, Private> 428} 429 430// ----- 431 432spirv.module Logical GLSL450 { 433 // expected-error @+1 {{storage class cannot be 'Generic'}} 434 spirv.GlobalVariable @var0 : !spirv.ptr<f32, Generic> 435} 436 437// ----- 438 439spirv.module Logical GLSL450 { 440 // expected-error @+1 {{storage class cannot be 'Function'}} 441 spirv.GlobalVariable @var0 : !spirv.ptr<f32, Function> 442} 443 444// ----- 445 446spirv.module Logical GLSL450 { 447 spirv.func @foo() "None" { 448 // expected-error @+1 {{op must appear in a module-like op's block}} 449 spirv.GlobalVariable @var0 : !spirv.ptr<f32, Input> 450 spirv.Return 451 } 452} 453 454// ----- 455 456//===----------------------------------------------------------------------===// 457// spirv.module 458//===----------------------------------------------------------------------===// 459 460// Module without capability and extension 461// CHECK: spirv.module Logical GLSL450 462spirv.module Logical GLSL450 { } 463 464// Module with a name 465// CHECK: spirv.module @{{.*}} Logical GLSL450 466spirv.module @name Logical GLSL450 { } 467 468// Module with (version, capabilities, extensions) triple 469// CHECK: spirv.module Logical GLSL450 requires #spirv.vce<v1.0, [Shader], [SPV_KHR_16bit_storage]> 470spirv.module Logical GLSL450 requires #spirv.vce<v1.0, [Shader], [SPV_KHR_16bit_storage]> { } 471 472// Module with additional attributes 473// CHECK: spirv.module Logical GLSL450 attributes {foo = "bar"} 474spirv.module Logical GLSL450 attributes {foo = "bar"} { } 475 476// Module with VCE triple and additional attributes 477// CHECK: spirv.module Logical GLSL450 requires #spirv.vce<v1.0, [Shader], [SPV_KHR_16bit_storage]> attributes {foo = "bar"} 478spirv.module Logical GLSL450 479 requires #spirv.vce<v1.0, [Shader], [SPV_KHR_16bit_storage]> 480 attributes {foo = "bar"} { } 481 482// Module with function 483// CHECK: spirv.module 484spirv.module Logical GLSL450 { 485 spirv.func @do_nothing() -> () "None" { 486 spirv.Return 487 } 488} 489 490// ----- 491 492// Missing addressing model 493// expected-error@+1 {{'spirv.module' expected valid keyword}} 494spirv.module { } 495 496// ----- 497 498// Wrong addressing model 499// expected-error@+1 {{'spirv.module' invalid addressing_model attribute specification: Physical}} 500spirv.module Physical { } 501 502// ----- 503 504// Missing memory model 505// expected-error@+1 {{'spirv.module' expected valid keyword}} 506spirv.module Logical { } 507 508// ----- 509 510// Wrong memory model 511// expected-error@+1 {{'spirv.module' invalid memory_model attribute specification: Bla}} 512spirv.module Logical Bla { } 513 514// ----- 515 516// Module with multiple blocks 517// expected-error @+1 {{expects region #0 to have 0 or 1 blocks}} 518spirv.module Logical GLSL450 { 519^first: 520 spirv.Return 521^second: 522 spirv.Return 523} 524 525// ----- 526 527// Use non SPIR-V op inside module 528spirv.module Logical GLSL450 { 529 // expected-error @+1 {{'spirv.module' can only contain spirv.* ops}} 530 "dialect.op"() : () -> () 531} 532 533// ----- 534 535// Use non SPIR-V op inside function 536spirv.module Logical GLSL450 { 537 spirv.func @do_nothing() -> () "None" { 538 // expected-error @+1 {{functions in 'spirv.module' can only contain spirv.* ops}} 539 "dialect.op"() : () -> () 540 } 541} 542 543// ----- 544 545// Use external function 546spirv.module Logical GLSL450 { 547 // expected-error @+1 {{'spirv.module' cannot contain external functions}} 548 spirv.func @extern() -> () "None" 549} 550 551// ----- 552 553//===----------------------------------------------------------------------===// 554// spirv.mlir.referenceof 555//===----------------------------------------------------------------------===// 556 557spirv.module Logical GLSL450 { 558 spirv.SpecConstant @sc1 = false 559 spirv.SpecConstant @sc2 = 42 : i64 560 spirv.SpecConstant @sc3 = 1.5 : f32 561 562 spirv.SpecConstantComposite @scc (@sc1, @sc2, @sc3) : !spirv.struct<(i1, i64, f32)> 563 564 // CHECK-LABEL: @reference 565 spirv.func @reference() -> i1 "None" { 566 // CHECK: spirv.mlir.referenceof @sc1 : i1 567 %0 = spirv.mlir.referenceof @sc1 : i1 568 spirv.ReturnValue %0 : i1 569 } 570 571 // CHECK-LABEL: @reference_composite 572 spirv.func @reference_composite() -> i1 "None" { 573 // CHECK: spirv.mlir.referenceof @scc : !spirv.struct<(i1, i64, f32)> 574 %0 = spirv.mlir.referenceof @scc : !spirv.struct<(i1, i64, f32)> 575 %1 = spirv.CompositeExtract %0[0 : i32] : !spirv.struct<(i1, i64, f32)> 576 spirv.ReturnValue %1 : i1 577 } 578 579 // CHECK-LABEL: @initialize 580 spirv.func @initialize() -> i64 "None" { 581 // CHECK: spirv.mlir.referenceof @sc2 : i64 582 %0 = spirv.mlir.referenceof @sc2 : i64 583 %1 = spirv.Variable init(%0) : !spirv.ptr<i64, Function> 584 %2 = spirv.Load "Function" %1 : i64 585 spirv.ReturnValue %2 : i64 586 } 587 588 // CHECK-LABEL: @compute 589 spirv.func @compute() -> f32 "None" { 590 // CHECK: spirv.mlir.referenceof @sc3 : f32 591 %0 = spirv.mlir.referenceof @sc3 : f32 592 %1 = spirv.Constant 6.0 : f32 593 %2 = spirv.FAdd %0, %1 : f32 594 spirv.ReturnValue %2 : f32 595 } 596} 597 598// ----- 599 600// Allow taking reference of spec constant in other module-like ops 601spirv.SpecConstant @sc = 5 : i32 602func.func @reference_of() { 603 // CHECK: spirv.mlir.referenceof @sc 604 %0 = spirv.mlir.referenceof @sc : i32 605 return 606} 607 608// ----- 609 610spirv.SpecConstant @sc = 5 : i32 611spirv.SpecConstantComposite @scc (@sc) : !spirv.array<1 x i32> 612 613func.func @reference_of_composite() { 614 // CHECK: spirv.mlir.referenceof @scc : !spirv.array<1 x i32> 615 %0 = spirv.mlir.referenceof @scc : !spirv.array<1 x i32> 616 %1 = spirv.CompositeExtract %0[0 : i32] : !spirv.array<1 x i32> 617 return 618} 619 620// ----- 621 622spirv.module Logical GLSL450 { 623 spirv.func @foo() -> () "None" { 624 // expected-error @+1 {{expected spirv.SpecConstant or spirv.SpecConstantComposite symbol}} 625 %0 = spirv.mlir.referenceof @sc : i32 626 spirv.Return 627 } 628} 629 630// ----- 631 632spirv.module Logical GLSL450 { 633 spirv.SpecConstant @sc = 42 : i32 634 spirv.func @foo() -> () "None" { 635 // expected-error @+1 {{result type mismatch with the referenced specialization constant's type}} 636 %0 = spirv.mlir.referenceof @sc : f32 637 spirv.Return 638 } 639} 640 641// ----- 642 643spirv.module Logical GLSL450 { 644 spirv.SpecConstant @sc = 42 : i32 645 spirv.SpecConstantComposite @scc (@sc) : !spirv.array<1 x i32> 646 spirv.func @foo() -> () "None" { 647 // expected-error @+1 {{result type mismatch with the referenced specialization constant's type}} 648 %0 = spirv.mlir.referenceof @scc : f32 649 spirv.Return 650 } 651} 652 653// ----- 654 655//===----------------------------------------------------------------------===// 656// spirv.SpecConstant 657//===----------------------------------------------------------------------===// 658 659spirv.module Logical GLSL450 { 660 // CHECK: spirv.SpecConstant @sc1 = false 661 spirv.SpecConstant @sc1 = false 662 // CHECK: spirv.SpecConstant @sc2 spec_id(5) = 42 : i64 663 spirv.SpecConstant @sc2 spec_id(5) = 42 : i64 664 // CHECK: spirv.SpecConstant @sc3 = 1.500000e+00 : f32 665 spirv.SpecConstant @sc3 = 1.5 : f32 666} 667 668// ----- 669 670spirv.module Logical GLSL450 { 671 // expected-error @+1 {{SpecId cannot be negative}} 672 spirv.SpecConstant @sc2 spec_id(-5) = 42 : i64 673} 674 675// ----- 676 677spirv.module Logical GLSL450 { 678 // expected-error @+1 {{default value bitwidth disallowed}} 679 spirv.SpecConstant @sc = 15 : i4 680} 681 682// ----- 683 684spirv.module Logical GLSL450 { 685 // expected-error @+1 {{default value can only be a bool, integer, or float scalar}} 686 spirv.SpecConstant @sc = dense<[2, 3]> : vector<2xi32> 687} 688 689// ----- 690 691func.func @use_in_function() -> () { 692 // expected-error @+1 {{op must appear in a module-like op's block}} 693 spirv.SpecConstant @sc = false 694 return 695} 696 697// ----- 698 699//===----------------------------------------------------------------------===// 700// spirv.SpecConstantComposite 701//===----------------------------------------------------------------------===// 702 703spirv.module Logical GLSL450 { 704 // expected-error @+1 {{result type must be a composite type}} 705 spirv.SpecConstantComposite @scc2 (@sc1, @sc2, @sc3) : i32 706} 707 708//===----------------------------------------------------------------------===// 709// spirv.SpecConstantComposite (spirv.array) 710//===----------------------------------------------------------------------===// 711 712// ----- 713 714spirv.module Logical GLSL450 { 715 spirv.SpecConstant @sc1 = 1.5 : f32 716 spirv.SpecConstant @sc2 = 2.5 : f32 717 spirv.SpecConstant @sc3 = 3.5 : f32 718 // CHECK: spirv.SpecConstantComposite @scc (@sc1, @sc2, @sc3) : !spirv.array<3 x f32> 719 spirv.SpecConstantComposite @scc (@sc1, @sc2, @sc3) : !spirv.array<3 x f32> 720} 721 722// ----- 723 724spirv.module Logical GLSL450 { 725 spirv.SpecConstant @sc1 = false 726 spirv.SpecConstant @sc2 spec_id(5) = 42 : i64 727 spirv.SpecConstant @sc3 = 1.5 : f32 728 // expected-error @+1 {{has incorrect number of operands: expected 4, but provided 3}} 729 spirv.SpecConstantComposite @scc (@sc1, @sc2, @sc3) : !spirv.array<4 x f32> 730 731} 732 733// ----- 734 735spirv.module Logical GLSL450 { 736 spirv.SpecConstant @sc1 = 1 : i32 737 spirv.SpecConstant @sc2 = 2.5 : f32 738 spirv.SpecConstant @sc3 = 3.5 : f32 739 // expected-error @+1 {{has incorrect types of operands: expected 'f32', but provided 'i32'}} 740 spirv.SpecConstantComposite @scc (@sc1, @sc2, @sc3) : !spirv.array<3 x f32> 741} 742 743//===----------------------------------------------------------------------===// 744// spirv.SpecConstantComposite (spirv.struct) 745//===----------------------------------------------------------------------===// 746 747// ----- 748 749spirv.module Logical GLSL450 { 750 spirv.SpecConstant @sc1 = 1 : i32 751 spirv.SpecConstant @sc2 = 2.5 : f32 752 spirv.SpecConstant @sc3 = 3.5 : f32 753 // CHECK: spirv.SpecConstantComposite @scc (@sc1, @sc2, @sc3) : !spirv.struct<(i32, f32, f32)> 754 spirv.SpecConstantComposite @scc (@sc1, @sc2, @sc3) : !spirv.struct<(i32, f32, f32)> 755} 756 757// ----- 758 759spirv.module Logical GLSL450 { 760 spirv.SpecConstant @sc1 = 1 : i32 761 spirv.SpecConstant @sc2 = 2.5 : f32 762 spirv.SpecConstant @sc3 = 3.5 : f32 763 // expected-error @+1 {{has incorrect number of operands: expected 2, but provided 3}} 764 spirv.SpecConstantComposite @scc (@sc1, @sc2, @sc3) : !spirv.struct<(i32, f32)> 765} 766 767// ----- 768 769spirv.module Logical GLSL450 { 770 spirv.SpecConstant @sc1 = 1.5 : f32 771 spirv.SpecConstant @sc2 = 2.5 : f32 772 spirv.SpecConstant @sc3 = 3.5 : f32 773 // expected-error @+1 {{has incorrect types of operands: expected 'i32', but provided 'f32'}} 774 spirv.SpecConstantComposite @scc (@sc1, @sc2, @sc3) : !spirv.struct<(i32, f32, f32)> 775} 776 777//===----------------------------------------------------------------------===// 778// spirv.SpecConstantComposite (vector) 779//===----------------------------------------------------------------------===// 780 781// ----- 782 783spirv.module Logical GLSL450 { 784 spirv.SpecConstant @sc1 = 1.5 : f32 785 spirv.SpecConstant @sc2 = 2.5 : f32 786 spirv.SpecConstant @sc3 = 3.5 : f32 787 // CHECK: spirv.SpecConstantComposite @scc (@sc1, @sc2, @sc3) : vector<3xf32> 788 spirv.SpecConstantComposite @scc (@sc1, @sc2, @sc3) : vector<3 x f32> 789} 790 791// ----- 792 793spirv.module Logical GLSL450 { 794 spirv.SpecConstant @sc1 = false 795 spirv.SpecConstant @sc2 spec_id(5) = 42 : i64 796 spirv.SpecConstant @sc3 = 1.5 : f32 797 // expected-error @+1 {{has incorrect number of operands: expected 4, but provided 3}} 798 spirv.SpecConstantComposite @scc (@sc1, @sc2, @sc3) : vector<4xf32> 799 800} 801 802// ----- 803 804spirv.module Logical GLSL450 { 805 spirv.SpecConstant @sc1 = 1 : i32 806 spirv.SpecConstant @sc2 = 2.5 : f32 807 spirv.SpecConstant @sc3 = 3.5 : f32 808 // expected-error @+1 {{has incorrect types of operands: expected 'f32', but provided 'i32'}} 809 spirv.SpecConstantComposite @scc (@sc1, @sc2, @sc3) : vector<3xf32> 810} 811 812//===----------------------------------------------------------------------===// 813// spirv.SpecConstantComposite (spirv.KHR.coopmatrix) 814//===----------------------------------------------------------------------===// 815 816// ----- 817 818spirv.module Logical GLSL450 { 819 spirv.SpecConstant @sc1 = 1.5 : f32 820 // expected-error @+1 {{unsupported composite type}} 821 spirv.SpecConstantComposite @scc (@sc1) : !spirv.coopmatrix<8x16xf32, Device, MatrixA> 822} 823 824//===----------------------------------------------------------------------===// 825// spirv.SpecConstantOperation 826//===----------------------------------------------------------------------===// 827 828// ----- 829 830spirv.module Logical GLSL450 { 831 spirv.func @foo() -> i32 "None" { 832 // CHECK: [[LHS:%.*]] = spirv.Constant 833 %0 = spirv.Constant 1: i32 834 // CHECK: [[RHS:%.*]] = spirv.Constant 835 %1 = spirv.Constant 1: i32 836 837 // CHECK: spirv.SpecConstantOperation wraps "spirv.IAdd"([[LHS]], [[RHS]]) : (i32, i32) -> i32 838 %2 = spirv.SpecConstantOperation wraps "spirv.IAdd"(%0, %1) : (i32, i32) -> i32 839 840 spirv.ReturnValue %2 : i32 841 } 842} 843 844// ----- 845 846spirv.module Logical GLSL450 { 847 spirv.SpecConstant @sc = 42 : i32 848 849 spirv.func @foo() -> i32 "None" { 850 // CHECK: [[SC:%.*]] = spirv.mlir.referenceof @sc 851 %0 = spirv.mlir.referenceof @sc : i32 852 // CHECK: spirv.SpecConstantOperation wraps "spirv.ISub"([[SC]], [[SC]]) : (i32, i32) -> i32 853 %1 = spirv.SpecConstantOperation wraps "spirv.ISub"(%0, %0) : (i32, i32) -> i32 854 spirv.ReturnValue %1 : i32 855 } 856} 857 858// ----- 859 860spirv.module Logical GLSL450 { 861 spirv.func @foo() -> i32 "None" { 862 %0 = spirv.Constant 1: i32 863 // expected-error @+1 {{op expects parent op 'spirv.SpecConstantOperation'}} 864 spirv.mlir.yield %0 : i32 865 } 866} 867 868// ----- 869 870spirv.module Logical GLSL450 { 871 spirv.func @foo() -> () "None" { 872 %0 = spirv.Variable : !spirv.ptr<i32, Function> 873 874 // expected-error @+1 {{invalid enclosed op}} 875 %1 = spirv.SpecConstantOperation wraps "spirv.Load"(%0) {memory_access = #spirv.memory_access<None>} : (!spirv.ptr<i32, Function>) -> i32 876 spirv.Return 877 } 878} 879 880// ----- 881 882spirv.module Logical GLSL450 { 883 spirv.func @foo() -> () "None" { 884 %0 = spirv.Variable : !spirv.ptr<i32, Function> 885 %1 = spirv.Load "Function" %0 : i32 886 887 // expected-error @+1 {{invalid operand, must be defined by a constant operation}} 888 %2 = spirv.SpecConstantOperation wraps "spirv.IAdd"(%1, %1) : (i32, i32) -> i32 889 890 spirv.Return 891 } 892} 893