diff --git a/cmake/functions.cmake b/cmake/functions.cmake
index cbd173e75..16cab3e15 100644
--- a/cmake/functions.cmake
+++ b/cmake/functions.cmake
@@ -111,10 +111,17 @@ endfunction()
 
 function(gc_add_mlir_dialect_library name)
     add_mlir_dialect_library(${ARGV})
-    target_link_libraries(obj.${name} PUBLIC GcInterface)
     set_property(GLOBAL APPEND PROPERTY GC_DIALECT_LIBS ${name})
 
     if(GcInterface IN_LIST ARGN)
         target_link_libraries(obj.${name} PUBLIC GcInterface)
     endif()
-endfunction()
\ No newline at end of file
+endfunction()
+
+function(gc_add_mlir_translation_library name)
+    add_mlir_translation_library(${ARGV})
+    set_property(GLOBAL APPEND PROPERTY GC_MLIR_LIBS ${name})
+    if(GcInterface IN_LIST ARGN)
+        target_link_libraries(obj.${name} PUBLIC GcInterface)
+    endif()
+endfunction()
diff --git a/include/gc/Dialect/CMakeLists.txt b/include/gc/Dialect/CMakeLists.txt
index db17a6f99..c2fe46c4d 100644
--- a/include/gc/Dialect/CMakeLists.txt
+++ b/include/gc/Dialect/CMakeLists.txt
@@ -2,3 +2,4 @@ add_subdirectory(CPURuntime)
 add_subdirectory(OneDNNGraph)
 add_subdirectory(Microkernel)
 add_subdirectory(Linalgx)
+add_subdirectory(LLVMIR)
diff --git a/include/gc/Dialect/LLVMIR/CMakeLists.txt b/include/gc/Dialect/LLVMIR/CMakeLists.txt
new file mode 100644
index 000000000..72d155166
--- /dev/null
+++ b/include/gc/Dialect/LLVMIR/CMakeLists.txt
@@ -0,0 +1,6 @@
+add_mlir_dialect(GenOps gen)
+add_mlir_doc(GenOps GENDialect Dialects/ -gen-dialect-doc -dialect=gen)
+set(LLVM_TARGET_DEFINITIONS GenOps.td)
+mlir_tablegen(GenOpsAttributes.h.inc -gen-attrdef-decls -attrdefs-dialect=gen)
+mlir_tablegen(GenOpsAttributes.cpp.inc -gen-attrdef-defs -attrdefs-dialect=gen)
+add_public_tablegen_target(MLIRGENConversionsIncGen)
diff --git a/include/gc/Dialect/LLVMIR/GENDialect.h b/include/gc/Dialect/LLVMIR/GENDialect.h
new file mode 100644
index 000000000..4c9149fab
--- /dev/null
+++ b/include/gc/Dialect/LLVMIR/GENDialect.h
@@ -0,0 +1,22 @@
+//===-- GENDialect.h - MLIR GEN target definitions --------------*- C++ -*-===//
+//
+// This file is licensed 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
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef MLIR_DIALECT_LLVMIR_GENDIALECT_H_
+#define MLIR_DIALECT_LLVMIR_GENDIALECT_H_
+
+#include "mlir/Bytecode/BytecodeOpInterface.h"
+#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
+#include "mlir/IR/Dialect.h"
+#include "mlir/IR/OpDefinition.h"
+
+#define GET_ATTRDEF_CLASSES
+#include "gc/Dialect/LLVMIR/GenOpsAttributes.h.inc"
+
+#include "gc/Dialect/LLVMIR/GenOpsDialect.h.inc"
+
+#endif /* MLIR_DIALECT_LLVMIR_XEDEFS_H_ */
diff --git a/include/gc/Dialect/LLVMIR/GenOps.td b/include/gc/Dialect/LLVMIR/GenOps.td
new file mode 100644
index 000000000..7139d217f
--- /dev/null
+++ b/include/gc/Dialect/LLVMIR/GenOps.td
@@ -0,0 +1,75 @@
+//===-- GenOps.td - Gen dialect definition -----------------*- tablegen -*-===//
+//
+// This file is licensed 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
+//
+//===----------------------------------------------------------------------===//
+#ifndef GENIR_OPS
+#define GENIR_OPS
+
+include "mlir/Dialect/GPU/IR/CompilationAttrInterfaces.td"
+include "mlir/Dialect/LLVMIR/LLVMOpBase.td"
+include "mlir/Dialect/SPIRV/IR/SPIRVBase.td"
+include "mlir/Interfaces/SideEffectInterfaces.td"
+
+def GEN_Dialect : Dialect {
+  let name = "gen";
+  let cppNamespace = "::mlir::gen";
+  let dependentDialects = ["LLVM::LLVMDialect"];
+  let hasOperationAttrVerify = 1;
+
+  let extraClassDeclaration = [{
+    /// Get the name of the attribute used to annotate external kernel
+    /// functions.
+    static StringRef getKernelFuncAttrName() { return "gen.kernel"; }
+    /// The address space value that represents global memory.
+    static constexpr unsigned kGlobalMemoryAddressSpace = 1;
+    /// The address space value that represents shared memory.
+    static constexpr unsigned kSharedMemoryAddressSpace = 3;
+    /// The address space value that represents private memory.
+    static constexpr unsigned kPrivateMemoryAddressSpace = 0;
+  }];
+
+  let useDefaultAttributePrinterParser = 1;
+}
+
+class GEN_Attr<string attrName, string attrMnemonic, list<Trait> traits = []>
+    : AttrDef<GEN_Dialect, attrName, traits> {
+  let mnemonic = attrMnemonic;
+}
+
+def GEN_TargettAttr : GEN_Attr<"GenTarget", "target"> {
+  let description = [{
+    GPU target attribute for controlling compilation of targets. All
+    parameters decay into default values if not present.
+
+    Examples:
+
+    1. Target with default values.
+    ```
+      gpu.module @mymodule [#gen.target] attributes {...} {
+        ...
+      }
+    ```
+  }];
+  let parameters = (ins
+    DefaultValuedParameter<"int", "2", "Optimization level to apply.">:$O,
+    StringRefParameter<"Target triple.", "\"spirv64-unknown-unknown\"">:$triple,
+    StringRefParameter<"Target chip.", "\"pvc\"">:$chip
+  );
+  let assemblyFormat = [{
+    (`<` struct($O, $triple, $chip)^ `>`)?
+  }];
+  let builders = [
+    AttrBuilder<(ins CArg<"int", "2">:$optLevel,
+                     CArg<"StringRef", "\"spirv64-unknown-unknown\"">:$triple,
+                     CArg<"StringRef", "\"pvc\"">:$chip), [{
+      return Base::get($_ctxt, optLevel, triple, chip);
+    }]>
+  ];
+  let skipDefaultBuilders = 1;
+  let genVerifyDecl = 1;
+}
+
+#endif // GENIR_OPS
diff --git a/include/gc/Target/LLVM/GEN/Target.h b/include/gc/Target/LLVM/GEN/Target.h
new file mode 100644
index 000000000..462bc78e5
--- /dev/null
+++ b/include/gc/Target/LLVM/GEN/Target.h
@@ -0,0 +1,30 @@
+//===-- Target.h - MLIR GEN target registration -----------------*- C++ -*-===//
+//
+// This file is licensed 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 provides registration calls for attaching the Gen target interface.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef MLIR_TARGET_GEN_TARGET_H
+#define MLIR_TARGET_GEN_TARGET_H
+
+namespace mlir {
+class DialectRegistry;
+class MLIRContext;
+namespace gen {
+/// Registers the `TargetAttrInterface` for the `#gen.target` attribute in
+/// the given registry.
+void registerGenTargetInterfaceExternalModels(DialectRegistry &registry);
+
+/// Registers the `TargetAttrInterface` for the `#gen.target` attribute in
+/// the registry associated with the given context.
+void registerGenTargetInterfaceExternalModels(MLIRContext &context);
+} // namespace gen
+} // namespace mlir
+
+#endif // MLIR_TARGET_GEN_TARGET_H
diff --git a/include/gc/Target/LLVM/GEN/Utils.h b/include/gc/Target/LLVM/GEN/Utils.h
new file mode 100644
index 000000000..5017dcf67
--- /dev/null
+++ b/include/gc/Target/LLVM/GEN/Utils.h
@@ -0,0 +1,53 @@
+//===-- Utils.h - MLIR GEN target utils -------------------------*- C++ -*-===//
+//
+// This file is licensed 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 files declares GEN target related utility classes and functions.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef MLIR_TARGET_LLVM_GEN_UTILS_H
+#define MLIR_TARGET_LLVM_GEN_UTILS_H
+
+#include "gc/Dialect/LLVMIR/GENDialect.h"
+#include "mlir/Dialect/GPU/IR/CompilationInterfaces.h"
+#include "mlir/Target/LLVM/ModuleToObject.h"
+
+namespace mlir {
+namespace gen {
+
+StringRef getONEAPIToolkitPath();
+
+/// Base class for all GEN serializations from GPU modules into binary strings.
+/// By default this class serializes into LLVM bitcode.
+class SerializeGPUModuleBase : public LLVM::ModuleToObject {
+public:
+  /// Initializes the `toolkitPath` with the path in `targetOptions` or if empty
+  /// with the path in `getONEAPIToolkitPath`.
+  SerializeGPUModuleBase(Operation &module, GenTargetAttr target,
+                         const gpu::TargetOptions &targetOptions = {});
+
+  // Initialize intermediate spirv target llvm backend
+  static void init();
+
+  /// Returns the target attribute.
+  GenTargetAttr getTarget() const;
+
+  /// Returns the ONEAPI toolkit path.
+  StringRef getToolkitPath() const;
+
+protected:
+  /// GEN target attribute.
+  GenTargetAttr target;
+
+  /// ONEAPI toolkit path.
+  std::string toolkitPath;
+};
+} // namespace gen
+} // namespace mlir
+
+#endif // MLIR_TARGET_LLVM_GEN_UTILS_H
diff --git a/include/gc/Target/LLVMIR/Dialect/GEN/GENToLLVMIRTranslation.h b/include/gc/Target/LLVMIR/Dialect/GEN/GENToLLVMIRTranslation.h
new file mode 100644
index 000000000..ac0a1f134
--- /dev/null
+++ b/include/gc/Target/LLVMIR/Dialect/GEN/GENToLLVMIRTranslation.h
@@ -0,0 +1,31 @@
+//===-- GENToLLVMIRTranslation.h - GEN to LLVM IR ---------------*- C++ -*-===//
+//
+// This file is licensed 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 provides registration calls for GEN dialect to LLVM IR translation.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef MLIR_TARGET_LLVMIR_DIALECT_GEN_GENTOLLVMIRTRANSLATION_H
+#define MLIR_TARGET_LLVMIR_DIALECT_GEN_GENTOLLVMIRTRANSLATION_H
+
+namespace mlir {
+
+class DialectRegistry;
+class MLIRContext;
+
+/// Register the GEN dialect and the translation from it to the LLVM IR in the
+/// given registry;
+void registerGENDialectTranslation(DialectRegistry &registry);
+
+/// Register the GEN dialect and the translation from it in the registry
+/// associated with the given context.
+void registerGENDialectTranslation(MLIRContext &context);
+
+} // namespace mlir
+
+#endif // MLIR_TARGET_LLVMIR_DIALECT_GEN_GENTOLLVMIRTRANSLATION_H
diff --git a/include/gc/Transforms/Passes.td b/include/gc/Transforms/Passes.td
index 0d5274eb9..8525ad53c 100644
--- a/include/gc/Transforms/Passes.td
+++ b/include/gc/Transforms/Passes.td
@@ -61,4 +61,23 @@ def LinalgToXeGPU : Pass<"linalg-to-xegpu", "func::FuncOp"> {
 }
 #endif
 
