mlir/test/Conversion/GPUToSPIRV/module-opencl.mlir

   1 // RUN: mlir-opt -allow-unregistered-dialect -convert-gpu-to-spirv="use-64bit-index=true" -verify-diagnostics -split-input-file %s -o - | FileCheck %s
   2
   3 module attributes {
   4   gpu.container_module,
   5   spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Kernel, Addresses], []>, #spirv.resource_limits<>>
   6 } {
   7   gpu.module @kernels {
   8     // CHECK-LABEL: spirv.module @{{.*}} Physical64 OpenCL
   9     //       CHECK:   spirv.func
  10     //  CHECK-SAME:     {{%.*}}: f32
  11     //   CHECK-NOT:     spirv.interface_var_abi
  12     //  CHECK-SAME:     {{%.*}}: !spirv.ptr<!spirv.array<12 x f32>, CrossWorkgroup>
  13     //   CHECK-NOT:     spirv.interface_var_abi
  14     //  CHECK-SAME:     spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>
  15     // CHECK-LABEL:   func.func @basic_module_structure
  16     //  CHECK-SAME:     attributes {gpu.kernel}
  17     gpu.func @basic_module_structure(%arg0 : f32, %arg1 : memref<12xf32, #spirv.storage_class<CrossWorkgroup>>) kernel
  18         attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>} {
  19       gpu.return
  20     }
  21   }
  22
  23   func.func @main() {
  24     %0 = "op"() : () -> (f32)
  25     %1 = "op"() : () -> (memref<12xf32, #spirv.storage_class<CrossWorkgroup>>)
  26     %cst = arith.constant 1 : index
  27     gpu.launch_func @kernels::@basic_module_structure
  28         blocks in (%cst, %cst, %cst) threads in (%cst, %cst, %cst)
  29         args(%0 : f32, %1 : memref<12xf32, #spirv.storage_class<CrossWorkgroup>>)
  30     return
  31   }
  32 }
  33
  34 // -----
  35
  36 module attributes {
  37   gpu.container_module
  38 } {
  39   gpu.module @kernels attributes {
  40     spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Kernel, Addresses], []>, #spirv.resource_limits<>>
  41   } {
  42     // CHECK-LABEL: spirv.module @{{.*}} Physical64 OpenCL
  43     //  CHECK-SAME: spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Kernel, Addresses], []>, #spirv.resource_limits<>>
  44     //       CHECK:   spirv.func
  45     //  CHECK-SAME:     {{%.*}}: f32
  46     //   CHECK-NOT:     spirv.interface_var_abi
  47     //  CHECK-SAME:     {{%.*}}: !spirv.ptr<!spirv.array<12 x f32>, CrossWorkgroup>
  48     //   CHECK-NOT:     spirv.interface_var_abi
  49     //  CHECK-SAME:     spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>
  50     // CHECK-LABEL:   func.func @basic_module_structure
  51     //  CHECK-SAME:     attributes {gpu.kernel}
  52     gpu.func @basic_module_structure(%arg0 : f32, %arg1 : memref<12xf32, #spirv.storage_class<CrossWorkgroup>>) kernel
  53         attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 4, 1]>} {
  54       gpu.return
  55     }
  56   }
  57
  58   func.func @main() {
  59     %0 = "op"() : () -> (f32)
  60     %1 = "op"() : () -> (memref<12xf32, #spirv.storage_class<CrossWorkgroup>>)
  61     %cst = arith.constant 1 : index
  62     gpu.launch_func @kernels::@basic_module_structure
  63         blocks in (%cst, %cst, %cst) threads in (%cst, %cst, %cst)
  64         args(%0 : f32, %1 : memref<12xf32, #spirv.storage_class<CrossWorkgroup>>)
  65     return
  66   }
  67 }