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

   1 // RUN: mlir-opt --lower-host-to-llvm %s -split-input-file | FileCheck %s
   2
   3 module attributes {gpu.container_module, spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader], [SPV_KHR_variable_pointers]>, #spirv.resource_limits<max_compute_workgroup_invocations = 128, max_compute_workgroup_size = [128, 128, 64]>>} {
   4
   5   //       CHECK: llvm.mlir.global linkonce @__spv__foo_bar_arg_0_descriptor_set0_binding0() {addr_space = 0 : i32} : !llvm.struct<(array<6 x i32>)>
   6   //       CHECK: llvm.func @__spv__foo_bar()
   7
   8   //       CHECK: spirv.module @__spv__foo
   9   //       CHECK:   spirv.GlobalVariable @bar_arg_0 bind(0, 0) : !spirv.ptr<!spirv.struct<(!spirv.array<6 x i32, stride=4> [0])>, StorageBuffer>
  10   //       CHECK:   spirv.func @__spv__foo_bar
  11
  12   //       CHECK:   spirv.EntryPoint "GLCompute" @__spv__foo_bar
  13   //       CHECK:   spirv.ExecutionMode @__spv__foo_bar "LocalSize", 1, 1, 1
  14
  15   // CHECK-LABEL: @main
  16   //       CHECK:   %[[SRC:.*]] = llvm.extractvalue %{{.*}}[0] : !llvm.struct<(ptr, ptr, 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
  18   //  CHECK-NEXT:   "llvm.intr.memcpy"(%[[DEST]], %[[SRC]], %[[SIZE:.*]]) <{isVolatile = false}> : (!llvm.ptr, !llvm.ptr, i64) -> ()
  19   //  CHECK-NEXT:   llvm.call @__spv__foo_bar() : () -> ()
  20   //  CHECK-NEXT:   "llvm.intr.memcpy"(%[[SRC]], %[[DEST]], %[[SIZE]]) <{isVolatile = false}> : (!llvm.ptr, !llvm.ptr, i64) -> ()
  21
  22   spirv.module @__spv__foo Logical GLSL450 requires #spirv.vce<v1.0, [Shader], [SPV_KHR_variable_pointers]> {
  23     spirv.GlobalVariable @bar_arg_0 bind(0, 0) : !spirv.ptr<!spirv.struct<(!spirv.array<6 x i32, stride=4> [0])>, StorageBuffer>
  24     spirv.func @bar() "None" attributes {workgroup_attributions = 0 : i64} {
  25       %0 = spirv.mlir.addressof @bar_arg_0 : !spirv.ptr<!spirv.struct<(!spirv.array<6 x i32, stride=4> [0])>, StorageBuffer>
  26       spirv.Return
  27     }
  28     spirv.EntryPoint "GLCompute" @bar
  29     spirv.ExecutionMode @bar "LocalSize", 1, 1, 1
  30   }
  31
  32   gpu.module @foo {
  33     gpu.func @bar(%arg0: memref<6xi32>) kernel attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [1, 1, 1]>} {
  34       gpu.return
  35     }
  36   }
  37
  38   func.func @main() {
  39     %buffer = memref.alloc() : memref<6xi32>
  40     %one = arith.constant 1 : index
  41     gpu.launch_func @foo::@bar blocks in (%one, %one, %one)
  42         threads in (%one, %one, %one) args(%buffer : memref<6xi32>)
  43     return
  44   }
  45 }
  46
  47 // -----
  48
  49 // Check using a specified sym_name attribute.
  50 module {
  51   spirv.module Logical GLSL450 attributes { sym_name = "spirv.sym" } {
  52     // CHECK: spirv.func @spirv.sym_bar
  53     // CHECK: spirv.EntryPoint "GLCompute" @spirv.sym_bar
  54     // CHECK: spirv.ExecutionMode @spirv.sym_bar "LocalSize", 1, 1, 1
  55     spirv.func @bar() "None" {
  56       spirv.Return
  57     }
  58     spirv.EntryPoint "GLCompute" @bar
  59     spirv.ExecutionMode @bar "LocalSize", 1, 1, 1
  60   }
  61 }
  62
  63 // -----
  64
  65 // Check using the default sym_name attribute.
  66 module {
  67   spirv.module Logical GLSL450 {
  68     // CHECK: spirv.func @__spv___bar
  69     // CHECK: spirv.EntryPoint "GLCompute" @__spv___bar
  70     // CHECK: spirv.ExecutionMode @__spv___bar "LocalSize", 1, 1, 1
  71     spirv.func @bar() "None" {
  72       spirv.Return
  73     }
  74     spirv.EntryPoint "GLCompute" @bar
  75     spirv.ExecutionMode @bar "LocalSize", 1, 1, 1
  76   }
  77 }