diff --git a/mlir/include/mlir/Dialect/GPU/IR/CompilationAttrInterfaces.td b/mlir/include/mlir/Dialect/GPU/IR/CompilationAttrInterfaces.td
--- a/mlir/include/mlir/Dialect/GPU/IR/CompilationAttrInterfaces.td
+++ b/mlir/include/mlir/Dialect/GPU/IR/CompilationAttrInterfaces.td
@@ -6,7 +6,8 @@
 //
 //===----------------------------------------------------------------------===//
 //
-// This file defines interfaces for GPU target attributes.
+// This file defines interfaces for GPU target attributes & GPU object manager
+// attributes.
 //
 //===----------------------------------------------------------------------===//
 
@@ -59,4 +60,55 @@
 def GPUNonEmptyTargetArrayAttr :
   ConfinedAttr<GPUTargetArrayAttr, [ArrayMinCount<1>]>;
 
+//===----------------------------------------------------------------------===//
+// GPU object manager attribute interface.
+//===----------------------------------------------------------------------===//
+
+def GPUBinaryHandlerLLVMTranslationAttrInterface :
+    AttrInterface<"BinaryHandlerLLVMTranslationAttrInterface"> {
+  let description = [{
+    Interface for GPU object manager attributes. Attributes implementing this
+    interface manage the interaction between GPU objects and host IR.
+  }];
+  let cppNamespace = "::mlir::gpu";
+  let methods = [
+    InterfaceMethod<[{
+        Embeds a GPU object into a host LLVM module. The operation expected by
+        this method must be a GPU BinaryOp.
+
+        All attributes implementing this interface must implement this method.
+        If the method fails then it must return `failure`.
+      }],
+      "LogicalResult", "embedBinary",
+      (ins "Operation*":$binaryOp, "llvm::IRBuilderBase&":$hostBuilder,
+           "LLVM::ModuleTranslation&":$hostModuleTranslation)
+    >,
+    InterfaceMethod<[{
+        Launches a kernel inside a binary. The first argument must be a GPU
+        LaunchFuncOp, while the second one a GPU BinaryOp.
+
+        All attributes implementing this interface must implement this method.
+        If the method fails then it must return `failure`.
+      }],
+      "LogicalResult", "launchKernel",
+      (ins "Operation*":$launchFunc, "Operation*":$binaryOp,
+           "llvm::IRBuilderBase&":$hostBuilder,
+           "LLVM::ModuleTranslation&":$hostModuleTranslation)
+    >
+  ];
+}
+
+def ImplementsBinaryHandlerLLVMTranslationAttrInterface : AttrConstraint<
+  CPred<"isa<::mlir::gpu::BinaryHandlerLLVMTranslationAttrInterface>($_self)">,
+  "Attribute implementing the `BinaryHandlerLLVMTranslationAttrInterface` interface."
+>;
+
+def GPUBinaryHandlerLLVMTranslationAttr :
+    ConfinedAttr<AnyAttr, [ImplementsBinaryHandlerLLVMTranslationAttrInterface]> {
+  let description = [{
+    Generic compilation attribute implementing the `BinaryHandlerLLVMTranslationAttrInterface`
+    interface.
+  }];
+}
+
 #endif // GPU_COMPILATIONATTRINTERFACES
diff --git a/mlir/include/mlir/Dialect/GPU/IR/GPUDialect.h b/mlir/include/mlir/Dialect/GPU/IR/GPUDialect.h
--- a/mlir/include/mlir/Dialect/GPU/IR/GPUDialect.h
+++ b/mlir/include/mlir/Dialect/GPU/IR/GPUDialect.h
@@ -29,7 +29,14 @@
 #include "mlir/Interfaces/SideEffectInterfaces.h"
 #include "llvm/ADT/STLExtras.h"
 
+namespace llvm {
+class IRBuilderBase;
+}
+
 namespace mlir {
+namespace LLVM {
+class ModuleTranslation;
+}
 namespace gpu {
 
 /// Utility class for the GPU dialect to represent triples of `Value`s