mlir/test/Dialect/SPIRV/IR/structure-ops.mlir

   1 // RUN: mlir-opt -allow-unregistered-dialect -split-input-file -verify-diagnostics %s | FileCheck %s
   2
   3 //===----------------------------------------------------------------------===//
   4 // spirv.mlir.addressof
   5 //===----------------------------------------------------------------------===//
   6
   7 spirv.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
  22 spirv.GlobalVariable @var : !spirv.ptr<!spirv.struct<(f32, !spirv.array<4xf32>)>, Input>
  23 func.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
  31 spirv.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
  41 spirv.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
  55 func.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
  81 func.func @unaccepted_std_attr() -> () {
  82   // expected-error @+1 {{cannot have attribute: unit}}
  83   %0 = spirv.Constant unit : none
  84   return
  85 }
  86
  87 // -----
  88
  89 func.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
  97 func.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
 105 func.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
 113 func.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
 120 func.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
 127 func.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
 139 spirv.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
 147 spirv.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
 161 spirv.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
 171 spirv.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
 186 spirv.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
 195 spirv.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
 206 spirv.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
 221 spirv.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
 230 spirv.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
 241 spirv.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"
 257 spirv.func @foo() "None"
 258
 259 // CHECK: spirv.func @bar(%{{.+}}: i32) -> i32 "Inline|Pure" {
 260 spirv.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}
 267 spirv.func @baz(%arg: i32) "DontInline" attributes {
 268   additional_stuff = 64
 269 } { spirv.Return }
 270
 271 // -----
 272
 273 spirv.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
 285 spirv.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}}
 294 spirv.func @missing_function_control() { spirv.Return }
 295
 296 // -----
 297
 298 // expected-error @+1 {{cannot have more than one result}}
 299 spirv.func @cannot_have_more_than_one_result(%arg: i32) -> (i32, i32) "None"
 300
 301 // -----
 302
 303 // expected-error @+1 {{expected SSA identifier}}
 304 spirv.func @cannot_have_variadic_arguments(%arg: i32, ...) "None"
 305
 306 // -----
 307
 308 // Nested function
 309 spirv.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
 325 spirv.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
 339 spirv.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
 348 spirv.SpecConstant @sc = 4.0 : f32
 349 // CHECK: spirv.GlobalVariable @var initializer(@sc)
 350 spirv.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
 367 spirv.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
 381 spirv.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
 391 module {
 392   // CHECK: spirv.GlobalVariable
 393   spirv.GlobalVariable @var0 : !spirv.ptr<f32, Input>
 394 }
 395
 396 // -----
 397
 398 spirv.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
 411 spirv.module Logical GLSL450 {
 412   // expected-error @+1 {{expected spirv.ptr type}}
 413   spirv.GlobalVariable @var0 : f32
 414 }
 415
 416 // -----
 417
 418 spirv.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
 425 spirv.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
 432 spirv.module Logical GLSL450 {
 433   // expected-error @+1 {{storage class cannot be 'Generic'}}
 434   spirv.GlobalVariable @var0 : !spirv.ptr<f32, Generic>
 435 }
 436
 437 // -----
 438
 439 spirv.module Logical GLSL450 {
 440   // expected-error @+1 {{storage class cannot be 'Function'}}
 441   spirv.GlobalVariable @var0 : !spirv.ptr<f32, Function>
 442 }
 443
 444 // -----
 445
 446 spirv.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
 462 spirv.module Logical GLSL450 { }
 463
 464 // Module with a name
 465 // CHECK: spirv.module @{{.*}} Logical GLSL450
 466 spirv.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]>
 470 spirv.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"}
 474 spirv.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"}
 478 spirv.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
 484 spirv.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}}
 494 spirv.module { }
 495
 496 // -----
 497
 498 // Wrong addressing model
 499 // expected-error@+1 {{'spirv.module' invalid addressing_model attribute specification: Physical}}
 500 spirv.module Physical { }
 501
 502 // -----
 503
 504 // Missing memory model
 505 // expected-error@+1 {{'spirv.module' expected valid keyword}}
 506 spirv.module Logical { }
 507
 508 // -----
 509
 510 // Wrong memory model
 511 // expected-error@+1 {{'spirv.module' invalid memory_model attribute specification: Bla}}
 512 spirv.module Logical Bla { }
 513
 514 // -----
 515
 516 // Module with multiple blocks
 517 // expected-error @+1 {{expects region #0 to have 0 or 1 blocks}}
 518 spirv.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
 528 spirv.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
 536 spirv.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
 546 spirv.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
 557 spirv.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
 601 spirv.SpecConstant @sc = 5 : i32
 602 func.func @reference_of() {
 603   // CHECK: spirv.mlir.referenceof @sc
 604   %0 = spirv.mlir.referenceof @sc : i32
 605   return
 606 }
 607
 608 // -----
 609
 610 spirv.SpecConstant @sc = 5 : i32
 611 spirv.SpecConstantComposite @scc (@sc) : !spirv.array<1 x i32>
 612
 613 func.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
 622 spirv.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
 632 spirv.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
 643 spirv.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
 659 spirv.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
 670 spirv.module Logical GLSL450 {
 671   // expected-error @+1 {{SpecId cannot be negative}}
 672   spirv.SpecConstant @sc2 spec_id(-5) = 42 : i64
 673 }
 674
 675 // -----
 676
 677 spirv.module Logical GLSL450 {
 678   // expected-error @+1 {{default value bitwidth disallowed}}
 679   spirv.SpecConstant @sc = 15 : i4
 680 }
 681
 682 // -----
 683
 684 spirv.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
 691 func.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
 703 spirv.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
 714 spirv.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
 724 spirv.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
 735 spirv.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
 749 spirv.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
 759 spirv.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
 769 spirv.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
 783 spirv.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
 793 spirv.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
 804 spirv.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
 818 spirv.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
 830 spirv.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
 846 spirv.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
 860 spirv.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
 870 spirv.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
 882 spirv.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 }