[MLIR][mlir-spirv-cpu-runner] A pass to emulate a call to kernel in LLVM

[georgemitenkov]

Summary

This patch introduces a pass for running
mlir-spirv-cpu-runner - LowerHostCodeToLLVMPass.

This pass emulates gpu.launch_func call in LLVM dialect and lowers
the host module code to LLVM. It removes the gpu.module, creates a
sequence of global variables that are later linked to the varables
in the kernel module, as well as a series of copies to/from
them to emulate the memory transfer to/from the host or to/from the
device sides. It also converts the remaining Standard dialect into
LLVM dialect, emitting C wrappers.

[mravishankar]

I did an initial pass on this. Will come back and look in more detail, but have some high-level comments.

[mravishankar]

mlir/include/mlir/Conversion/Passes.td

100

Looking through this, I am not sure why this is split into two. It might be better to combine this into one pass. AFAICS you have all the information you need to convert a gpu.launch_func to

• a series of copies from argument to the global variables and then the
• Actual call to the kernel function
• a series of a copies to copy the result back.

mlir/lib/Conversion/GPUCommon/ConvertLaunchFuncToLLVMCalls.cpp

182

The GPU Dialect adds the gpu.container_module to identify kernel side. Further the LLVM -> SPIR-V lowering preserve the spv.target_env attribute. Its better to check for one or both of these attributes to be more targeted in what is treated as a kernel module?

[georgemitenkov]

mlir/include/mlir/Conversion/Passes.td

100

The reason to split comes from the fact that we cannot convert Standard to LLVM before our pass (Then gpu.launch_func arguments' types do not match as memrefs are converted into memref descriptors in LLVM). So we would have something like this:

"gpu.launch_func"(%24, %24, %24, %24, %24, %24, %16) {kernel = @kernels::@simple} : (!llvm.i64, !llvm.i64, !llvm.i64, !llvm.i64, !llvm.i64, !llvm.i64, !llvm.struct<(ptr<i32>, ptr<i32>, i64, array<1 x i64>, array<1 x i64>)>) -> ()

Nor we can run std-to-llvm after. If we replace the launch op with llvm.call, then its arguments will also be of LLVMType. This means we cannot really write the series of copies from argument to the global variables. This is because arguments' types are in LLVM but actual values passed are still in standard.

My solution to this was to preprocess gpu.launch_func so that it can be lowered to LLVM as a normal function, and then replace it as intended when everything is in LLVM. There might be something I am missing here, I will have a look at how to combine 2 passes in one (which is a definitely a nicer way) :)

mlir/lib/Conversion/GPUCommon/ConvertLaunchFuncToLLVMCalls.cpp

182

Thanks! I am not sure how gpu.container_module in general helps because it only identifies that it contains a kernel module, but may also have other modules in it?
In this case, we can check for the module that is != gpu.container_module as we consider only 2 modules.

I think reusing spv.target_env and spv.module's attributes is a great way, but for that we have to have a way of passing this info in SPIR-V to LLVM conversion. Something like