+def GpuLegalizeModule: Pass<"gc-gpu-legalize-module", ""> {
+  let summary = "Legalizes a GPU module for spirv conversion.";
+  let description = [{
+    Update all the nested gpu modules with an appropriate spirv target
+    information that is used further down in the pipeline.
+  }];
+  let dependentDialects = ["gpu::GPUDialect", "spirv::SPIRVDialect"];
+}
+
+def ConvertGpuSignaturesToLLVM: Pass<"gc-gpu-signatures-to-llvm", "gpu::GPUModuleOp"> {
+  let summary = "Legalize GPU kernel signatures for runtime code conversion.";
+  let dependentDialects = ["gpu::GPUDialect", "memref::MemRefDialect"];
+}
+
+def GpuGenAttachTarget: Pass<"gc-attach-gen-target", ""> {
+  let summary = "Attaches Gen target to a GPU module.";
+  let dependentDialects = ["gpu::GPUDialect", "gen::GENDialect"];
+}
+
 #endif // GC_DIALECT_GC_PASSES
diff --git a/lib/gc/CAPI/Passes.cpp b/lib/gc/CAPI/Passes.cpp
index 07ff402e5..977b8706b 100644
--- a/lib/gc/CAPI/Passes.cpp
+++ b/lib/gc/CAPI/Passes.cpp
@@ -18,6 +18,7 @@ using namespace mlir::cpuruntime;
 
 namespace mlir::gc {
 void registerCPUPipeline();
+void registerGPUPipeline();
 } // namespace mlir::gc
 
 #ifdef __cplusplus
