mlir/test/Conversion/GPUToSPIRV/gpu-to-spirv.mlir

   1 // RUN: mlir-opt -allow-unregistered-dialect -split-input-file -convert-gpu-to-spirv -verify-diagnostics %s -o - | FileCheck %s
   2
   3 module attributes {gpu.container_module} {
   4   gpu.module @kernels {
   5     // CHECK:       spirv.module @{{.*}} Logical GLSL450 {
   6     // CHECK-LABEL: spirv.func @basic_module_structure
   7     // CHECK-SAME: {{%.*}}: f32 {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 0), StorageBuffer>}
   8     // CHECK-SAME: {{%.*}}: !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32, stride=4> [0])>, StorageBuffer> {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 1)>}
   9     // CHECK-SAME: spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>
  10     gpu.func @basic_module_structure(%arg0 : f32, %arg1 : memref<12xf32, #spirv.storage_class<StorageBuffer>>) kernel
  11       attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>} {
  12       // CHECK: spirv.Return
  13       gpu.return
  14     }
  15   }
  16
  17   func.func @main() {
  18     %0 = "op"() : () -> (f32)
  19     %1 = "op"() : () -> (memref<12xf32, #spirv.storage_class<StorageBuffer>>)
  20     %cst = arith.constant 1 : index
  21     gpu.launch_func @kernels::@basic_module_structure
  22         blocks in (%cst, %cst, %cst) threads in (%cst, %cst, %cst)
  23         args(%0 : f32, %1 : memref<12xf32, #spirv.storage_class<StorageBuffer>>)
  24     return
  25   }
  26 }
  27
  28 // -----
  29
  30 module attributes {gpu.container_module} {
  31   gpu.module @kernels {
  32     // CHECK:       spirv.module @{{.*}} Logical GLSL450 {
  33     // CHECK-LABEL: spirv.func @basic_module_structure_preset_ABI
  34     // CHECK-SAME: {{%[a-zA-Z0-9_]*}}: f32
  35     // CHECK-SAME: spirv.interface_var_abi = #spirv.interface_var_abi<(1, 2), StorageBuffer>
  36     // CHECK-SAME: !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32, stride=4> [0])>, StorageBuffer>
  37     // CHECK-SAME: spirv.interface_var_abi = #spirv.interface_var_abi<(3, 0)>
  38     // CHECK-SAME: spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>
  39     gpu.func @basic_module_structure_preset_ABI(
  40       %arg0 : f32
  41         {spirv.interface_var_abi = #spirv.interface_var_abi<(1, 2), StorageBuffer>},
  42       %arg1 : memref<12xf32, #spirv.storage_class<StorageBuffer>>
  43         {spirv.interface_var_abi = #spirv.interface_var_abi<(3, 0)>}) kernel
  44       attributes
  45         {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>} {
  46       // CHECK: spirv.Return
  47       gpu.return
  48     }
  49   }
  50 }
  51
  52 // -----
  53
  54 module attributes {gpu.container_module} {
  55   gpu.module @kernels {
  56     // expected-error @below {{failed to legalize operation 'gpu.func'}}
  57     // expected-remark @below {{match failure: missing 'spirv.entry_point_abi' attribute}}
  58     gpu.func @missing_entry_point_abi(%arg0 : f32, %arg1 : memref<12xf32, #spirv.storage_class<StorageBuffer>>) kernel {
  59       gpu.return
  60     }
  61   }
  62
  63   func.func @main() {
  64     %0 = "op"() : () -> (f32)
  65     %1 = "op"() : () -> (memref<12xf32, #spirv.storage_class<StorageBuffer>>)
  66     %cst = arith.constant 1 : index
  67     gpu.launch_func @kernels::@missing_entry_point_abi
  68         blocks in (%cst, %cst, %cst) threads in (%cst, %cst, %cst)
  69         args(%0 : f32, %1 : memref<12xf32, #spirv.storage_class<StorageBuffer>>)
  70     return
  71   }
  72 }
  73
  74 // -----
  75
  76 module attributes {gpu.container_module} {
  77   gpu.module @kernels {
  78     // expected-error @below {{failed to legalize operation 'gpu.func'}}
  79     // expected-remark @below {{match failure: missing 'spirv.interface_var_abi' attribute at argument 1}}
  80     gpu.func @missing_entry_point_abi(
  81       %arg0 : f32
  82         {spirv.interface_var_abi = #spirv.interface_var_abi<(1, 2), StorageBuffer>},
  83       %arg1 : memref<12xf32, #spirv.storage_class<StorageBuffer>>) kernel
  84     attributes
  85       {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>} {
  86       gpu.return
  87     }
  88   }
  89 }
  90
  91 // -----
  92
  93 module attributes {gpu.container_module} {
  94   gpu.module @kernels {
  95     // expected-error @below {{failed to legalize operation 'gpu.func'}}
  96     // expected-remark @below {{match failure: missing 'spirv.interface_var_abi' attribute at argument 0}}
  97     gpu.func @missing_entry_point_abi(
  98       %arg0 : f32,
  99       %arg1 : memref<12xf32, #spirv.storage_class<StorageBuffer>>
 100         {spirv.interface_var_abi = #spirv.interface_var_abi<(3, 0)>}) kernel
 101     attributes
 102       {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>} {
 103       gpu.return
 104     }
 105   }
 106 }
 107
 108 // -----
 109
 110 module attributes {gpu.container_module} {
 111   gpu.module @kernels {
 112     // CHECK-LABEL: spirv.func @barrier
 113     gpu.func @barrier(%arg0 : f32, %arg1 : memref<12xf32, #spirv.storage_class<StorageBuffer>>) kernel
 114       attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>} {
 115       // CHECK: spirv.ControlBarrier <Workgroup>, <Workgroup>, <AcquireRelease|WorkgroupMemory>
 116       gpu.barrier
 117       gpu.return
 118     }
 119   }
 120
 121   func.func @main() {
 122     %0 = "op"() : () -> (f32)
 123     %1 = "op"() : () -> (memref<12xf32, #spirv.storage_class<StorageBuffer>>)
 124     %cst = arith.constant 1 : index
 125     gpu.launch_func @kernels::@barrier
 126         blocks in (%cst, %cst, %cst) threads in (%cst, %cst, %cst)
 127         args(%0 : f32, %1 : memref<12xf32, #spirv.storage_class<StorageBuffer>>)
 128     return
 129   }
 130 }