349 lines · plain
1// RUN: mlir-opt -split-input-file -verify-diagnostics %s | FileCheck %s2 3// expected-error @+1 {{found unsupported 'spirv.something' attribute on operation}}4func.func @unknown_attr_on_op() attributes {5 spirv.something = 646} { 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 return13}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 : f3220 return %0: f3221}22 23// -----24 25//===----------------------------------------------------------------------===//26// spirv.entry_point_abi27//===----------------------------------------------------------------------===//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 = 6432} { 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_abi53//===----------------------------------------------------------------------===//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// CHECK: {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 0)>}105// CHECK: {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 1)>}106spirv.ARM.Graph @interface_var(%arg: !spirv.arm.tensor<1xf32> {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 0)>}) -> (107 !spirv.arm.tensor<1xf32> {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 1)>}108) { spirv.ARM.GraphOutputs %arg : !spirv.arm.tensor<1xf32> }109 110// -----111 112//===----------------------------------------------------------------------===//113// spirv.resource_limits114//===----------------------------------------------------------------------===//115 116// CHECK-LABEL: func @resource_limits_all_default()117func.func @resource_limits_all_default() attributes {118 // CHECK-SAME: #spirv.resource_limits<>119 limits = #spirv.resource_limits<>120} { return }121 122// -----123 124// CHECK-LABEL: func @resource_limits_min_max_subgroup_size()125func.func @resource_limits_min_max_subgroup_size() attributes {126 // CHECK-SAME: #spirv.resource_limits<min_subgroup_size = 32, max_subgroup_size = 64>127 limits = #spirv.resource_limits<min_subgroup_size = 32, max_subgroup_size=64>128} { return }129 130// -----131 132//===----------------------------------------------------------------------===//133// spirv.target_env134//===----------------------------------------------------------------------===//135 136func.func @target_env() attributes {137 // CHECK: spirv.target_env = #spirv.target_env<138 // CHECK-SAME: #spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>,139 // CHECK-SAME: #spirv.resource_limits<max_compute_workgroup_size = [128, 64, 64]>>140 spirv.target_env = #spirv.target_env<141 #spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>,142 #spirv.resource_limits<143 max_compute_workgroup_size = [128, 64, 64]144 >>145} { return }146 147// -----148 149func.func @target_env_client_api() attributes {150 // CHECK: spirv.target_env = #spirv.target_env<151 // CHECK-SAME: #spirv.vce<v1.0, [], []>,152 // CHECK-SAME: api=Metal,153 // CHECK-SAME: #spirv.resource_limits<>>154 spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [], []>, api=Metal, #spirv.resource_limits<>>155} { return }156 157// -----158 159func.func @target_env_client_api() attributes {160 // CHECK: spirv.target_env = #spirv.target_env161 // CHECK-NOT: api=162 spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [], []>, api=Unknown, #spirv.resource_limits<>>163} { return }164 165// -----166 167func.func @target_env_vendor_id() attributes {168 // CHECK: spirv.target_env = #spirv.target_env<169 // CHECK-SAME: #spirv.vce<v1.0, [], []>,170 // CHECK-SAME: NVIDIA,171 // CHECK-SAME: #spirv.resource_limits<>>172 spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [], []>, NVIDIA, #spirv.resource_limits<>>173} { return }174 175// -----176 177func.func @target_env_vendor_id_device_type() attributes {178 // CHECK: spirv.target_env = #spirv.target_env<179 // CHECK-SAME: #spirv.vce<v1.0, [], []>,180 // CHECK-SAME: AMD:DiscreteGPU,181 // CHECK-SAME: #spirv.resource_limits<>>182 spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [], []>, AMD:DiscreteGPU, #spirv.resource_limits<>>183} { return }184 185// -----186 187func.func @target_env_vendor_id_device_type_device_id() attributes {188 // CHECK: spirv.target_env = #spirv.target_env<189 // CHECK-SAME: #spirv.vce<v1.0, [], []>,190 // CHECK-SAME: Qualcomm:IntegratedGPU:100925441,191 // CHECK-SAME: #spirv.resource_limits<>>192 spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [], []>, Qualcomm:IntegratedGPU:0x6040001, #spirv.resource_limits<>>193} { return }194 195// -----196 197func.func @target_env_client_api_vendor_id_device_type_device_id() attributes {198 // CHECK: spirv.target_env = #spirv.target_env<199 // CHECK-SAME: #spirv.vce<v1.0, [], []>,200 // CHECK-SAME: api=Vulkan,201 // CHECK-SAME: Qualcomm:IntegratedGPU:100925441,202 // CHECK-SAME: #spirv.resource_limits<>>203 spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [], []>, api=Vulkan, Qualcomm:IntegratedGPU:0x6040001, #spirv.resource_limits<>>204} { return }205 206// -----207 208func.func @target_env_extra_fields() attributes {209 // expected-error @+3 {{expected '>'}}210 spirv.target_env = #spirv.target_env<211 #spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>,212 #spirv.resource_limits<>,213 more_stuff214 >215} { return }216 217// -----218 219func.func @target_env_cooperative_matrix_khr() attributes{220 // CHECK: spirv.target_env = #spirv.target_env<221 // CHECK-SAME: SPV_KHR_cooperative_matrix222 // CHECK-SAME: #spirv.coop_matrix_props_khr<223 // CHECK-SAME: m_size = 8, n_size = 8, k_size = 32,224 // CHECK-SAME: a_type = i8, b_type = i8, c_type = i32,225 // CHECK-SAME: result_type = i32, acc_sat = true, scope = <Subgroup>>226 // CHECK-SAME: #spirv.coop_matrix_props_khr<227 // CHECK-SAME: m_size = 8, n_size = 8, k_size = 16,228 // CHECK-SAME: a_type = f16, b_type = f16, c_type = f16,229 // CHECK-SAME: result_type = f16, acc_sat = false, scope = <Subgroup>>230 spirv.target_env = #spirv.target_env<231 #spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class,232 SPV_KHR_cooperative_matrix]>,233 #spirv.resource_limits<234 cooperative_matrix_properties_khr = [#spirv.coop_matrix_props_khr<235 m_size = 8,236 n_size = 8,237 k_size = 32,238 a_type = i8,239 b_type = i8,240 c_type = i32,241 result_type = i32,242 acc_sat = true,243 scope = #spirv.scope<Subgroup>244 >, #spirv.coop_matrix_props_khr<245 m_size = 8,246 n_size = 8,247 k_size = 16,248 a_type = f16,249 b_type = f16,250 c_type = f16,251 result_type = f16,252 acc_sat = false,253 scope = #spirv.scope<Subgroup>254 >]255 >>256} { return }257 258// -----259 260func.func @target_env_cooperative_matrix_nv() attributes{261 // CHECK: spirv.target_env = #spirv.target_env<262 // CHECK-SAME: SPV_NV_cooperative_matrix263 // CHECK-SAME: #spirv.coop_matrix_props_nv<264 // CHECK-SAME: m_size = 8, n_size = 8, k_size = 32,265 // CHECK-SAME: a_type = i8, b_type = i8, c_type = i32,266 // CHECK-SAME: result_type = i32, scope = <Subgroup>>267 // CHECK-SAME: #spirv.coop_matrix_props_nv<268 // CHECK-SAME: m_size = 8, n_size = 8, k_size = 16,269 // CHECK-SAME: a_type = f16, b_type = f16, c_type = f16,270 // CHECK-SAME: result_type = f16, scope = <Subgroup>>271 spirv.target_env = #spirv.target_env<272 #spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class,273 SPV_NV_cooperative_matrix]>,274 #spirv.resource_limits<275 cooperative_matrix_properties_nv = [#spirv.coop_matrix_props_nv<276 m_size = 8,277 n_size = 8,278 k_size = 32,279 a_type = i8,280 b_type = i8,281 c_type = i32,282 result_type = i32,283 scope = #spirv.scope<Subgroup>284 >, #spirv.coop_matrix_props_nv<285 m_size = 8,286 n_size = 8,287 k_size = 16,288 a_type = f16,289 b_type = f16,290 c_type = f16,291 result_type = f16,292 scope = #spirv.scope<Subgroup>293 >]294 >>295} { return }296 297// -----298 299//===----------------------------------------------------------------------===//300// spirv.vce301//===----------------------------------------------------------------------===//302 303func.func @vce_wrong_type() attributes {304 // expected-error @+1 {{expected valid keyword}}305 vce = #spirv.vce<64>306} { return }307 308// -----309 310func.func @vce_missing_fields() attributes {311 // expected-error @+1 {{expected ','}}312 vce = #spirv.vce<v1.0>313} { return }314 315// -----316 317func.func @vce_wrong_version() attributes {318 // expected-error @+1 {{unknown version: V_x_y}}319 vce = #spirv.vce<V_x_y, []>320} { return }321 322// -----323 324func.func @vce_wrong_extension_type() attributes {325 // expected-error @+1 {{expected valid keyword}}326 vce = #spirv.vce<v1.0, [32: i32], [Shader]>327} { return }328 329// -----330 331func.func @vce_wrong_extension() attributes {332 // expected-error @+1 {{unknown extension: SPIRV_Something}}333 vce = #spirv.vce<v1.0, [Shader], [SPIRV_Something]>334} { return }335 336// -----337 338func.func @vce_wrong_capability() attributes {339 // expected-error @+1 {{unknown capability: Something}}340 vce = #spirv.vce<v1.0, [Something], []>341} { return }342 343// -----344 345func.func @vce() attributes {346 // CHECK: #spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>347 vce = #spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>348} { return }349