1// RUN: mlir-opt -split-input-file -verify-diagnostics %s | FileCheck %s 2 3// expected-error @+1 {{found unsupported 'spirv.something' attribute on operation}} 4func.func @unknown_attr_on_op() attributes { 5 spirv.something = 64 6} { return } 7 8// ----- 9 10// expected-error @+1 {{found unsupported 'spirv.something' attribute on region argument}} 11func.func @unknown_attr_on_region(%arg: i32 {spirv.something}) { 12 return 13} 14 15// ----- 16 17// expected-error @+1 {{cannot attach SPIR-V attributes to region result}} 18func.func @unknown_attr_on_region() -> (i32 {spirv.something}) { 19 %0 = arith.constant 10.0 : f32 20 return %0: f32 21} 22 23// ----- 24 25//===----------------------------------------------------------------------===// 26// spirv.entry_point_abi 27//===----------------------------------------------------------------------===// 28 29// expected-error @+1 {{'spirv.entry_point_abi' attribute must be an entry point ABI attribute}} 30func.func @spv_entry_point() attributes { 31 spirv.entry_point_abi = 64 32} { return } 33 34// ----- 35 36func.func @spv_entry_point() attributes { 37 // expected-error @+2 {{failed to parse SPIRV_EntryPointABIAttr parameter 'workgroup_size' which is to be a `DenseI32ArrayAttr`}} 38 // expected-error @+1 {{expected '['}} 39 spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = 64> 40} { return } 41 42// ----- 43 44func.func @spv_entry_point() attributes { 45 // CHECK: {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [64, 1, 1]>} 46 spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [64, 1, 1]> 47} { return } 48 49// ----- 50 51//===----------------------------------------------------------------------===// 52// spirv.interface_var_abi 53//===----------------------------------------------------------------------===// 54 55// expected-error @+1 {{'spirv.interface_var_abi' must be a spirv::InterfaceVarABIAttr}} 56func.func @interface_var( 57 %arg0 : f32 {spirv.interface_var_abi = 64} 58) { return } 59 60// ----- 61 62func.func @interface_var( 63// expected-error @+1 {{missing descriptor set}} 64 %arg0 : f32 {spirv.interface_var_abi = #spirv.interface_var_abi<()>} 65) { return } 66 67// ----- 68 69func.func @interface_var( 70// expected-error @+1 {{missing binding}} 71 %arg0 : f32 {spirv.interface_var_abi = #spirv.interface_var_abi<(1,)>} 72) { return } 73 74// ----- 75 76func.func @interface_var( 77// expected-error @+1 {{unknown storage class: }} 78 %arg0 : f32 {spirv.interface_var_abi = #spirv.interface_var_abi<(1,2), Foo>} 79) { return } 80 81// ----- 82 83// CHECK: {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 1), Uniform>} 84func.func @interface_var( 85 %arg0 : f32 {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 1), Uniform>} 86) { return } 87 88// ----- 89 90// CHECK: {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 1)>} 91func.func @interface_var( 92 %arg0 : f32 {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 1)>} 93) { return } 94 95// ----- 96 97// expected-error @+1 {{'spirv.interface_var_abi' attribute cannot specify storage class when attaching to a non-scalar value}} 98func.func @interface_var( 99 %arg0 : memref<4xf32> {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 1), Uniform>} 100) { return } 101 102// ----- 103 104//===----------------------------------------------------------------------===// 105// spirv.resource_limits 106//===----------------------------------------------------------------------===// 107 108// CHECK-LABEL: func @resource_limits_all_default() 109func.func @resource_limits_all_default() attributes { 110 // CHECK-SAME: #spirv.resource_limits<> 111 limits = #spirv.resource_limits<> 112} { return } 113 114// ----- 115 116// CHECK-LABEL: func @resource_limits_min_max_subgroup_size() 117func.func @resource_limits_min_max_subgroup_size() attributes { 118 // CHECK-SAME: #spirv.resource_limits<min_subgroup_size = 32, max_subgroup_size = 64> 119 limits = #spirv.resource_limits<min_subgroup_size = 32, max_subgroup_size=64> 120} { return } 121 122// ----- 123 124//===----------------------------------------------------------------------===// 125// spirv.target_env 126//===----------------------------------------------------------------------===// 127 128func.func @target_env() attributes { 129 // CHECK: spirv.target_env = #spirv.target_env< 130 // CHECK-SAME: #spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, 131 // CHECK-SAME: #spirv.resource_limits<max_compute_workgroup_size = [128, 64, 64]>> 132 spirv.target_env = #spirv.target_env< 133 #spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, 134 #spirv.resource_limits< 135 max_compute_workgroup_size = [128, 64, 64] 136 >> 137} { return } 138 139// ----- 140 141func.func @target_env_client_api() attributes { 142 // CHECK: spirv.target_env = #spirv.target_env< 143 // CHECK-SAME: #spirv.vce<v1.0, [], []>, 144 // CHECK-SAME: api=Metal, 145 // CHECK-SAME: #spirv.resource_limits<>> 146 spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [], []>, api=Metal, #spirv.resource_limits<>> 147} { return } 148 149// ----- 150 151func.func @target_env_client_api() attributes { 152 // CHECK: spirv.target_env = #spirv.target_env 153 // CHECK-NOT: api= 154 spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [], []>, api=Unknown, #spirv.resource_limits<>> 155} { return } 156 157// ----- 158 159func.func @target_env_vendor_id() attributes { 160 // CHECK: spirv.target_env = #spirv.target_env< 161 // CHECK-SAME: #spirv.vce<v1.0, [], []>, 162 // CHECK-SAME: NVIDIA, 163 // CHECK-SAME: #spirv.resource_limits<>> 164 spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [], []>, NVIDIA, #spirv.resource_limits<>> 165} { return } 166 167// ----- 168 169func.func @target_env_vendor_id_device_type() attributes { 170 // CHECK: spirv.target_env = #spirv.target_env< 171 // CHECK-SAME: #spirv.vce<v1.0, [], []>, 172 // CHECK-SAME: AMD:DiscreteGPU, 173 // CHECK-SAME: #spirv.resource_limits<>> 174 spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [], []>, AMD:DiscreteGPU, #spirv.resource_limits<>> 175} { return } 176 177// ----- 178 179func.func @target_env_vendor_id_device_type_device_id() attributes { 180 // CHECK: spirv.target_env = #spirv.target_env< 181 // CHECK-SAME: #spirv.vce<v1.0, [], []>, 182 // CHECK-SAME: Qualcomm:IntegratedGPU:100925441, 183 // CHECK-SAME: #spirv.resource_limits<>> 184 spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [], []>, Qualcomm:IntegratedGPU:0x6040001, #spirv.resource_limits<>> 185} { return } 186 187// ----- 188 189func.func @target_env_client_api_vendor_id_device_type_device_id() attributes { 190 // CHECK: spirv.target_env = #spirv.target_env< 191 // CHECK-SAME: #spirv.vce<v1.0, [], []>, 192 // CHECK-SAME: api=Vulkan, 193 // CHECK-SAME: Qualcomm:IntegratedGPU:100925441, 194 // CHECK-SAME: #spirv.resource_limits<>> 195 spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [], []>, api=Vulkan, Qualcomm:IntegratedGPU:0x6040001, #spirv.resource_limits<>> 196} { return } 197 198// ----- 199 200func.func @target_env_extra_fields() attributes { 201 // expected-error @+3 {{expected '>'}} 202 spirv.target_env = #spirv.target_env< 203 #spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, 204 #spirv.resource_limits<>, 205 more_stuff 206 > 207} { return } 208 209// ----- 210 211func.func @target_env_cooperative_matrix_khr() attributes{ 212 // CHECK: spirv.target_env = #spirv.target_env< 213 // CHECK-SAME: SPV_KHR_cooperative_matrix 214 // CHECK-SAME: #spirv.coop_matrix_props_khr< 215 // CHECK-SAME: m_size = 8, n_size = 8, k_size = 32, 216 // CHECK-SAME: a_type = i8, b_type = i8, c_type = i32, 217 // CHECK-SAME: result_type = i32, acc_sat = true, scope = <Subgroup>> 218 // CHECK-SAME: #spirv.coop_matrix_props_khr< 219 // CHECK-SAME: m_size = 8, n_size = 8, k_size = 16, 220 // CHECK-SAME: a_type = f16, b_type = f16, c_type = f16, 221 // CHECK-SAME: result_type = f16, acc_sat = false, scope = <Subgroup>> 222 spirv.target_env = #spirv.target_env< 223 #spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class, 224 SPV_KHR_cooperative_matrix]>, 225 #spirv.resource_limits< 226 cooperative_matrix_properties_khr = [#spirv.coop_matrix_props_khr< 227 m_size = 8, 228 n_size = 8, 229 k_size = 32, 230 a_type = i8, 231 b_type = i8, 232 c_type = i32, 233 result_type = i32, 234 acc_sat = true, 235 scope = #spirv.scope<Subgroup> 236 >, #spirv.coop_matrix_props_khr< 237 m_size = 8, 238 n_size = 8, 239 k_size = 16, 240 a_type = f16, 241 b_type = f16, 242 c_type = f16, 243 result_type = f16, 244 acc_sat = false, 245 scope = #spirv.scope<Subgroup> 246 >] 247 >> 248} { return } 249 250// ----- 251 252func.func @target_env_cooperative_matrix_nv() attributes{ 253 // CHECK: spirv.target_env = #spirv.target_env< 254 // CHECK-SAME: SPV_NV_cooperative_matrix 255 // CHECK-SAME: #spirv.coop_matrix_props_nv< 256 // CHECK-SAME: m_size = 8, n_size = 8, k_size = 32, 257 // CHECK-SAME: a_type = i8, b_type = i8, c_type = i32, 258 // CHECK-SAME: result_type = i32, scope = <Subgroup>> 259 // CHECK-SAME: #spirv.coop_matrix_props_nv< 260 // CHECK-SAME: m_size = 8, n_size = 8, k_size = 16, 261 // CHECK-SAME: a_type = f16, b_type = f16, c_type = f16, 262 // CHECK-SAME: result_type = f16, scope = <Subgroup>> 263 spirv.target_env = #spirv.target_env< 264 #spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class, 265 SPV_NV_cooperative_matrix]>, 266 #spirv.resource_limits< 267 cooperative_matrix_properties_nv = [#spirv.coop_matrix_props_nv< 268 m_size = 8, 269 n_size = 8, 270 k_size = 32, 271 a_type = i8, 272 b_type = i8, 273 c_type = i32, 274 result_type = i32, 275 scope = #spirv.scope<Subgroup> 276 >, #spirv.coop_matrix_props_nv< 277 m_size = 8, 278 n_size = 8, 279 k_size = 16, 280 a_type = f16, 281 b_type = f16, 282 c_type = f16, 283 result_type = f16, 284 scope = #spirv.scope<Subgroup> 285 >] 286 >> 287} { return } 288 289// ----- 290 291//===----------------------------------------------------------------------===// 292// spirv.vce 293//===----------------------------------------------------------------------===// 294 295func.func @vce_wrong_type() attributes { 296 // expected-error @+1 {{expected valid keyword}} 297 vce = #spirv.vce<64> 298} { return } 299 300// ----- 301 302func.func @vce_missing_fields() attributes { 303 // expected-error @+1 {{expected ','}} 304 vce = #spirv.vce<v1.0> 305} { return } 306 307// ----- 308 309func.func @vce_wrong_version() attributes { 310 // expected-error @+1 {{unknown version: V_x_y}} 311 vce = #spirv.vce<V_x_y, []> 312} { return } 313 314// ----- 315 316func.func @vce_wrong_extension_type() attributes { 317 // expected-error @+1 {{expected valid keyword}} 318 vce = #spirv.vce<v1.0, [32: i32], [Shader]> 319} { return } 320 321// ----- 322 323func.func @vce_wrong_extension() attributes { 324 // expected-error @+1 {{unknown extension: SPIRV_Something}} 325 vce = #spirv.vce<v1.0, [Shader], [SPIRV_Something]> 326} { return } 327 328// ----- 329 330func.func @vce_wrong_capability() attributes { 331 // expected-error @+1 {{unknown capability: Something}} 332 vce = #spirv.vce<v1.0, [Something], []> 333} { return } 334 335// ----- 336 337func.func @vce() attributes { 338 // CHECK: #spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]> 339 vce = #spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]> 340} { return } 341