module attires { "kernel" } { ...

Also, I am wondering why GPU-> SPIR-V module conversion doesn't pass the symbolic name? If we would have a SPIR-V module with the name "spv.{gpu module name here}" for example, we can find the kernel straight away by looking up in symbol table?

[georgemitenkov]

mlir/include/mlir/Conversion/Passes.td

100

Actually, there is a way to do it by pulling in std-to-llvm patterns. I will update the patch with a new combined version of the passes.

[georgemitenkov]

Combined 2 passes into 1.
Finding kernel module is the same for now (to be changed).

[georgemitenkov]

Added a utilty function for kernel global variable lookup.

[georgemitenkov]

Added missing change.

[mravishankar]

mlir/lib/Conversion/GPUCommon/ConvertLaunchFuncToLLVMCalls.cpp

111

Correct me if I am wrong, but it seems like the assumption here is that the structure is actually

module {
   spv.module { 
      spv.func @foo 
   }
   module @kernel_module_name {gpu.container_module} {
      gpu.func @foo
   }
   
   gpu.launch_func @kernel_module_name::foo
}

This is the case now, but that could change. We should be really looking at at the spv.module directly. There some steps though to get there

1. You probably need to add SymAttr to spv.module. As you mentioned, that when converting the gpu module to spir-v module, we can use the same symbol name
2. spv.EntryPoint actually is supposed to have a list of spv.globalVariables that are accessed within the entry point. We dont emit it currently cause it isnt needed by Vulkan side. But that is useful here to get the "arguments". The spv.EntryPoint is added by the LowerABIAttributes pass. Probably need to do it when lowering a func to spv.func where you have all the information.

This is a fairly big change though. A intermediate stage would be given that you check for a single spv.Module and a single entry point function, under that assumption, all non builtin spv.globalVariables are kernel arguments, and they have the [binding,descriptor_set] as [0, 0], [1,0], [2, 0]... and so on. So argument 0 -> [0, 0], argument 1 -> [1, 0], etc. Use that to define/find the symbol that is used to do the copy. Does this make sense.

146

Can we also append the gpu.container_module name and the kernel function name to avoid conflicts.

182

Also, I am wondering why GPU-> SPIR-V module conversion doesn't pass the symbolic name? If we would have a SPIR-V module with the name "spv.{gpu module name here}" for example, we can find the kernel straight away by looking up in symbol table?

Good point. We didnt need it, but totally agree that the spv.module should have a symbol name as well. (See comment above).

[georgemitenkov]

mlir/lib/Conversion/GPUCommon/ConvertLaunchFuncToLLVMCalls.cpp

111

Actually the current structure was

module {gpu.container_module} {
  module {
    llvm.func() @foo
  }
  gpu.module @kernel_module_name {
      gpu.func @foo
  }

  gpu.launch_func @kernel_module_name::foo
}

because we first convert SPIR-V to LLVM for easier global variable type deduction. But dealing with gpu.launch_func op before makes more sense as we do not lose spv.EntryPoint and binding/descriptor set info.

Probably need to do it when lowering a func to spv.func where you have all the information.

I think that this can be done within LowerABIAttributes where we convert spv.func(args) to spv.func()?

Then the ideal flow would be:

• look up the needed SPIR-V module by __spv_{kernel_module_name}
• look up an entry point for this kernel and get the list of spv.globalVariables associated with this kernel. I think these globals are named by

auto name = funcOp.getName().str() + "_arg_" + std::to_string(argIndex);

I think it will be better to append the module name in the beginning as well?

• Create the llvm.global (need to convert the type here) with the name:

name + "_binding0_descriptor_set_1"

The same naming convention can be used when we lower SPIR-V global's to LLVM. This would remove the need of EncodeDescriptorSetsPass then cause it can be done in ConvertSPIRVToLLVM?

[georgemitenkov]

This is a big update to the pass structure. Now it is run after GPU to SPIR-V
conversion (and all ABI lowering, etc.). This allows to use SPIR-V module,
binds of global variables for kernel call emulation. Also, we can have multiple
kernels in the program (To fully run these the linking has to be updated).

Also, kernel function and kernel operands (global variables) are renamed during
the pass to have a kernel name in front of their names.

EncodeDescriptorSetsPass is removed as binding encoding can be done in SPIR-V
to LLVM conversion (TODO).

[georgemitenkov]

Added missing file.

[georgemitenkov]

Removed EncodeDescriptorSetsPass.

[mravishankar]

mlir/lib/Conversion/GPUCommon/ConvertLaunchFuncToLLVMCalls.cpp

111

Thanks for clarifying. This makes it clear whats happening.

The LowerABIAttributes pass already does what you want for the spv.globalVariable (here). I think encoding the name as you mentioned seems fine to me. Removing an extra pass is good as well.

[georgemitenkov]

Rebase on top of other patches and minor changes: variables/function names now keep __spv__ from SPIR-V module.

[rriddle]

mlir/include/mlir/Conversion/Passes.td

101

Is this wrapped at 80 characters?

mlir/lib/Conversion/GPUCommon/ConvertLaunchFuncToLLVMCalls.cpp

97

Please avoid using distance as it is O(N) in size. You can use llvm::hasSingleElement() for this check.

98

nit: You can return this directly, it auto-converts to failure().

114

nit: Please just use module.sym_name() instead.

126

You can set an attribute via module.sym_name(newName).

253

nit: Spell out auto.

289

nit: Drop trivial braces.

[georgemitenkov]

mlir/lib/Conversion/GPUCommon/ConvertLaunchFuncToLLVMCalls.cpp

126

I am not sure this can be done like that? Do you mean using sym_nameAttr() and passing a string attribute instead?

[georgemitenkov]

Addressed comments.

[mravishankar]

This looks good! I just have a few final nits.

Btw, It would be better for me to commit this for you. Please let me know when this is ready, and I will commit this. I am setting this as "request change" only so that once you update this it shows up on my dashboard.

Thanks for addressing all the issues raised! It looks fairly clean to me now.

mlir/lib/Conversion/GPUCommon/ConvertLaunchFuncToLLVMCalls.cpp

30

Nit: Instead of macros I would rather have a static method that returns a std::string.

162

Nit: Just add { ... } around statements that span multiple lines...

177

Use OpBuilder::InsertionGuard(rewriter). It will reset the insertion point after the scope.

[georgemitenkov]

Addressed comments and rebased on master to pick up committed patches.

[georgemitenkov]

Fixed StringRef error.

[georgemitenkov]

Added a missing change.

[mravishankar]

Coming around to landing this change now. This was causing some issues with builds (internal builds within Google) and I figured a better placement of the files works better. See this patch : https://reviews.llvm.org/D89448 for the changes. I will commit that change instead. The dependent change should be modified to suit this though. I can take care of that as well

[georgemitenkov]

Moved files around as suggested by @mravishankar.

[georgemitenkov]

In D86112#2331714, @mravishankar wrote:

Coming around to landing this change now. This was causing some issues with builds (internal builds within Google) and I figured a better placement of the files works better. See this patch : https://reviews.llvm.org/D89448 for the changes. I will commit that change instead. The dependent change should be modified to suit this though. I can take care of that as well

Great, thanks! I will update the dependent runner patch accordingly.

[georgemitenkov]

Updated struct types according to https://reviews.llvm.org/D87206.

[georgemitenkov]

@mravishankar So this patch should be good to land then?

mlir/include/mlir/Conversion/GPUCommon/GPUCommonPass.h

 } // namespace LLVM
 
 using OwnedBlob = std::unique_ptr<std::vector<char>>;
 using BlobGenerator =
     std::function<OwnedBlob(const std::string &, Location, StringRef)>;
 using LoweringCallback = std::function<std::unique_ptr<llvm::Module>(
     Operation *, llvm::LLVMContext &, StringRef)>;
 
