mlir/test/Conversion/GPUToSPIRV/entry-point.mlir

   1 // RUN: mlir-opt -test-spirv-entry-point-abi %s | FileCheck %s -check-prefix=DEFAULT
   2 // RUN: mlir-opt -test-spirv-entry-point-abi="workgroup-size=32" %s | FileCheck %s -check-prefix=WG32
   3 // RUN: mlir-opt -test-spirv-entry-point-abi="subgroup-size=4" %s | FileCheck %s -check-prefix=SG4
   4 // RUN: mlir-opt -test-spirv-entry-point-abi="target-width=32" %s | FileCheck %s -check-prefix=TW32
   5 // RUN: mlir-opt -test-spirv-entry-point-abi="workgroup-size=32,8 subgroup-size=4 target-width=32" %s | FileCheck %s -check-prefix=WG32_8-SG4-TW32
   6
   7 //      DEFAULT: gpu.func @foo()
   8 // DEFAULT-SAME: spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [1, 1, 1]>
   9
  10 //      WG32: gpu.func @foo()
  11 // WG32-SAME:  spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 1, 1]>
  12
  13 //      SG4: gpu.func @foo()
  14 // SG4-SAME:  spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [1, 1, 1], subgroup_size = 4>
  15
  16 //      TW32: gpu.func @foo()
  17 // TW32-SAME:  spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [1, 1, 1], target_width = 32>
  18
  19 //      WG32_8-SG4-TW32: gpu.func @foo()
  20 // WG32_8-SG4-TW32-SAME:  spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 8, 1], subgroup_size = 4, target_width = 32>
  21
  22 gpu.module @kernels {
  23   gpu.func @foo() kernel {
  24     gpu.return
  25   }
  26 }