@@ -29,6 +30,7 @@ extern "C" {
 
 MLIR_CAPI_EXPORTED void mlirRegisterAllGCPassesAndPipelines() {
   registerCPUPipeline();
+  registerGPUPipeline();
   mlirRegisterCPURuntimePasses();
   mlirRegisterGraphCompilerPasses();
 }
diff --git a/lib/gc/CMakeLists.txt b/lib/gc/CMakeLists.txt
index 7e955ffe9..441d43b8c 100644
--- a/lib/gc/CMakeLists.txt
+++ b/lib/gc/CMakeLists.txt
@@ -2,3 +2,4 @@ add_subdirectory(CAPI)
 add_subdirectory(Dialect)
 add_subdirectory(Transforms)
 add_subdirectory(ExecutionEngine)
+add_subdirectory(Target)
diff --git a/lib/gc/Dialect/CMakeLists.txt b/lib/gc/Dialect/CMakeLists.txt
index fe07dda0d..3d1dea1ab 100644
--- a/lib/gc/Dialect/CMakeLists.txt
+++ b/lib/gc/Dialect/CMakeLists.txt
@@ -2,3 +2,4 @@ add_subdirectory(CPURuntime)
 add_subdirectory(Linalgx)
 add_subdirectory(Microkernel)
 add_subdirectory(OneDNNGraph)
+add_subdirectory(LLVMIR)
diff --git a/lib/gc/Dialect/LLVMIR/CMakeLists.txt b/lib/gc/Dialect/LLVMIR/CMakeLists.txt
new file mode 100644
index 000000000..5c46f5d99
--- /dev/null
+++ b/lib/gc/Dialect/LLVMIR/CMakeLists.txt
@@ -0,0 +1,20 @@
+gc_add_mlir_dialect_library(MLIRGENDialect
+  IR/GENDialect.cpp
+
+  ADDITIONAL_HEADER_DIRS
+  ${MLIR_MAIN_INCLUDE_DIR}/mlir/Dialect/LLVMIR
+  ${PROJECT_SOURCE_DIR}/include/gc/Dialect/LLVMIR
+
+  DEPENDS
+  MLIRGENConversionsIncGen
+
+  LINK_COMPONENTS
+  AsmParser
+  Core
+
+  LINK_LIBS PUBLIC
+  MLIRIR
+  MLIRLLVMDialect
+  MLIRSideEffectInterfaces
+  GcInterface
+)
diff --git a/lib/gc/Dialect/LLVMIR/IR/GENDialect.cpp b/lib/gc/Dialect/LLVMIR/IR/GENDialect.cpp
new file mode 100644
index 000000000..326d561e4
--- /dev/null
+++ b/lib/gc/Dialect/LLVMIR/IR/GENDialect.cpp
@@ -0,0 +1,57 @@
+//===-- GENDialect.cpp - GEN Attrs and dialect registration -----*- C++ -*-===//
+//
+// This file is licensed 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
+//
+//===----------------------------------------------------------------------===//
+#include "gc/Dialect/LLVMIR/GENDialect.h"
+
+#include "mlir/Dialect/GPU/IR/CompilationInterfaces.h"
+#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
+#include "mlir/IR/DialectImplementation.h"
+#include "llvm/ADT/TypeSwitch.h"
+
+using namespace mlir;
+using namespace gen;
+
+#include "gc/Dialect/LLVMIR/GenOpsDialect.cpp.inc"
+
+LogicalResult
+GenTargetAttr::verify(function_ref<InFlightDiagnostic()> emitError, int O,
+                      StringRef triple, StringRef chip) {
+  if (O < 0 || O > 3) {
+    emitError() << "The optimization level must be a number between 0 and 3.";
+    return failure();
+  }
+  if (triple.empty()) {
+    emitError() << "The target triple cannot be empty.";
+    return failure();
+  }
+  if (chip.empty()) {
+    emitError() << "The target chip cannot be empty.";
+    return failure();
+  }
+  return success();
+}
+
+LogicalResult GENDialect::verifyOperationAttribute(Operation *op,
+                                                   NamedAttribute attr) {
+  return success();
+}
+
+void GENDialect::initialize() {
+  // clang-tidy is confused by the registration mechanism
+  // NOLINTBEGIN
+  addAttributes<
+#define GET_ATTRDEF_LIST
+#include "gc/Dialect/LLVMIR/GenOpsAttributes.cpp.inc"
+      >();
+  // NOLINTEND
+
+  allowUnknownOperations();
+  declarePromisedInterface<gpu::TargetAttrInterface, GenTargetAttr>();
+}
+
+#define GET_ATTRDEF_CLASSES
+#include "gc/Dialect/LLVMIR/GenOpsAttributes.cpp.inc"
diff --git a/lib/gc/Target/CMakeLists.txt b/lib/gc/Target/CMakeLists.txt
new file mode 100644
index 000000000..3a8c89369
--- /dev/null
+++ b/lib/gc/Target/CMakeLists.txt
@@ -0,0 +1,2 @@
+add_subdirectory(LLVMIR)
+add_subdirectory(LLVM)
diff --git a/lib/gc/Target/LLVM/CMakeLists.txt b/lib/gc/Target/LLVM/CMakeLists.txt
new file mode 100644
index 000000000..fcc2a06b8
--- /dev/null
+++ b/lib/gc/Target/LLVM/CMakeLists.txt
@@ -0,0 +1,19 @@
+gc_add_mlir_dialect_library(MLIRGENTarget
+  GEN/Target.cpp
+
+  OBJECT
+
+  ADDITIONAL_HEADER_DIRS
+  ${MLIR_MAIN_INCLUDE_DIR}/mlir/Dialect/LLVMIR
+  ${PROJECT_SOURCE_DIR}/include/gc/Dialect/LLVMIR
+
+  LINK_LIBS PUBLIC
+  MLIRIR
+  MLIRExecutionEngineUtils
+  MLIRSupport
+  MLIRGPUDialect
+  MLIRTargetLLVM
+  LLVMSPIRVCodeGen
+  MLIRGENToLLVMIRTranslation
+  GcInterface
+  )
diff --git a/lib/gc/Target/LLVM/GEN/Target.cpp b/lib/gc/Target/LLVM/GEN/Target.cpp
new file mode 100644
index 000000000..afa172ac2
--- /dev/null
+++ b/lib/gc/Target/LLVM/GEN/Target.cpp
@@ -0,0 +1,311 @@
+//===-- Target.cpp - MLIR LLVM GEN target compilation -----------*- C++ -*-===//
+//
+// This file is licensed 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 defines GEN target related functions including registration
+// calls for the `#gen.target` compilation attribute.
+//
+//===----------------------------------------------------------------------===//
+
+#include "gc/Target/LLVM/GEN/Target.h"
+
+#include "gc/Dialect/LLVMIR/GENDialect.h"
+#include "gc/Target/LLVM/GEN/Utils.h"
+#include "mlir/Dialect/GPU/IR/CompilationInterfaces.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
+#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
+#include "mlir/IR/ExtensibleDialect.h"
+#include "mlir/Target/LLVMIR/Dialect/GPU/GPUToLLVMIRTranslation.h"
+#include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
+
+#include "llvm/IR/LegacyPassManager.h"
+#include "llvm/Support/FileSystem.h"
+#include "llvm/Support/FileUtilities.h"
+#include "llvm/Support/FormatVariadic.h"
+#include "llvm/Support/MemoryBuffer.h"
+#include "llvm/Support/Path.h"
+#include "llvm/Support/Process.h"
+#include "llvm/Support/Program.h"
+#include "llvm/Support/TargetSelect.h"
+#include "llvm/Target/TargetMachine.h"
+
+using namespace mlir;
+using namespace mlir::gen;
+
+namespace {
+// Gen implementation of the gpu:TargetAttrInterface.
+class GenTargetAttrImpl
+    : public gpu::TargetAttrInterface::FallbackModel<GenTargetAttrImpl> {
+public:
+  std::optional<SmallVector<char, 0>>
+  serializeToObject(Attribute attribute, Operation *module,
+                    const gpu::TargetOptions &options) const;
+
+  Attribute createObject(Attribute attribute,
+                         const SmallVector<char, 0> &object,
+                         const gpu::TargetOptions &options) const;
+};
+} // namespace
+
+void mlir::gen::registerGenTargetInterfaceExternalModels(
+    DialectRegistry &registry) {
+  registry.addExtension(+[](MLIRContext *ctx, gen::GENDialect *dialect) {
+    GenTargetAttr::attachInterface<GenTargetAttrImpl>(*ctx);
+  });
+}
+
+void mlir::gen::registerGenTargetInterfaceExternalModels(MLIRContext &context) {
+  DialectRegistry registry;
+  registerGenTargetInterfaceExternalModels(registry);
+  context.appendDialectRegistry(registry);
+}
+
+StringRef mlir::gen::getONEAPIToolkitPath() {
+  if (const char *var = std::getenv("ONEAPI_ROOT"))
+    return var;
+  return "/usr/";
+}
+
+SerializeGPUModuleBase::SerializeGPUModuleBase(
+    Operation &module, GenTargetAttr target,
+    const gpu::TargetOptions &targetOptions)
+    : ModuleToObject(module, target.getTriple(), target.getChip(), {},
+                     target.getO()),
+      target(target), toolkitPath(targetOptions.getToolkitPath()) {
+  if (toolkitPath.empty())
+    toolkitPath = getONEAPIToolkitPath();
+}
+
+void SerializeGPUModuleBase::init() {
+  static llvm::once_flag initializeBackendOnce;
+  llvm::call_once(initializeBackendOnce, []() {
+#if LLVM_HAS_SPIRV_TARGET
+    LLVMInitializeSPIRVTarget();
+    LLVMInitializeSPIRVTargetInfo();
+    LLVMInitializeSPIRVTargetMC();
+    LLVMInitializeSPIRVAsmPrinter();
+#endif
+  });
+}
+
+GenTargetAttr SerializeGPUModuleBase::getTarget() const { return target; }
+
+StringRef SerializeGPUModuleBase::getToolkitPath() const { return toolkitPath; }
+
+namespace {
+class GenSerializer : public SerializeGPUModuleBase {
+public:
+  GenSerializer(Operation &module, GenTargetAttr target,
+                const gpu::TargetOptions &targetOptions);
+
+  gpu::GPUModuleOp getOperation();
+
+  std::optional<SmallVector<char, 0>>
+  compileToBinary(const std::string &serializedSPV);
+
+  std::optional<SmallVector<char, 0>>
+  moduleToObject(llvm::Module &llvmModule) override;
+
+  std::optional<std::string> findTool(StringRef tool);
+
+private:
+  using TmpFile = std::pair<llvm::SmallString<128>, llvm::FileRemover>;
+  std::optional<TmpFile> createTemp(StringRef name, StringRef suffix);
+
+  std::optional<std::string>
+  translateToSPIRVBinary(llvm::Module &llvmModule,
+                         llvm::TargetMachine &targetMachine);
+  gpu::TargetOptions targetOptions;
+};
+} // namespace
+
+GenSerializer::GenSerializer(Operation &module, GenTargetAttr target,
+                             const gpu::TargetOptions &targetOptions)
+    : SerializeGPUModuleBase(module, target, targetOptions) {}
+
+gpu::GPUModuleOp GenSerializer::getOperation() {
+  return dyn_cast<gpu::GPUModuleOp>(&SerializeGPUModuleBase::getOperation());
+}
+
+std::optional<GenSerializer::TmpFile>
+GenSerializer::createTemp(StringRef name, StringRef suffix) {
+  llvm::SmallString<128> filename;
+  std::error_code ec =
+      llvm::sys::fs::createTemporaryFile(name, suffix, filename);
+  if (ec) {
+    getOperation().emitError() << "Couldn't create the temp file: `" << filename
+                               << "`, error message: " << ec.message();
+    return std::nullopt;
+  }
+  return TmpFile(filename, llvm::FileRemover(filename.c_str()));
+}
+
+std::optional<std::string> GenSerializer::findTool(StringRef tool) {
+  if (std::optional<std::string> toolPath =
+          llvm::sys::Process::FindInEnvPath("PATH", tool))
+    return *toolPath;
+  getOperation().emitError()
+      << "Couldn't find the `" << tool
+      << "` binary. Please specify the toolkit "
+         "path, add the compiler to $PATH, or set one of the environment "
+         "variables in `gen::getGENToolkitPath()`.";
+  return std::nullopt;
+}
+
+std::optional<SmallVector<char, 0>>
+GenSerializer::moduleToObject(llvm::Module &llvmModule) {
+  // Return LLVM IR if the compilation target is `offload`.
+  if (targetOptions.getCompilationTarget() == gpu::CompilationTarget::Offload)
+    return SerializeGPUModuleBase::moduleToObject(llvmModule);
+
+#if !LLVM_HAS_SPIRV_TARGET
+  getOperation()->emitError(
+      "The `SPIRV` target was not built. Please enable it when building LLVM.");
+  return std::nullopt;
+#endif // LLVM_HAS_SPIRV_TARGET
+
+  std::optional<llvm::TargetMachine *> targetMachine =
+      getOrCreateTargetMachine();
+  if (!targetMachine) {
+    getOperation().emitError() << "Target Machine unavailable for triple "
+                               << triple << ", can't compile with LLVM\n";
+    return std::nullopt;
+  }
+
+  // Return SPIRV if the compilation target is `assembly`.
+  if (targetOptions.getCompilationTarget() ==
+      gpu::CompilationTarget::Assembly) {
+    std::optional<std::string> serializedISA =
+        translateToISA(llvmModule, **targetMachine);
+    if (!serializedISA) {
+      getOperation().emitError() << "Failed translating the module to ISA.";
+      return std::nullopt;
+    }
+    // Make sure to include the null terminator.
+    StringRef bin(serializedISA->c_str(), serializedISA->size() + 1);
+    return SmallVector<char, 0>(bin.begin(), bin.end());
+  }
+
+  std::optional<std::string> serializedSPIRVBinary =
+      translateToSPIRVBinary(llvmModule, **targetMachine);
+  if (!serializedSPIRVBinary) {
+    getOperation().emitError() << "Failed translating the module to Binary.";
+    return std::nullopt;
+  }
+
+  return compileToBinary(*serializedSPIRVBinary);
+}
+
+std::optional<SmallVector<char, 0>>
+GenSerializer::compileToBinary(const std::string &serializedSPV) {
+  std::optional<std::string> ocloc = findTool("ocloc");
+  if (!ocloc)
+    return std::nullopt;
+
+  std::string basename =
+      llvm::formatv("mlir-{0}-{1}-{2}", getOperation().getNameAttr().getValue(),
+                    getTarget().getTriple(), getTarget().getChip());
+
+  std::optional<TmpFile> spvFile = createTemp(basename, "spv");
+  if (!spvFile)
+    return std::nullopt;
+  std::optional<TmpFile> binaryFile = createTemp(basename, "bin");
+  if (!binaryFile)
+    return std::nullopt;
+
+  Location loc = getOperation().getLoc();
+  std::error_code ec;
+  {
+    llvm::raw_fd_ostream spvStream(spvFile->first, ec);
+    if (ec) {
+      emitError(loc) << "Couldn't open the file: `" << spvFile->first
+                     << "`, error message: " << ec.message();
+      return std::nullopt;
+    }
+    spvStream << serializedSPV;
+    if (spvStream.has_error()) {
+      emitError(loc) << "An error occurred while writing the SPIRV to: `"
+                     << spvFile->first << "`.";
+      return std::nullopt;
+    }
+    spvStream.flush();
+  }
+
+  SmallVector<StringRef, 12> oclocArgs(
+      {StringRef("compile"), StringRef("-device"), getTarget().getChip(),
+       StringRef("-spirv_input"), StringRef("-file"), StringRef(spvFile->first),
+       StringRef("-o"), StringRef(binaryFile->first)});
+
+  std::string message;
+  if (llvm::sys::ExecuteAndWait(ocloc.value(), oclocArgs,
+                                /*Env=*/std::nullopt,
+                                /*Redirects=*/std::nullopt,
+                                /*SecondsToWait=*/0,
+                                /*MemoryLimit=*/0,
+                                /*ErrMsg=*/&message)) {
+    emitError(loc) << " ocloc invocation failed. Message:\n" << message;
+    return std::nullopt;
+  }
+  llvm::ErrorOr<std::unique_ptr<llvm::MemoryBuffer>> binaryBuffer =
+      llvm::MemoryBuffer::getFile(binaryFile->first);
+  if (!binaryBuffer) {
+    emitError(loc) << "Couldn't open the file: `" << binaryFile->first
+                   << "`, error message: " << binaryBuffer.getError().message();
+    return std::nullopt;
+  }
+  StringRef result = (*binaryBuffer)->getBuffer();
+  return SmallVector<char, 0>(result.begin(), result.end());
+}
+
+std::optional<std::string>
+GenSerializer::translateToSPIRVBinary(llvm::Module &llvmModule,
+                                      llvm::TargetMachine &targetMachine) {
+  std::string targetISA;
+  llvm::raw_string_ostream stream(targetISA);
+
+  { // Drop pstream after this to prevent the ISA from being stuck buffering
+    llvm::buffer_ostream pstream(stream);
+    llvm::legacy::PassManager codegenPasses;
+
+    if (targetMachine.addPassesToEmitFile(codegenPasses, pstream, nullptr,
+                                          llvm::CodeGenFileType::ObjectFile))
+      return std::nullopt;
+
+    codegenPasses.run(llvmModule);
+  }
+  return stream.str();
+}
+
+std::optional<SmallVector<char, 0>>
+GenTargetAttrImpl::serializeToObject(Attribute attribute, Operation *module,
+                                     const gpu::TargetOptions &options) const {
+  if (!module)
+    return std::nullopt;
+  auto gpuMod = dyn_cast<gpu::GPUModuleOp>(module);
+  if (!gpuMod) {
+    module->emitError("expected to be a gpu.module op");
+    return std::nullopt;
+  }
+
+  GenSerializer serializer(*module, cast<GenTargetAttr>(attribute), options);
+  serializer.init();
+
+  return serializer.run();
+}
+
+Attribute
+GenTargetAttrImpl::createObject(Attribute attribute,
+                                const SmallVector<char, 0> &object,
+                                const gpu::TargetOptions &options) const {
+  gpu::CompilationTarget format = options.getCompilationTarget();
+  DictionaryAttr objectProps;
+  Builder builder(attribute.getContext());
+  return builder.getAttr<gpu::ObjectAttr>(
+      attribute, format,
+      builder.getStringAttr(StringRef(object.data(), object.size())),
+      objectProps);
+}
diff --git a/lib/gc/Target/LLVMIR/CMakeLists.txt b/lib/gc/Target/LLVMIR/CMakeLists.txt
new file mode 100644
index 000000000..51b942d57
--- /dev/null
+++ b/lib/gc/Target/LLVMIR/CMakeLists.txt
@@ -0,0 +1 @@
+add_subdirectory(Dialect/GEN)
diff --git a/lib/gc/Target/LLVMIR/Dialect/GEN/CMakeLists.txt b/lib/gc/Target/LLVMIR/Dialect/GEN/CMakeLists.txt
new file mode 100644
index 000000000..40630306d
--- /dev/null
+++ b/lib/gc/Target/LLVMIR/Dialect/GEN/CMakeLists.txt
@@ -0,0 +1,17 @@
+gc_add_mlir_translation_library(MLIRGENToLLVMIRTranslation
+  GENToLLVMIRTranslation.cpp
+
+  DEPENDS
+  MLIRGENConversionsIncGen
+
+  LINK_COMPONENTS
+  Core
+
+  LINK_LIBS PUBLIC
+  MLIRIR
+  MLIRLLVMDialect
+  MLIRGENDialect
+  MLIRSupport
+  MLIRTargetLLVMIRExport
+  GcInterface
+  )
diff --git a/lib/gc/Target/LLVMIR/Dialect/GEN/GENToLLVMIRTranslation.cpp b/lib/gc/Target/LLVMIR/Dialect/GEN/GENToLLVMIRTranslation.cpp
new file mode 100644
index 000000000..6a5466455
--- /dev/null
+++ b/lib/gc/Target/LLVMIR/Dialect/GEN/GENToLLVMIRTranslation.cpp
@@ -0,0 +1,71 @@
+//===-- GENToLLVMIRTranslation.cpp - Translate GEN to LLVM IR ---*- C++ -*-===//
+//
+// This file is licensed 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 a translation between the MLIR GEN dialect and
+// LLVM IR.
+//
+//===----------------------------------------------------------------------===//
+
+#include "gc/Target/LLVMIR/Dialect/GEN/GENToLLVMIRTranslation.h"
+#include "gc/Dialect/LLVMIR/GENDialect.h"
+#include "mlir/IR/BuiltinAttributes.h"
+#include "mlir/IR/Operation.h"
+#include "mlir/Target/LLVMIR/ModuleTranslation.h"
+
+#include "llvm/IR/ConstantRange.h"
+#include "llvm/IR/IRBuilder.h"
+#include "llvm/Support/raw_ostream.h"
+
+using namespace mlir;
+using namespace mlir::LLVM;
+
+namespace {
+/// Implementation of the dialect interface that converts operations belonging
+/// to the GEN dialect to LLVM IR.
+class GENDialectLLVMIRTranslationInterface
+    : public LLVMTranslationDialectInterface {
+public:
+  using LLVMTranslationDialectInterface::LLVMTranslationDialectInterface;
+
+  /// Translates the given operation to LLVM IR using the provided IR builder
+  /// and saving the state in `moduleTranslation`.
+  LogicalResult
+  convertOperation(Operation *op, llvm::IRBuilderBase &builder,
+                   LLVM::ModuleTranslation &moduleTranslation) const final {
+    // no operations, not supposed to be called
+    return failure();
+  }
+
+  /// Attaches module-level metadata for functions marked as kernels.
+  LogicalResult
+  amendOperation(Operation *op, ArrayRef<llvm::Instruction *> instructions,
+                 NamedAttribute attribute,
+                 LLVM::ModuleTranslation &moduleTranslation) const final {
+    auto func = dyn_cast<LLVM::LLVMFuncOp>(op);
+    if (!func)
+      return failure();
+    // todo; note: migth not need it as we'll have storage classes translated
+    // already
+
+    return success();
+  }
+};
+} // namespace
+
+void mlir::registerGENDialectTranslation(DialectRegistry &registry) {
+  registry.insert<gen::GENDialect>();
+  registry.addExtension(+[](MLIRContext *ctx, gen::GENDialect *dialect) {
+    dialect->addInterfaces<GENDialectLLVMIRTranslationInterface>();
+  });
+}
+
+void mlir::registerGENDialectTranslation(MLIRContext &context) {
+  DialectRegistry registry;
+  registerGENDialectTranslation(registry);
+  context.appendDialectRegistry(registry);
+}
diff --git a/lib/gc/Transforms/CMakeLists.txt b/lib/gc/Transforms/CMakeLists.txt
index 08ae24143..ab86311fe 100644
--- a/lib/gc/Transforms/CMakeLists.txt
+++ b/lib/gc/Transforms/CMakeLists.txt
@@ -25,6 +25,4 @@ gc_add_mlir_library(GcPasses
     GcInterface
   )
 
-if(GC_ENABLE_IMEX)
-  add_subdirectory(GPU)
-endif()
+add_subdirectory(GPU)
diff --git a/lib/gc/Transforms/GPU/CMakeLists.txt b/lib/gc/Transforms/GPU/CMakeLists.txt
index 18a7434e2..454cd4c92 100644
--- a/lib/gc/Transforms/GPU/CMakeLists.txt
+++ b/lib/gc/Transforms/GPU/CMakeLists.txt
@@ -1,10 +1,15 @@
 gc_add_mlir_library(GcGpuPasses
-  LinalgToXeGPU.cpp
+  GPULegalizeModule.cpp
+  ConvertGpuSignaturesToLLVM.cpp
+  GPUAttachGenTarget.cpp
 
+  PARTIAL_SOURCES_INTENDED
   DEPENDS
     GraphCompilerPassIncGen
 
   LINK_LIBS PUBLIC
+    MLIRGENDialect
+    MLIRGENTarget
     MLIRGPUDialect
     MLIRXeGPUDialect
     MLIRGPUTransforms
@@ -18,3 +23,25 @@ gc_add_mlir_library(GcGpuPasses
     GcUtilsIR
 )
 
+if(GC_ENABLE_IMEX)
+gc_add_mlir_library(GcIMEXPasses
+  LinalgToXeGPU.cpp
+
+  PARTIAL_SOURCES_INTENDED
+  DEPENDS
+    GraphCompilerPassIncGen
+
+  LINK_LIBS PUBLIC
+    MLIRGPUDialect
+    MLIRXeGPUDialect
+    MLIRGPUTransforms
+    MLIRGPUToSPIRV
+    MLIRSCFToGPU
+    MLIRSCFToSPIRV
+    MLIRMathToSPIRV
+    MLIRControlFlowToSPIRV
+    MLIRMemRefTransforms
+    GcInterface
+    GcUtilsIR
+)
+endif()
diff --git a/lib/gc/Transforms/GPU/ConvertGpuSignaturesToLLVM.cpp b/lib/gc/Transforms/GPU/ConvertGpuSignaturesToLLVM.cpp
new file mode 100644
index 000000000..0bcdc22bf
--- /dev/null
+++ b/lib/gc/Transforms/GPU/ConvertGpuSignaturesToLLVM.cpp
@@ -0,0 +1,61 @@
+//===- ConvertGpuSignaturesToLLVM.cpp - Legalize signatures -----*- C++ -*-===//
+//
+// This file is licensed 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "gc/Transforms/Passes.h"
+
+#include "gc/Dialect/LLVMIR/GENDialect.h"
+#include "mlir/Conversion/LLVMCommon/ConversionTarget.h"
+#include "mlir/Conversion/LLVMCommon/TypeConverter.h"
+#include "mlir/Conversion/MemRefToLLVM/MemRefToLLVM.h"
+#include "mlir/Dialect/Func/IR/FuncOps.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
+#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
+#include "mlir/Dialect/MemRef/IR/MemRef.h"
+
+// TODO: replace once upstream support signature conversion
+#include "GPUOpsLowering.h"
+
+using namespace mlir;
+
+namespace mlir {
+namespace gc {
+#define GEN_PASS_DEF_CONVERTGPUSIGNATURESTOLLVM
+#include "gc/Transforms/Passes.h.inc"
+} // namespace gc
+} // namespace mlir
+
+struct ConvertGpuSignaturesToLLVM
+    : public gc::impl::ConvertGpuSignaturesToLLVMBase<
+          ConvertGpuSignaturesToLLVM> {
+  using ConvertGpuSignaturesToLLVMBase::ConvertGpuSignaturesToLLVMBase;
+
+  void runOnOperation() override;
+};
+
+void ConvertGpuSignaturesToLLVM::runOnOperation() {
+  gpu::GPUModuleOp gpuModule = getOperation();
+
+  for (auto func : gpuModule.getOps<func::FuncOp>()) {
+    func->setAttr(LLVM::LLVMDialect::getEmitCWrapperAttrName(),
+                  UnitAttr::get(&getContext()));
+  }
+
+  LLVMTypeConverter converter(gpuModule.getContext());
+  RewritePatternSet patterns(gpuModule.getContext());
+  LLVMConversionTarget target(getContext());
+
+  patterns.add<GPUReturnOpLowering>(converter);
+  patterns.add<GPUFuncOpLowering>(
+      converter, gen::GENDialect::kPrivateMemoryAddressSpace /*local*/,
+      gen::GENDialect::kSharedMemoryAddressSpace /*shared*/,
+      StringAttr::get(&converter.getContext(),
+                      gen::GENDialect::getKernelFuncAttrName()));
+
+  if (failed(applyPartialConversion(gpuModule, target, std::move(patterns))))
+    signalPassFailure();
+}
diff --git a/lib/gc/Transforms/GPU/GPUAttachGenTarget.cpp b/lib/gc/Transforms/GPU/GPUAttachGenTarget.cpp
new file mode 100644
index 000000000..8e84e66ae
--- /dev/null
+++ b/lib/gc/Transforms/GPU/GPUAttachGenTarget.cpp
@@ -0,0 +1,46 @@
+//===- GPUAttachGenTarget.cpp - Attach Gen target to gpu module -*- C++ -*-===//
+//
+// This file is licensed 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "gc/Transforms/Passes.h"
+
+#include "gc/Dialect/LLVMIR/GENDialect.h"
+#include "gc/Target/LLVM/GEN/Target.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
+#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
+
+using namespace mlir;
+
+namespace mlir {
+namespace gc {
+#define GEN_PASS_DEF_GPUGENATTACHTARGET
+#include "gc/Transforms/Passes.h.inc"
+} // namespace gc
+} // namespace mlir
+
+struct GpuGenAttachTarget
+    : public gc::impl::GpuGenAttachTargetBase<GpuGenAttachTarget> {
+  using GpuGenAttachTargetBase::GpuGenAttachTargetBase;
+
+  void runOnOperation() override;
+};
+
+void GpuGenAttachTarget::runOnOperation() {
+  OpBuilder builder(&getContext());
+  auto target =
+      builder.getAttr<gen::GenTargetAttr>(2, "spirv64-unknown-unknown");
+  getOperation()->walk([&](gpu::GPUModuleOp gpuModule) {
+    SmallVector<Attribute> targets;
+    // Temporary solution to avoid an attempt to create a spirv binary
+    // if (std::optional<ArrayAttr> attrs = gpuModule.getTargets())
+    //   targets.append(attrs->getValue().begin(), attrs->getValue().end());
+    targets.push_back(target);
+    // Remove any duplicate targets.
+    targets.erase(llvm::unique(targets), targets.end());
+    gpuModule.setTargetsAttr(builder.getArrayAttr(targets));
+  });
+}
diff --git a/lib/gc/Transforms/GPU/GPULegalizeModule.cpp b/lib/gc/Transforms/GPU/GPULegalizeModule.cpp
new file mode 100644
index 000000000..f2e697f82
--- /dev/null
+++ b/lib/gc/Transforms/GPU/GPULegalizeModule.cpp
@@ -0,0 +1,67 @@
+//===- GPULegalizeModule.cpp - Legalize target for gpu module ---*- C++ -*-===//
+//
+// This file is licensed 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "gc/Transforms/Passes.h"
+
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
+#include "mlir/Dialect/SPIRV/IR/SPIRVAttributes.h"
+#include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h"
+
+using namespace mlir;
+
+namespace mlir {
+namespace gc {
+#define GEN_PASS_DEF_GPULEGALIZEMODULE
+#include "gc/Transforms/Passes.h.inc"
+} // namespace gc
+} // namespace mlir
+
+struct GpuLegalizeModule
+    : public gc::impl::GpuLegalizeModuleBase<GpuLegalizeModule> {
+  using GpuLegalizeModuleBase::GpuLegalizeModuleBase;
+
+  void runOnOperation() override;
+};
+
+void GpuLegalizeModule::runOnOperation() {
+  OpBuilder builder(&getContext());
+  using namespace mlir::spirv;
+
+  auto version = Version::V_1_0;
+  SmallVector<Capability, 4> capabilities = {
+      Capability::Addresses, Capability::Int64, Capability::Kernel};
+  SmallVector<Extension> extensions{};
+
+  auto caps = ArrayRef<Capability>(capabilities);
+  auto exts = ArrayRef<Extension>(extensions);
+  VerCapExtAttr vce = VerCapExtAttr::get(version, caps, exts, &getContext());
+
+  auto limits = ResourceLimitsAttr::get(
+      &getContext(), /*max_compute_shared_memory_size=*/16384,
+      /*max_compute_workgroup_invocations=*/128,
+      /*max_compute_workgroup_size=*/builder.getI32ArrayAttr({128, 128, 64}),
+      /*subgroup_size=*/16,
+      /*min_subgroup_size=*/std::nullopt,
+      /*max_subgroup_size=*/std::nullopt,
+      /*cooperative_matrix_properties_khr=*/ArrayAttr{},
+      /*cooperative_matrix_properties_nv=*/ArrayAttr{});
+
+  auto target = spirv::TargetEnvAttr::get(
+      vce, limits, ClientAPI::OpenCL, Vendor::Intel, DeviceType::DiscreteGPU,
+      TargetEnvAttr::kUnknownDeviceID);
+
+  getOperation()->walk([&](gpu::GPUModuleOp gpuModule) {
+    SmallVector<Attribute> targets;
+    if (std::optional<ArrayAttr> attrs = gpuModule.getTargets())
+      targets.append(attrs->getValue().begin(), attrs->getValue().end());
+    targets.push_back(target);
+    // Remove any duplicate targets.
+    targets.erase(llvm::unique(targets), targets.end());
+    gpuModule.setTargetsAttr(builder.getArrayAttr(targets));
+  });
+}
diff --git a/lib/gc/Transforms/GPU/GPUOpsLowering.h b/lib/gc/Transforms/GPU/GPUOpsLowering.h
new file mode 100644
index 000000000..23638981c
--- /dev/null
+++ b/lib/gc/Transforms/GPU/GPUOpsLowering.h
@@ -0,0 +1,140 @@
+//===- GPUOpsLowering.h - GPU FuncOp / ReturnOp lowering --------*- C++ -*-===//
+//
+// This file is licensed 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
+//
+//===----------------------------------------------------------------------===//
+#ifndef MLIR_CONVERSION_GPUCOMMON_GPUOPSLOWERING_H_
+#define MLIR_CONVERSION_GPUCOMMON_GPUOPSLOWERING_H_
+
+#include "mlir/Conversion/LLVMCommon/Pattern.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
+#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
+
+namespace mlir {
+
+/// Lowering for gpu.dynamic.shared.memory to LLVM dialect. The pattern first
+/// create a 0-sized global array symbol similar as LLVM expects. It constructs
+/// a memref descriptor with these values and return it.
+struct GPUDynamicSharedMemoryOpLowering
+    : public ConvertOpToLLVMPattern<gpu::DynamicSharedMemoryOp> {
+  using ConvertOpToLLVMPattern<
+      gpu::DynamicSharedMemoryOp>::ConvertOpToLLVMPattern;
+  GPUDynamicSharedMemoryOpLowering(const LLVMTypeConverter &converter,
+                                   unsigned alignmentBit = 0)
+      : ConvertOpToLLVMPattern<gpu::DynamicSharedMemoryOp>(converter),
+        alignmentBit(alignmentBit) {}
+
+  LogicalResult
+  matchAndRewrite(gpu::DynamicSharedMemoryOp op, OpAdaptor adaptor,
+                  ConversionPatternRewriter &rewriter) const override;
+
+private:
+  // Alignment bit
+  unsigned alignmentBit;
+};
+
+struct GPUFuncOpLowering : ConvertOpToLLVMPattern<gpu::GPUFuncOp> {
+  GPUFuncOpLowering(
+      const LLVMTypeConverter &converter, unsigned allocaAddrSpace,
+      unsigned workgroupAddrSpace, StringAttr kernelAttributeName,
+      std::optional<StringAttr> kernelBlockSizeAttributeName = std::nullopt)
+      : ConvertOpToLLVMPattern<gpu::GPUFuncOp>(converter),
+        allocaAddrSpace(allocaAddrSpace),
+        workgroupAddrSpace(workgroupAddrSpace),
+        kernelAttributeName(kernelAttributeName),
+        kernelBlockSizeAttributeName(kernelBlockSizeAttributeName) {}
+
+  LogicalResult
+  matchAndRewrite(gpu::GPUFuncOp gpuFuncOp, OpAdaptor adaptor,
+                  ConversionPatternRewriter &rewriter) const override;
+
+private:
+  /// The address space to use for `alloca`s in private memory.
+  unsigned allocaAddrSpace;
+  /// The address space to use declaring workgroup memory.
+  unsigned workgroupAddrSpace;
+
+  /// The attribute name to use instead of `gpu.kernel`.
+  StringAttr kernelAttributeName;
+
+  /// The attribute name to to set block size
+  std::optional<StringAttr> kernelBlockSizeAttributeName;
+};
+
+/// The lowering of gpu.printf to a call to HIP hostcalls
+///
+/// Simplifies llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp, as we don't have
+/// to deal with %s (even if there were first-class strings in MLIR, they're not
+/// legal input to gpu.printf) or non-constant format strings
+struct GPUPrintfOpToHIPLowering : public ConvertOpToLLVMPattern<gpu::PrintfOp> {
+  using ConvertOpToLLVMPattern<gpu::PrintfOp>::ConvertOpToLLVMPattern;
+
+  LogicalResult
+  matchAndRewrite(gpu::PrintfOp gpuPrintfOp, gpu::PrintfOpAdaptor adaptor,
+                  ConversionPatternRewriter &rewriter) const override;
+};
+
+/// The lowering of gpu.printf to a call to an external printf() function
+///
+/// This pass will add a declaration of printf() to the GPUModule if needed
+/// and separate out the format strings into global constants. For some
+/// runtimes, such as OpenCL on AMD, this is sufficient setup, as the compiler
+/// will lower printf calls to appropriate device-side code
+struct GPUPrintfOpToLLVMCallLowering
+    : public ConvertOpToLLVMPattern<gpu::PrintfOp> {
+  GPUPrintfOpToLLVMCallLowering(const LLVMTypeConverter &converter,
+                                int addressSpace = 0)
+      : ConvertOpToLLVMPattern<gpu::PrintfOp>(converter),
+        addressSpace(addressSpace) {}
+
+  LogicalResult
+  matchAndRewrite(gpu::PrintfOp gpuPrintfOp, gpu::PrintfOpAdaptor adaptor,
+                  ConversionPatternRewriter &rewriter) const override;
+
+private:
+  int addressSpace;
+};
+
+/// Lowering of gpu.printf to a vprintf standard library.
+struct GPUPrintfOpToVPrintfLowering
+    : public ConvertOpToLLVMPattern<gpu::PrintfOp> {
+  using ConvertOpToLLVMPattern<gpu::PrintfOp>::ConvertOpToLLVMPattern;
+
+  LogicalResult
+  matchAndRewrite(gpu::PrintfOp gpuPrintfOp, gpu::PrintfOpAdaptor adaptor,
+                  ConversionPatternRewriter &rewriter) const override;
+};
+
+struct GPUReturnOpLowering : public ConvertOpToLLVMPattern<gpu::ReturnOp> {
+  using ConvertOpToLLVMPattern<gpu::ReturnOp>::ConvertOpToLLVMPattern;
+
+  LogicalResult
+  matchAndRewrite(gpu::ReturnOp op, OpAdaptor adaptor,
+                  ConversionPatternRewriter &rewriter) const override;
+};
+
+namespace impl {
+/// Unrolls op if it's operating on vectors.
+LogicalResult scalarizeVectorOp(Operation *op, ValueRange operands,
+                                ConversionPatternRewriter &rewriter,
+                                const LLVMTypeConverter &converter);
+} // namespace impl
+
+/// Rewriting that unrolls SourceOp to scalars if it's operating on vectors.
+template <typename SourceOp>
+struct ScalarizeVectorOpLowering : public ConvertOpToLLVMPattern<SourceOp> {
+public:
+  using ConvertOpToLLVMPattern<SourceOp>::ConvertOpToLLVMPattern;
+
+  LogicalResult
+  matchAndRewrite(SourceOp op, typename SourceOp::Adaptor adaptor,
+                  ConversionPatternRewriter &rewriter) const override {
+    return impl::scalarizeVectorOp(op, adaptor.getOperands(), rewriter,
+                                   *this->getTypeConverter());
+  }
+};
+} // namespace mlir
+
+#endif // MLIR_CONVERSION_GPUCOMMON_GPUOPSLOWERING_H_
diff --git a/lib/gc/Transforms/Pipeline.cpp b/lib/gc/Transforms/Pipeline.cpp
index 7d487f149..999ab0d0c 100644
--- a/lib/gc/Transforms/Pipeline.cpp
+++ b/lib/gc/Transforms/Pipeline.cpp
@@ -10,6 +10,8 @@
 #include "mlir/Dialect/Arith/Transforms/Passes.h"
 #include "mlir/Dialect/Bufferization/Transforms/OneShotAnalysis.h"
 #include "mlir/Dialect/Bufferization/Transforms/Passes.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
+#include "mlir/Dialect/GPU/Transforms/Passes.h"
 #include "mlir/Dialect/LLVMIR/LLVMDialect.h"
 #include "mlir/Dialect/LLVMIR/Transforms/Passes.h"
 #include "mlir/Dialect/Linalg/Passes.h"
@@ -149,10 +151,45 @@ void populateCPUPipeline(mlir::OpPassManager &pm) {
   populateLLVMPasses(pm);
 }
 
+void populateGPUPipeline(mlir::OpPassManager &pm) {
+  pm.addNestedPass<func::FuncOp>(createLinalgGeneralizeNamedOpsPass());
+  bufferization::OneShotBufferizationOptions options;
+  options.bufferizeFunctionBoundaries = true;
+  options.setFunctionBoundaryTypeConversion(
+      bufferization::LayoutMapOption::IdentityLayoutMap);
+  pm.addPass(bufferization::createOneShotBufferizePass(options));
+  pm.addPass(createCSEPass());
+  pm.addNestedPass<func::FuncOp>(createConvertLinalgToParallelLoopsPass());
+  pm.addNestedPass<func::FuncOp>(createGpuMapParallelLoopsPass());
+  pm.addNestedPass<func::FuncOp>(createParallelLoopToGpuPass());
+  pm.addNestedPass<func::FuncOp>(createLowerAffinePass());
+  pm.addPass(memref::createNormalizeMemRefsPass());
+  pm.addPass(createGpuKernelOutliningPass());
+  pm.addPass(createGpuLegalizeModule());
+  pm.addPass(memref::createFoldMemRefAliasOpsPass());
+  ConvertIndexToLLVMPassOptions idxOptions;
+  idxOptions.indexBitwidth = 32;
+  pm.addNestedPass<gpu::GPUModuleOp>(createConvertIndexToLLVMPass(idxOptions));
+  pm.addNestedPass<gpu::GPUModuleOp>(createConvertGpuOpsToLLVMSPVOps());
+  pm.addPass(createCanonicalizerPass());
+  pm.addNestedPass<gpu::GPUModuleOp>(createConvertGpuSignaturesToLLVM());
+  pm.addPass(createGpuToLLVMConversionPass());
+  pm.addPass(createCanonicalizerPass());
+  pm.addPass(createGpuGenAttachTarget());
+  GpuModuleToBinaryPassOptions gpuModuleToBinaryPassOptions;
+  pm.addPass(createGpuModuleToBinaryPass(gpuModuleToBinaryPassOptions));
+}
+
 void registerCPUPipeline() {
   PassPipelineRegistration<>("gc-cpu-pipeline",
                              "The CPU pipeline for Graph Compiler",
                              populateCPUPipeline);
 }
 
+void registerGPUPipeline() {
+  PassPipelineRegistration<>("gc-gpu-pipeline",
+                             "The GPU pipeline for Graph Compiler",
+                             populateGPUPipeline);
+}
+
 } // namespace mlir::gc
diff --git a/python/CMakeLists.txt b/python/CMakeLists.txt
index 68d522d26..bf88854d0 100644
--- a/python/CMakeLists.txt
+++ b/python/CMakeLists.txt
@@ -99,7 +99,9 @@ add_mlir_python_common_capi_library(GcPythonCAPI
     MLIRPythonExtension.RegisterEverything
     MLIRPythonSources.Core
 )
-target_link_libraries(GcPythonCAPI PUBLIC GcInterface)
+# todo: replace with a gpu rutnime library once we have an appropriate target
+get_property(gc_passes_libs GLOBAL PROPERTY GC_PASS_LIBS)
+target_link_libraries(GcPythonCAPI PUBLIC GcInterface ${gc_passes_libs})
 
 ################################################################################
 # Instantiation of all Python modules
diff --git a/src/dnnl/CMakeLists.txt b/src/dnnl/CMakeLists.txt
index 69e8cb29f..cea142dc2 100644
--- a/src/dnnl/CMakeLists.txt
+++ b/src/dnnl/CMakeLists.txt
@@ -24,9 +24,12 @@ set(GC_DNNL_SOURCES
   JsonParser.cpp
   dnnl_graph_compiler.cpp
 )
+# todo: replace with a gpu rutnime library once we have an appropriate target
+get_property(gc_passes_libs GLOBAL PROPERTY GC_PASS_LIBS)
 set(GC_DNNL_LINKED_LIBS
   GcJitWrapper
   GcCpuRuntime
+  ${gc_passes_libs}
 )
 gc_add_mlir_library(GcDnnl SHARED
   ${GC_DNNL_SOURCES}
diff --git a/src/gc-opt/CMakeLists.txt b/src/gc-opt/CMakeLists.txt
index 96d5ae860..54d7bf9b5 100644
--- a/src/gc-opt/CMakeLists.txt
+++ b/src/gc-opt/CMakeLists.txt
@@ -20,20 +20,24 @@ if(NOT GC_ENABLE_OPT)
   return()
 endif()
 
+# todo: this needs further cleanup
 if(GC_DEV_LINK_LLVM_DYLIB)
   set(MLIR_LINK_COMPONENTS
     MLIR
   )
   get_property(dialect_libs GLOBAL PROPERTY GC_DIALECT_LIBS)
-  get_property(conversion_libs GLOBAL PROPERTY GC_PASS_LIBS)
 else()
   set(MLIR_LINK_COMPONENTS
     MLIROptLib
+    MLIRToLLVMIRTranslationRegistration
   )
   get_property(dialect_libs GLOBAL PROPERTY MLIR_DIALECT_LIBS)
-  get_property(conversion_libs GLOBAL PROPERTY MLIR_CONVERSION_LIBS)
 endif()
 
+get_property(conversion_libs GLOBAL PROPERTY MLIR_CONVERSION_LIBS)
+get_property(extension_libs GLOBAL PROPERTY MLIR_EXTENSION_LIBS)
+get_property(gc_passes_libs GLOBAL PROPERTY GC_PASS_LIBS)
+
 add_llvm_executable(gc-opt gc-opt.cpp)
 llvm_update_compile_flags(gc-opt)
 mlir_check_all_link_libraries(gc-opt)
@@ -42,8 +46,9 @@ target_link_libraries(gc-opt PUBLIC GcInterface)
 target_link_libraries(gc-opt PRIVATE
   ${dialect_libs}
   ${conversion_libs}
+  ${extension_libs}
   ${MLIR_LINK_COMPONENTS}
-  GcPasses
+  ${gc_passes_libs}
 )
 
 if(GC_ENABLE_IMEX)
@@ -52,7 +57,7 @@ if(GC_ENABLE_IMEX)
   get_property(IMEX_INCLUDES GLOBAL PROPERTY IMEX_INCLUDES)
   target_include_directories(gc-opt PRIVATE ${IMEX_INCLUDES})
   target_link_libraries(gc-opt PRIVATE
-    GcGpuPasses
+    GcGpuIMEXPasses
     IMEXGPUXDialect
     IMEXXeTileDialect
     IMEXRegionDialect
diff --git a/src/gc-opt/gc-opt.cpp b/src/gc-opt/gc-opt.cpp
index 7526106e4..4f046f4ae 100644
--- a/src/gc-opt/gc-opt.cpp
+++ b/src/gc-opt/gc-opt.cpp
@@ -18,14 +18,19 @@
  */
 
 #include "gc/Dialect/CPURuntime/Transforms/CPURuntimePasses.h"
+#include "gc/Dialect/LLVMIR/GENDialect.h"
 #include "gc/Dialect/Linalgx/LinalgxDialect.h"
 #include "gc/Dialect/Microkernel/MicrokernelDialect.h"
 #ifdef GC_HAS_ONEDNN_DIALECT
 #include "gc/Dialect/OneDNNGraph/OneDNNGraphDialect.h"
 #endif
+#include "gc/Target/LLVM/GEN/Target.h"
+#include "gc/Target/LLVMIR/Dialect/GEN/GENToLLVMIRTranslation.h"
 #include "gc/Transforms/Passes.h"
 #include "mlir/InitAllDialects.h"
+#include "mlir/InitAllExtensions.h"
 #include "mlir/InitAllPasses.h"
+#include "mlir/Target/LLVMIR/Dialect/All.h"
 #include "mlir/Tools/mlir-opt/MlirOptMain.h"
 
 #ifdef GC_USE_IMEX
@@ -35,6 +40,7 @@
 
 namespace mlir::gc {
 void registerCPUPipeline();
+void registerGPUPipeline();
 } // namespace mlir::gc
 
 int main(int argc, char *argv[]) {
@@ -49,6 +55,7 @@ int main(int argc, char *argv[]) {
 #endif
   mlir::registerAllPasses();
   mlir::gc::registerCPUPipeline();
+  mlir::gc::registerGPUPipeline();
   mlir::gc::registerGraphCompilerPasses();
   mlir::cpuruntime::registerCPURuntimePasses();
   mlir::DialectRegistry registry;
@@ -58,7 +65,16 @@ int main(int argc, char *argv[]) {
   registry.insert<mlir::cpuruntime::CPURuntimeDialect>();
   registry.insert<mlir::linalgx::LinalgxDialect>();
   registry.insert<mlir::microkernel::MicrokernelDialect>();
+  registry.insert<mlir::gen::GENDialect>();
   mlir::registerAllDialects(registry);
+  // covers lowerings for weird dialects like ub
+  // TODO: avoid `registerALL` to remove this
+  mlir::registerAllExtensions(registry);
+  // Adds missing `LLVMTranslationDialectInterface` registration for dialect for
+  // gpu.module op
+  mlir::registerAllToLLVMIRTranslations(registry);
+  mlir::gen::registerGenTargetInterfaceExternalModels(registry);
+  mlir::registerGENDialectTranslation(registry);
 #ifdef GC_USE_IMEX
   registry.insert<::imex::xetile::XeTileDialect, ::imex::gpux::GPUXDialect>();
 #endif
diff --git a/test/mlir/unittests/ExecutionEngine/CMakeLists.txt b/test/mlir/unittests/ExecutionEngine/CMakeLists.txt
index 2cfe3f77e..a063bbb3b 100644
--- a/test/mlir/unittests/ExecutionEngine/CMakeLists.txt
+++ b/test/mlir/unittests/ExecutionEngine/CMakeLists.txt
@@ -1,7 +1,10 @@
 add_mlir_unittest(GCExecutionEngineTests
   JitWrapper.cpp
 )
+# todo: remove once we have a gpu runtime library
+get_property(gc_passes_libs GLOBAL PROPERTY GC_PASS_LIBS)
 target_link_libraries(GCExecutionEngineTests
   PRIVATE
   GcJitWrapper
-  GcCpuRuntime)
+  GcCpuRuntime
+  ${gc_passes_libs})