mlir/test/Conversion/SPIRVToLLVM/lower-host-to-llvm-calls.mlir

   1 // RUN: mlir-opt --lower-host-to-llvm %s | FileCheck %s
   2
   3 module attributes {gpu.container_module, spv.target_env = #spv.target_env<#spv.vce<v1.0, [Shader], [SPV_KHR_variable_pointers]>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>} {
   4
   5   //       CHECK: llvm.mlir.global linkonce @__spv__foo_bar_arg_0_descriptor_set0_binding0() : !llvm.struct<(array<6 x i32>)>
   6   //       CHECK: llvm.func @__spv__foo_bar()
   7
   8   //       CHECK: spv.module @__spv__foo
   9   //       CHECK:   spv.GlobalVariable @bar_arg_0 bind(0, 0) : !spv.ptr<!spv.struct<(!spv.array<6 x i32, stride=4> [0])>, StorageBuffer>
  10   //       CHECK:   spv.func @__spv__foo_bar
  11
  12   //       CHECK:   spv.EntryPoint "GLCompute" @__spv__foo_bar
  13   //       CHECK:   spv.ExecutionMode @__spv__foo_bar "LocalSize", 1, 1, 1
  14
  15   // CHECK-LABEL: @main
  16   //       CHECK:   %[[SRC:.*]] = llvm.extractvalue %{{.*}}[0] : !llvm.struct<(ptr<i32>, ptr<i32>, i64, array<1 x i64>, array<1 x i64>)>
  17   //  CHECK-NEXT:   %[[DEST:.*]] = llvm.mlir.addressof @__spv__foo_bar_arg_0_descriptor_set0_binding0 : !llvm.ptr<struct<(array<6 x i32>)>>
  18   //  CHECK-NEXT:   llvm.mlir.constant(false) : i1
  19   //  CHECK-NEXT:   "llvm.intr.memcpy"(%[[DEST]], %[[SRC]], %[[SIZE:.*]], %{{.*}}) : (!llvm.ptr<struct<(array<6 x i32>)>>, !llvm.ptr<i32>, i64, i1) -> ()
  20   //  CHECK-NEXT:   llvm.call @__spv__foo_bar() : () -> ()
  21   //  CHECK-NEXT:   llvm.mlir.constant(false) : i1
  22   //  CHECK-NEXT:   "llvm.intr.memcpy"(%[[SRC]], %[[DEST]], %[[SIZE]], %{{.*}}) : (!llvm.ptr<i32>, !llvm.ptr<struct<(array<6 x i32>)>>, i64, i1) -> ()
  23
  24   spv.module @__spv__foo Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_variable_pointers]> {
  25     spv.GlobalVariable @bar_arg_0 bind(0, 0) : !spv.ptr<!spv.struct<(!spv.array<6 x i32, stride=4> [0])>, StorageBuffer>
  26     spv.func @bar() "None" attributes {workgroup_attributions = 0 : i64} {
  27       %0 = spv.mlir.addressof @bar_arg_0 : !spv.ptr<!spv.struct<(!spv.array<6 x i32, stride=4> [0])>, StorageBuffer>
  28       spv.Return
  29     }
  30     spv.EntryPoint "GLCompute" @bar
  31     spv.ExecutionMode @bar "LocalSize", 1, 1, 1
  32   }
  33
  34   gpu.module @foo {
  35     gpu.func @bar(%arg0: memref<6xi32>) kernel attributes {spv.entry_point_abi = {local_size = dense<1> : vector<3xi32>}} {
  36       gpu.return
  37     }
  38   }
  39
  40   func @main() {
  41     %buffer = memref.alloc() : memref<6xi32>
  42     %one = arith.constant 1 : index
  43     gpu.launch_func @foo::@bar blocks in (%one, %one, %one)
  44         threads in (%one, %one, %one) args(%buffer : memref<6xi32>)
  45     return
  46   }
  47 }