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