1// RUN: mlir-opt -split-input-file -verify-diagnostics %s | FileCheck %s
2
3// expected-error @+1 {{found unsupported 'spv.something' attribute on operation}}
4func @unknown_attr_on_op() attributes {
5  spv.something = 64
6} { return }
7
8// -----
9
10// expected-error @+1 {{found unsupported 'spv.something' attribute on region argument}}
11func @unknown_attr_on_region(%arg: i32 {spv.something}) {
12  return
13}
14
15// -----
16
17// expected-error @+1 {{cannot attach SPIR-V attributes to region result}}
18func @unknown_attr_on_region() -> (i32 {spv.something}) {
19  %0 = constant 10.0 : f32
20  return %0: f32
21}
22
23// -----
24
25//===----------------------------------------------------------------------===//
26// spv.entry_point_abi
27//===----------------------------------------------------------------------===//
28
29// expected-error @+1 {{'spv.entry_point_abi' attribute must be a dictionary attribute containing one 32-bit integer elements attribute: 'local_size'}}
30func @spv_entry_point() attributes {
31  spv.entry_point_abi = 64
32} { return }
33
34// -----
35
36// expected-error @+1 {{'spv.entry_point_abi' attribute must be a dictionary attribute containing one 32-bit integer elements attribute: 'local_size'}}
37func @spv_entry_point() attributes {
38  spv.entry_point_abi = {local_size = 64}
39} { return }
40
41// -----
42
43func @spv_entry_point() attributes {
44  // CHECK: {spv.entry_point_abi = {local_size = dense<[64, 1, 1]> : vector<3xi32>}}
45  spv.entry_point_abi = {local_size = dense<[64, 1, 1]>: vector<3xi32>}
46} { return }
47
48// -----
49
50//===----------------------------------------------------------------------===//
51// spv.interface_var_abi
52//===----------------------------------------------------------------------===//
53
54// expected-error @+1 {{'spv.interface_var_abi' must be a spirv::InterfaceVarABIAttr}}
55func @interface_var(
56  %arg0 : f32 {spv.interface_var_abi = 64}
57) { return }
58
59// -----
60
61func @interface_var(
62// expected-error @+1 {{missing descriptor set}}
63  %arg0 : f32 {spv.interface_var_abi = #spv.interface_var_abi<()>}
64) { return }
65
66// -----
67
68func @interface_var(
69// expected-error @+1 {{missing binding}}
70  %arg0 : f32 {spv.interface_var_abi = #spv.interface_var_abi<(1,)>}
71) { return }
72
73// -----
74
75func @interface_var(
76// expected-error @+1 {{unknown storage class: }}
77  %arg0 : f32 {spv.interface_var_abi = #spv.interface_var_abi<(1,2), Foo>}
78) { return }
79
80// -----
81
82// CHECK: {spv.interface_var_abi = #spv.interface_var_abi<(0, 1), Uniform>}
83func @interface_var(
84    %arg0 : f32 {spv.interface_var_abi = #spv.interface_var_abi<(0, 1), Uniform>}
85) { return }
86
87// -----
88
89// CHECK: {spv.interface_var_abi = #spv.interface_var_abi<(0, 1)>}
90func @interface_var(
91    %arg0 : f32 {spv.interface_var_abi = #spv.interface_var_abi<(0, 1)>}
92) { return }
93
94// -----
95
96// expected-error @+1 {{'spv.interface_var_abi' attribute cannot specify storage class when attaching to a non-scalar value}}
97func @interface_var(
98  %arg0 : memref<4xf32> {spv.interface_var_abi = #spv.interface_var_abi<(0, 1), Uniform>}
99) { return }
100
101// -----
102
103//===----------------------------------------------------------------------===//
104// spv.target_env
105//===----------------------------------------------------------------------===//
106
107func @target_env_wrong_limits() attributes {
108  spv.target_env = #spv.target_env<
109    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>,
110    // expected-error @+1 {{limits must be a dictionary attribute containing two 32-bit integer attributes 'max_compute_workgroup_invocations' and 'max_compute_workgroup_size'}}
111    {max_compute_workgroup_invocations = 128 : i64, max_compute_workgroup_size = dense<[128, 64, 64]> : vector<3xi32>}>
112} { return }
113
114// -----
115
116func @target_env() attributes {
117  // CHECK:      spv.target_env = #spv.target_env<
118  // CHECK-SAME:   #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>,
119  // CHECK-SAME:   {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 64, 64]> : vector<3xi32>}>
120  spv.target_env = #spv.target_env<
121    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>,
122    {
123      max_compute_workgroup_invocations = 128 : i32,
124      max_compute_workgroup_size = dense<[128, 64, 64]> : vector<3xi32>
125    }>
126} { return }
127
128// -----
129
130func @target_env_vendor_id() attributes {
131  // CHECK:      spv.target_env = #spv.target_env<
132  // CHECK-SAME:   #spv.vce<v1.0, [], []>,
133  // CHECK-SAME:   NVIDIA,
134  // CHECK-SAME:   {}>
135  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, NVIDIA, {}>
136} { return }
137
138// -----
139
140func @target_env_vendor_id_device_type() attributes {
141  // CHECK:      spv.target_env = #spv.target_env<
142  // CHECK-SAME:   #spv.vce<v1.0, [], []>,
143  // CHECK-SAME:   AMD:DiscreteGPU,
144  // CHECK-SAME:   {}>
145  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, AMD:DiscreteGPU, {}>
146} { return }
147
148// -----
149
150func @target_env_vendor_id_device_type_device_id() attributes {
151  // CHECK:      spv.target_env = #spv.target_env<
152  // CHECK-SAME:   #spv.vce<v1.0, [], []>,
153  // CHECK-SAME:   Qualcomm:IntegratedGPU:100925441,
154  // CHECK-SAME:   {}>
155  spv.target_env = #spv.target_env<#spv.vce<v1.0, [], []>, Qualcomm:IntegratedGPU:0x6040001, {}>
156} { return }
157
158// -----
159
160func @target_env_extra_fields() attributes {
161  // expected-error @+6 {{expected '>'}}
162  spv.target_env = #spv.target_env<
163    #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>,
164    {
165      max_compute_workgroup_invocations = 128 : i32,
166      max_compute_workgroup_size = dense<[128, 64, 64]> : vector<3xi32>
167    },
168    more_stuff
169  >
170} { return }
171
172// -----
173
174func @target_env_cooperative_matrix() attributes{
175  // CHECK:      spv.target_env = #spv.target_env<
176  // CHECK-SAME:   SPV_NV_cooperative_matrix
177  // CHECK-SAME:   cooperative_matrix_properties_nv = [
178  // CHECK-SAME:     {a_type = i8, b_type = i8, c_type = i32,
179  // CHECK-SAME:      k_size = 32 : i32, m_size = 8 : i32, n_size = 8 : i32
180  // CHECK-SAME:      result_type = i32, scope = 3 : i32}
181  // CHECK-SAME:     {a_type = f16, b_type = f16, c_type = f16,
182  // CHECK-SAME:      k_size = 16 : i32, m_size = 8 : i32, n_size = 8 : i32
183  // CHECK-SAME:      result_type = f16, scope = 3 : i32}
184  spv.target_env = #spv.target_env<
185  #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class,
186                            SPV_NV_cooperative_matrix]>,
187  {
188    cooperative_matrix_properties_nv = [{
189      m_size = 8: i32,
190      n_size = 8: i32,
191      k_size = 32: i32,
192      a_type = i8,
193      b_type = i8,
194      c_type = i32,
195      result_type = i32,
196      scope = 3: i32
197    }, {
198      m_size = 8: i32,
199      n_size = 8: i32,
200      k_size = 16: i32,
201      a_type = f16,
202      b_type = f16,
203      c_type = f16,
204      result_type = f16,
205      scope = 3: i32
206    }]
207  }>
208} { return }
209
210// -----
211
212//===----------------------------------------------------------------------===//
213// spv.vce
214//===----------------------------------------------------------------------===//
215
216func @vce_wrong_type() attributes {
217  // expected-error @+1 {{expected valid keyword}}
218  vce = #spv.vce<64>
219} { return }
220
221// -----
222
223func @vce_missing_fields() attributes {
224  // expected-error @+1 {{expected ','}}
225  vce = #spv.vce<v1.0>
226} { return }
227
228// -----
229
230func @vce_wrong_version() attributes {
231  // expected-error @+1 {{unknown version: V_x_y}}
232  vce = #spv.vce<V_x_y, []>
233} { return }
234
235// -----
236
237func @vce_wrong_extension_type() attributes {
238  // expected-error @+1 {{expected valid keyword}}
239  vce = #spv.vce<v1.0, [32: i32], [Shader]>
240} { return }
241
242// -----
243
244func @vce_wrong_extension() attributes {
245  // expected-error @+1 {{unknown extension: SPV_Something}}
246  vce = #spv.vce<v1.0, [Shader], [SPV_Something]>
247} { return }
248
249// -----
250
251func @vce_wrong_capability() attributes {
252  // expected-error @+1 {{unknown capability: Something}}
253  vce = #spv.vce<v1.0, [Something], []>
254} { return }
255
256// -----
257
258func @vce() attributes {
259  // CHECK: #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>
260  vce = #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>
261} { return }
262