+/// Creates a pass to emulate `gpu.launch_func` call in LLVM dialect and lower
+/// the host module code to LLVM.
+///
+/// This transformation creates a sequence of global variables that are later
+/// linked to the varables in the kernel module, and a series of copies to/from
+/// them to emulate the memory transfer from the host or to the device sides. It
+/// also converts the remaining Standard dialect into LLVM dialect, emitting C
+/// wrappers.
+std::unique_ptr<OperationPass<ModuleOp>> createLowerHostCodeToLLVMPass();
+
 /// Creates a pass to convert a gpu.launch_func operation into a sequence of
 /// GPU runtime calls.
 ///
 /// This pass does not generate code to call GPU runtime APIs directly but
 /// instead uses a small wrapper library that exports a stable and conveniently
 /// typed ABI on top of GPU runtimes such as CUDA or ROCm (HIP).
 std::unique_ptr<OperationPass<ModuleOp>>
 createGpuToLLVMConversionPass(StringRef gpuBinaryAnnotation = "");

mlir/include/mlir/Conversion/Passes.td

def GpuToLLVMConversionPass : Pass<"gpu-to-llvm", "ModuleOp"> {
   let summary = "Convert GPU dialect to LLVM dialect with GPU runtime calls";
   let constructor = "mlir::createGpuToLLVMConversionPass()";
   let options = [
     Option<"gpuBinaryAnnotation", "gpu-binary-annotation", "std::string",
            "", "Annotation attribute string for GPU binary">,
   ];
 }
 
