mlir/test/Dialect/SPIRV/IR/target-and-abi.mlir

   1 // RUN: mlir-opt -split-input-file -verify-diagnostics %s | FileCheck %s
   2
   3 // expected-error @+1 {{found unsupported 'spirv.something' attribute on operation}}
   4 func.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}}
  11 func.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}}
  18 func.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}}
  30 func.func @spv_entry_point() attributes {
  31   spirv.entry_point_abi = 64
  32 } { return }
  33
  34 // -----
  35
  36 func.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
  44 func.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}}
  56 func.func @interface_var(
  57   %arg0 : f32 {spirv.interface_var_abi = 64}
  58 ) { return }
  59
  60 // -----
  61
  62 func.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
  69 func.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
  76 func.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>}
  84 func.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)>}
  91 func.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}}
  98 func.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()
 109 func.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()
 117 func.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
 128 func.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
 141 func.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
 151 func.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
 159 func.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
 169 func.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
 179 func.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
 189 func.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
 200 func.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
 211 func.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
 252 func.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
 295 func.func @vce_wrong_type() attributes {
 296   // expected-error @+1 {{expected valid keyword}}
 297   vce = #spirv.vce<64>
 298 } { return }
 299
 300 // -----
 301
 302 func.func @vce_missing_fields() attributes {
 303   // expected-error @+1 {{expected ','}}
 304   vce = #spirv.vce<v1.0>
 305 } { return }
 306
 307 // -----
 308
 309 func.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
 316 func.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
 323 func.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
 330 func.func @vce_wrong_capability() attributes {
 331   // expected-error @+1 {{unknown capability: Something}}
 332   vce = #spirv.vce<v1.0, [Something], []>
 333 } { return }
 334
 335 // -----
 336
 337 func.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 }