+def LowerHostCodeToLLVM : Pass<"lower-host-to-llvm", "ModuleOp"> {
+  let summary = "Lowers the host module code to LLVM dialect and converts "
+                "`gpu.launch_func` op into `llvm.call` op";
+  let constructor = "mlir::createLowerHostCodeToLLVMPass()";
+  let dependentDialects = ["LLVM::LLVMDialect"];
+}
+
 //===----------------------------------------------------------------------===//
 // GPUToNVVM
 //===----------------------------------------------------------------------===//
 
 def ConvertGpuOpsToNVVMOps : Pass<"convert-gpu-to-nvvm", "gpu::GPUModuleOp"> {
   let summary = "Generate NVVM operations for gpu operations";
   let constructor = "mlir::createLowerGpuOpsToNVVMOpsPass()";
   let dependentDialects = ["NVVM::NVVMDialect"];

mlir/lib/Conversion/GPUCommon/CMakeLists.txt

if (MLIR_ROCM_CONVERSIONS_ENABLED)
   set(AMDGPU_LIBS
     AMDGPUCodeGen
     AMDGPUDesc
     AMDGPUInfo
   )
 endif()
 
 add_mlir_conversion_library(MLIRGPUToGPURuntimeTransforms
+  ConvertLaunchFuncToLLVMCalls.cpp
   ConvertLaunchFuncToRuntimeCalls.cpp
   ConvertKernelFuncToBlob.cpp
 
   DEPENDS
   MLIRConversionPassIncGen
   intrinsics_gen
 
   LINK_COMPONENTS

mlir/lib/Conversion/GPUCommon/ConvertLaunchFuncToLLVMCalls.cpp

+//===- ConvertLaunchFuncToLLVMCalls.cpp - MLIR GPU launch to LLVM pass ----===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// This file implements passes to convert `gpu.launch_func` op into a sequence
+// of LLVM calls that emulate the host and device sides.
+//
+//===----------------------------------------------------------------------===//
+
+#include "../PassDetail.h"
+#include "mlir/Conversion/GPUCommon/GPUCommonPass.h"
+#include "mlir/Conversion/SPIRVToLLVM/ConvertSPIRVToLLVM.h"
+#include "mlir/Conversion/StandardToLLVM/ConvertStandardToLLVM.h"
+#include "mlir/Dialect/GPU/GPUDialect.h"
+#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
+#include "mlir/Dialect/SPIRV/SPIRVOps.h"
+#include "mlir/Dialect/StandardOps/IR/Ops.h"
+#include "mlir/IR/Module.h"
+#include "mlir/IR/SymbolTable.h"
+#include "mlir/Transforms/DialectConversion.h"
+
+#include "llvm/ADT/DenseMap.h"
+#include "llvm/Support/FormatVariadic.h"
+
+#define BINDING_NAME                                                           \
+  llvm::convertToSnakeFromCamelCase(                                           \
+      stringifyDecoration(spirv::Decoration::Binding))
+#define DESCRIPTOR_SET_NAME                                                    \
+  llvm::convertToSnakeFromCamelCase(                                           \
+      stringifyDecoration(spirv::Decoration::DescriptorSet))
+
+using namespace mlir;
+
+static constexpr const char kSPIRVModule[] = "__spv__";
+
+//===----------------------------------------------------------------------===//
+// Utility functions
+//===----------------------------------------------------------------------===//
+
+/// Calculates the index of the kernel's operand that is represented by the
+/// given global variable with the `bind` attribute. We assume that the index of
+/// each kernel's operand is mapped to (descriptorSet, binding) by the map:
+///   i -> (0, i)
+/// which is implemented under `LowerABIAttributesPass`.
+static unsigned calculateGlobalIndex(spirv::GlobalVariableOp op) {
+  IntegerAttr binding = op.getAttrOfType<IntegerAttr>(BINDING_NAME);
+  return binding.getInt();
+}
+
+/// Copies the given number of bytes from src to dst pointers.
+static void copy(Location loc, Value dst, Value src, Value size,
+                 OpBuilder &builder) {
+  MLIRContext *context = builder.getContext();
+  auto llvmI1Type = LLVM::LLVMType::getInt1Ty(context);
+  Value isVolatile = builder.create<LLVM::ConstantOp>(
+      loc, llvmI1Type, builder.getBoolAttr(false));
+  builder.create<LLVM::MemcpyOp>(loc, dst, src, size, isVolatile);
+}
+
+/// Encodes the binding and descriptor set numbers into a new symbolic name.
+/// The name is specified by
+///   {kernel_module_name}_{variable_name}_descriptor_set{ds}_binding{b}
+/// to avoid symbolic conflicts, where 'ds' and 'b' are descriptor set and
+/// binding numbers.
+static StringRef createGlobalVariableWithBindName(spirv::GlobalVariableOp op,
+                                                  StringRef kernelModuleName) {
+  IntegerAttr descriptorSet =
+      op.getAttrOfType<IntegerAttr>(DESCRIPTOR_SET_NAME);
+  IntegerAttr binding = op.getAttrOfType<IntegerAttr>(BINDING_NAME);
+  return StringRef(llvm::formatv("{0}_{1}_descriptor_set{2}_binding{3}",
+                                 kernelModuleName.str(), op.sym_name().str(),
+                                 std::to_string(descriptorSet.getInt()),
+                                 std::to_string(binding.getInt())));
+}
+
+/// Returns true if the given global variable has both a descriptor set number
+/// and a binding number.
+static bool hasDescriptorSetAndBinding(spirv::GlobalVariableOp op) {
+  IntegerAttr descriptorSet =
+      op.getAttrOfType<IntegerAttr>(DESCRIPTOR_SET_NAME);
+  IntegerAttr binding = op.getAttrOfType<IntegerAttr>(BINDING_NAME);
+  return descriptorSet && binding;
+}
+
+/// Fills `globalVariableMap` with SPIR-V global variables that represent kernel
+/// arguments from the given SPIR-V module. We assume that the module contains a
+/// single entry point function. Hence, all `spv.globalVariable`s with a bind
+/// attribute are kernel arguments.
+static LogicalResult getKernelGlobalVariables(
+    spirv::ModuleOp module,
+    DenseMap<uint32_t, spirv::GlobalVariableOp> &globalVariableMap) {
+  auto entryPoints = module.getOps<spirv::EntryPointOp>();
+  if (std::distance(entryPoints.begin(), entryPoints.end()) != 1) {
+    module.emitError(
+        "The module must contain exactly one entry point function");
+    return failure();
+  }
+  auto globalVariables = module.getOps<spirv::GlobalVariableOp>();
+  for (auto globalOp : globalVariables) {
+    if (hasDescriptorSetAndBinding(globalOp))
+      globalVariableMap[calculateGlobalIndex(globalOp)] = globalOp;
+  }
+  return success();
+}
+
+/// Encodes the SPIR-V module's symbolic name into the name of the entry point
+/// function.
+static LogicalResult encodeKernelName(spirv::ModuleOp module) {
+  StringRef spvModuleName =
+      module.getAttrOfType<StringAttr>(SymbolTable::getSymbolAttrName())
+          .getValue();
+  // We already know that the module contains exactly one entry point function
+  // based on `getKernelGlobalVariables()` call. Update this function's name
+  // to:
+  //   {spv_module_name}_{function_name}
+  auto entryPoint = *module.getOps<spirv::EntryPointOp>().begin();
+  StringRef funcName = entryPoint.fn();
+  auto funcOp = module.lookupSymbol<spirv::FuncOp>(funcName);
+  StringRef newFuncName = spvModuleName.str() + "_" + funcName.str();
+  if (failed(SymbolTable::replaceAllSymbolUses(funcOp, newFuncName, module)))
+    return failure();
+  SymbolTable::setSymbolName(funcOp, newFuncName);
+  return success();
+}
+
+//===----------------------------------------------------------------------===//
+// Conversion patterns
+//===----------------------------------------------------------------------===//
+
+namespace {
+
+/// Structure to group information about the variables being copied.
+struct CopyInfo {
+  Value dst;
+  Value src;
+  Value size;
+};
+
+/// This pattern emulates a call to the kernel in LLVM dialect. For that, we
+/// copy the data to the global variable (emulating device side), call the
+/// kernel as a normal void LLVM function, and copy the data back (emulating the
+/// host side).
+class GPULaunchLowering : public ConvertOpToLLVMPattern<gpu::LaunchFuncOp> {
+  using ConvertOpToLLVMPattern<gpu::LaunchFuncOp>::ConvertOpToLLVMPattern;
+
+  LogicalResult
+  matchAndRewrite(Operation *op, ArrayRef<Value> operands,
+                  ConversionPatternRewriter &rewriter) const override {
+    gpu::LaunchFuncOp launchOp = cast<gpu::LaunchFuncOp>(op);
+    MLIRContext *context = rewriter.getContext();
+    auto module = launchOp.getParentOfType<ModuleOp>();
+
+    // Get the SPIR-V module that represents the gpu kernel module. The module
+    // is named:
+    //   __spv__{kernel_module_name}
+    // based on GPU to SPIR-V conversion.
+    StringRef kernelModuleName = launchOp.getKernelModuleName();
+    std::string spvModuleName = kSPIRVModule + kernelModuleName.str();
+    auto spvModule = module.lookupSymbol<spirv::ModuleOp>(spvModuleName);
+    if (!spvModule)
+      return launchOp.emitOpError("SPIR-V kernel module '")
+             << spvModuleName << "' is not found";
+
+    // Declare kernel function in the main module so that it later can be linked
+    // with its definition from the kernel module. We know that the kernel
+    // function would have no arguments and the data is passed via global
+    // variables. The name of the kernel will be
+    //   {spv_module_name}_{kernel_function_name}
+    // to avoid symbolic name conflicts.
+    OpBuilder::InsertionGuard guard(rewriter);
+    StringRef kernelFuncName = launchOp.getKernelName();
+    StringRef newKernelFuncName = spvModuleName + "_" + kernelFuncName.str();
+    auto kernelFunc = module.lookupSymbol<LLVM::LLVMFuncOp>(newKernelFuncName);
+    if (!kernelFunc) {
+      rewriter.setInsertionPointToStart(module.getBody());
+      kernelFunc = rewriter.create<LLVM::LLVMFuncOp>(
+          rewriter.getUnknownLoc(), newKernelFuncName,
+          LLVM::LLVMType::getFunctionTy(LLVM::LLVMType::getVoidTy(context),
+                                        ArrayRef<LLVM::LLVMType>(),
+                                        /*isVarArg=*/false));
+      rewriter.setInsertionPoint(launchOp);
+    }
+
+    // Get all global variables associated with the kernel operands.
+    DenseMap<uint32_t, spirv::GlobalVariableOp> globalVariableMap;
+    if (failed(getKernelGlobalVariables(spvModule, globalVariableMap)))
+      return failure();
+
+    // Traverse kernel operands that were converted to MemRefDescriptors. For
+    // each operand, create a global variable and copy data from operand to it.
+    Location loc = launchOp.getLoc();
+    SmallVector<CopyInfo, 4> copyInfo;
+    auto numKernelOperands = launchOp.getNumKernelOperands();
+    auto kernelOperands = operands.take_back(numKernelOperands);
+    for (auto operand : llvm::enumerate(kernelOperands)) {
+      // Check if the kernel's opernad is a ranked memref.
+      auto memRefType = launchOp.getKernelOperand(operand.index())
+                            .getType()
+                            .dyn_cast<MemRefType>();
+      if (!memRefType)
+        return failure();
+
+      // Calculate the size of the memref and get the pointer to the allocated
+      // buffer.
+      SmallVector<Value, 4> sizes;
+      getMemRefDescriptorSizes(loc, memRefType, operand.value(), rewriter,
+                               sizes);
+      Value size = getCumulativeSizeInBytes(loc, memRefType.getElementType(),
+                                            sizes, rewriter);
+      MemRefDescriptor descriptor(operand.value());
+      Value src = descriptor.allocatedPtr(rewriter, loc);
+
+      // Get the global variable in the SPIR-V module that is associated with
+      // the kernel operand. Construct its new name and create a corresponding
+      // LLVM dialect global variable.
+      spirv::GlobalVariableOp spirvGlobal = globalVariableMap[operand.index()];
+      auto pointeeType =
+          spirvGlobal.type().cast<spirv::PointerType>().getPointeeType();
+      auto dstGlobalType = typeConverter.convertType(pointeeType);
+      if (!dstGlobalType)
+        return failure();
+      StringRef name =
+          createGlobalVariableWithBindName(spirvGlobal, spvModuleName);
+      // Check if this variable has already been created.
+      auto dstGlobal = module.lookupSymbol<LLVM::GlobalOp>(name);
+      if (!dstGlobal) {
+        rewriter.setInsertionPointToStart(module.getBody());
+        dstGlobal = rewriter.create<LLVM::GlobalOp>(
+            loc, dstGlobalType.cast<LLVM::LLVMType>(),
+            /*isConstant=*/false, LLVM::Linkage::Linkonce, name, Attribute());
+        rewriter.setInsertionPoint(launchOp);
+      }
+
+      // Copy the data from src operand pointer to dst global variable. Save
+      // src, dst and size so that we can copy data back after emulating the
+      // kernel call.
+      Value dst = rewriter.create<LLVM::AddressOfOp>(loc, dstGlobal);
+      copy(loc, dst, src, size, rewriter);
+
+      CopyInfo info;
+      info.dst = dst;
+      info.src = src;
+      info.size = size;
+      copyInfo.push_back(info);
+    }
+    // Create a call to the kernel and copy the data back.
+    rewriter.replaceOpWithNewOp<LLVM::CallOp>(op, kernelFunc,
+                                              ArrayRef<Value>());
+    for (auto info : copyInfo)
+      copy(loc, info.src, info.dst, info.size, rewriter);
+    return success();
+  }
+};
+
+class LowerHostCodeToLLVM
+    : public LowerHostCodeToLLVMBase<LowerHostCodeToLLVM> {
+public:
+  void runOnOperation() override {
+    ModuleOp module = getOperation();
+
+    // Erase the GPU module.
+    for (auto gpuModule :
+         llvm::make_early_inc_range(module.getOps<gpu::GPUModuleOp>()))
+      gpuModule.erase();
+
+    // Specify options to lower Standard to LLVM and pull in the conversion
+    // patterns.
+    LowerToLLVMOptions options = {
+        /*useBarePtrCallConv=*/false,
+        /*emitCWrappers=*/true,
+        /*indexBitwidth=*/kDeriveIndexBitwidthFromDataLayout};
+    auto *context = module.getContext();
+    OwningRewritePatternList patterns;
+    LLVMTypeConverter typeConverter(context, options);
+    populateStdToLLVMConversionPatterns(typeConverter, patterns);
+    patterns.insert<GPULaunchLowering>(typeConverter);
+
+    // Pull in SPIR-V type conversion patterns to convert SPIR-V global
+    // variable's type to LLVM dialect type.
+    populateSPIRVToLLVMTypeConversion(typeConverter);
+
+    ConversionTarget target(*context);
+    target.addLegalDialect<LLVM::LLVMDialect>();
+    if (failed(applyPartialConversion(module, target, patterns))) {
+      signalPassFailure();
+    }
+
+    // Finally, modify the kernel function in SPIR-V modules to avoid symbolic
+    // conflicts.
+    for (auto spvModule : module.getOps<spirv::ModuleOp>())
+      encodeKernelName(spvModule);
+  }
+};
+} // namespace
+
+std::unique_ptr<mlir::OperationPass<mlir::ModuleOp>>
+mlir::createLowerHostCodeToLLVMPass() {
+  return std::make_unique<LowerHostCodeToLLVM>();
+}

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

+// RUN: mlir-opt --lower-host-to-llvm %s | FileCheck %s
+  
+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>}>} {
+
+  //       CHECK: llvm.mlir.global linkonce @__spv__foo_bar_arg_0_descriptor_set0_binding0() : !llvm.struct<(array<6 x i32>)>
+  //       CHECK: llvm.func @__spv__foo_bar()
+
+  //       CHECK: spv.module @__spv__foo
+  //       CHECK:   spv.globalVariable @bar_arg_0 bind(0, 0) : !spv.ptr<!spv.struct<!spv.array<6 x i32, stride=4> [0]>, StorageBuffer>
+  //       CHECK:   spv.func @__spv__foo_bar
+  
+  //       CHECK:   spv.EntryPoint "GLCompute" @__spv__foo_bar
+  //       CHECK:   spv.ExecutionMode @__spv__foo_bar "LocalSize", 1, 1, 1
+
+  // CHECK-LABEL: @main
+  //       CHECK:   %[[SRC:.*]] = llvm.extractvalue %{{.*}}[0] : !llvm.struct<(ptr<i32>, ptr<i32>, i64, array<1 x i64>, array<1 x i64>)>
+  //  CHECK-NEXT:   %[[DEST:.*]] = llvm.mlir.addressof @__spv__foo_bar_arg_0_descriptor_set0_binding0 : !llvm.ptr<struct<(array<6 x i32>)>>
+  //  CHECK-NEXT:   llvm.mlir.constant(false) : !llvm.i1
+  //  CHECK-NEXT:   "llvm.intr.memcpy"(%[[DEST]], %[[SRC]], %[[SIZE:.*]], %{{.*}}) : (!llvm.ptr<struct<(array<6 x i32>)>>, !llvm.ptr<i32>, !llvm.i64, !llvm.i1) -> ()
+  //  CHECK-NEXT:   llvm.call @__spv__foo_bar() : () -> ()
+  //  CHECK-NEXT:   llvm.mlir.constant(false) : !llvm.i1
+  //  CHECK-NEXT:   "llvm.intr.memcpy"(%[[SRC]], %[[DEST]], %[[SIZE]], %{{.*}}) : (!llvm.ptr<i32>, !llvm.ptr<struct<(array<6 x i32>)>>, !llvm.i64, !llvm.i1) -> ()
+
+  spv.module @__spv__foo Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_variable_pointers]> {
+    spv.globalVariable @bar_arg_0 bind(0, 0) : !spv.ptr<!spv.struct<!spv.array<6 x i32, stride=4> [0]>, StorageBuffer>
+    spv.func @bar() "None" attributes {workgroup_attributions = 0 : i64} {
+      %0 = spv._address_of @bar_arg_0 : !spv.ptr<!spv.struct<!spv.array<6 x i32, stride=4> [0]>, StorageBuffer>
+      spv.Return
+    }
+    spv.EntryPoint "GLCompute" @bar
+    spv.ExecutionMode @bar "LocalSize", 1, 1, 1
+  }
+
+  gpu.module @foo {
+    gpu.func @bar(%arg0: memref<6xi32>) kernel attributes {spv.entry_point_abi = {local_size = dense<1> : vector<3xi32>}} {
+      gpu.return
+    }
+  }
+
+  func @main() {
+    %buffer = alloc() : memref<6xi32>
+    %one = constant 1 : index
+    "gpu.launch_func"(%one, %one, %one,
+                      %one, %one, %one,
+                      %buffer) {kernel = @foo::@bar} : (index, index, index, index, index, index, memref<6xi32>) -> ()
+    return
+  }
+}