diff --git a/mlir/include/mlir/Dialect/GPU/IR/CMakeLists.txt b/mlir/include/mlir/Dialect/GPU/IR/CMakeLists.txt
--- a/mlir/include/mlir/Dialect/GPU/IR/CMakeLists.txt
+++ b/mlir/include/mlir/Dialect/GPU/IR/CMakeLists.txt
@@ -16,6 +16,11 @@
 mlir_tablegen(GPUOpsEnums.cpp.inc -gen-enum-defs)
 add_public_tablegen_target(MLIRGPUOpsEnumsGen)
 
+set(LLVM_TARGET_DEFINITIONS CompilationAttrInterfaces.td)
+mlir_tablegen(CompilationAttrInterfaces.h.inc -gen-attr-interface-decls)
+mlir_tablegen(CompilationAttrInterfaces.cpp.inc -gen-attr-interface-defs)
+add_public_tablegen_target(MLIRGPUCompilationAttrInterfacesIncGen)
+
 set(LLVM_TARGET_DEFINITIONS GPUOps.td)
 mlir_tablegen(GPUOpsAttributes.h.inc -gen-attrdef-decls -attrdefs-dialect=gpu)
 mlir_tablegen(GPUOpsAttributes.cpp.inc -gen-attrdef-defs -attrdefs-dialect=gpu)
diff --git a/mlir/include/mlir/Dialect/GPU/IR/CompilationAttrInterfaces.td b/mlir/include/mlir/Dialect/GPU/IR/CompilationAttrInterfaces.td
new file mode 100644
--- /dev/null
+++ b/mlir/include/mlir/Dialect/GPU/IR/CompilationAttrInterfaces.td
@@ -0,0 +1,62 @@
+//===-- CompilationAttrInterfaces.td - GPU compilation interfaces ---------===//
+//
+// 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 defines interfaces for GPU target attributes.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef GPU_COMPILATIONATTRINTERFACES
+#define GPU_COMPILATIONATTRINTERFACES
+
+include "mlir/IR/AttrTypeBase.td"
+include "mlir/IR/OpBase.td"
+
+//===----------------------------------------------------------------------===//
+// GPU target attribute interface.
+//===----------------------------------------------------------------------===//
+
+def GPUTargetAttrInterface : AttrInterface<"TargetAttrInterface"> {
+  let description = [{
+    Interface for GPU target attributes. Attributes implementing this interface
+    compile GPU modules into binary objects, providing an opaque interface to
+    hide implementation details.
+  }];
+  let cppNamespace = "::mlir::gpu";
+  let methods = [
+    InterfaceMethod<[{
+        Serializes a GPU module to a string containing a representation of the
+        module.
+
+        If serialization fails then the method should return `std::nullopt`.
+
+        The `options` argument is meant to be used for passing additional
+        options that are not in the attribute.
+      }],
+      "std::optional<SmallVector<char, 0>>", "serializeToObject",
+      (ins "Operation*":$module, "const gpu::TargetOptions&":$options)>
+  ];
+}
+
+def ImplementsTargetAttrInterface : AttrConstraint<
+  CPred<"::mlir::isa<::mlir::gpu::TargetAttrInterface>($_self)">,
+  "Attribute implementing the `TargetAttrInterface` interface."
+>;
+
+def GPUTargetAttr : ConfinedAttr<AnyAttr, [ImplementsTargetAttrInterface]> {
+  let description = [{
+    Generic target attribute implementing the `TargetAttrInterface` interface.
+  }];
+}
+
+def GPUTargetArrayAttr :
+  TypedArrayAttrBase<GPUTargetAttr, "Array of GPU target attributes">;
+
+def GPUNonEmptyTargetArrayAttr :
+  ConfinedAttr<GPUTargetArrayAttr, [ArrayMinCount<1>]>;
+
+#endif // GPU_COMPILATIONATTRINTERFACES
diff --git a/mlir/include/mlir/Dialect/GPU/IR/CompilationInterfaces.h b/mlir/include/mlir/Dialect/GPU/IR/CompilationInterfaces.h
new file mode 100644
--- /dev/null
+++ b/mlir/include/mlir/Dialect/GPU/IR/CompilationInterfaces.h
@@ -0,0 +1,62 @@
+//===-- CompilationInterfaces.h - GPU compilation interfaces ----*- C++ -*-===//
+//
+// 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 defines interfaces for GPU target attributes.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef MLIR_DIALECT_GPU_IR_COMPILATIONINTERFACES_H
+#define MLIR_DIALECT_GPU_IR_COMPILATIONINTERFACES_H
+
+#include "mlir/IR/Attributes.h"
+#include "mlir/IR/Location.h"
+
+namespace mlir {
+namespace gpu {
+/// This class serves as an opaque interface for passing options to the
+/// `TargetAttrInterface` methods. Users of this class must implement the
+/// `classof` method as well as using the macros `MLIR_*_EXPLICIT_TYPE_ID` to
+/// ensure type safeness.
+class TargetOptions {
+public:
+  /// Constructor initializing the toolkit path and the list of bitcode files.
+  TargetOptions(StringRef toolkitPath = {},
+                ArrayRef<std::string> bitcodeFiles = {});
+
+  /// Returns the typeID.
+  TypeID getTypeID() const;
+
+  /// Returns the toolkit path.
+  StringRef getToolkitPath() const;
+
+  /// Returns the bitcode files to link to.
+  ArrayRef<std::string> getBitcodeFiles() const;
+
+protected:
+  /// Derived classes must use this constructor to initialize `typeID` to the
+  /// appropiate value: ie. `TargetOptions(TypeID::get<DerivedClass>())`.
+  TargetOptions(TypeID typeID, StringRef toolkitPath = {},
+                ArrayRef<std::string> bitcodeFiles = {});
+
+  /// Path to the target toolkit.
+  std::string toolkitPath;
+
+  /// List of files to link with the LLVM module.
+  SmallVector<std::string> bitcodeFiles;
+
+private:
+  TypeID typeID;
+};
+} // namespace gpu
+} // namespace mlir
+
+MLIR_DECLARE_EXPLICIT_TYPE_ID(::mlir::gpu::TargetOptions)
+
+#include "mlir/Dialect/GPU/IR/CompilationAttrInterfaces.h.inc"
+
+#endif
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
@@ -16,6 +16,7 @@
 
 #include "mlir/Bytecode/BytecodeOpInterface.h"
 #include "mlir/Dialect/DLTI/Traits.h"
+#include "mlir/Dialect/GPU/IR/CompilationInterfaces.h"
 #include "mlir/IR/Builders.h"
 #include "mlir/IR/BuiltinTypes.h"
 #include "mlir/IR/Dialect.h"
diff --git a/mlir/lib/Dialect/GPU/CMakeLists.txt b/mlir/lib/Dialect/GPU/CMakeLists.txt
--- a/mlir/lib/Dialect/GPU/CMakeLists.txt
+++ b/mlir/lib/Dialect/GPU/CMakeLists.txt
@@ -22,6 +22,7 @@
 
 add_mlir_dialect_library(MLIRGPUDialect
   IR/GPUDialect.cpp
+  IR/CompilationInterfaces.cpp
   IR/InferIntRangeInterfaceImpls.cpp
 
   ADDITIONAL_HEADER_DIRS
@@ -32,6 +33,7 @@
   MLIRGPUOpsAttributesIncGen
   MLIRGPUOpsEnumsGen
   MLIRGPUOpInterfacesIncGen
+  MLIRGPUCompilationAttrInterfacesIncGen
 
   LINK_LIBS PUBLIC
   MLIRArithDialect
diff --git a/mlir/lib/Dialect/GPU/IR/CompilationInterfaces.cpp b/mlir/lib/Dialect/GPU/IR/CompilationInterfaces.cpp
new file mode 100644
--- /dev/null
+++ b/mlir/lib/Dialect/GPU/IR/CompilationInterfaces.cpp
@@ -0,0 +1,41 @@
+//===-- CompilationInterfaces.cpp - GPU compilation interfaces --*- C++ -*-===//
+//
+// 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 defines interfaces for GPU compilation attributes.
+//
+//===----------------------------------------------------------------------===//
+
+#include "mlir/Dialect/GPU/IR/CompilationInterfaces.h"
+
+using namespace mlir;
+using namespace mlir::gpu;
+
+//===----------------------------------------------------------------------===//
+// GPU Compilation options
+//===----------------------------------------------------------------------===//
+
+TargetOptions::TargetOptions(StringRef toolkitPath,
+                             ArrayRef<std::string> bitcodeFiles)
+    : TargetOptions(TypeID::get<TargetOptions>(), toolkitPath, bitcodeFiles) {}
+
+TargetOptions::TargetOptions(TypeID typeID, StringRef toolkitPath,
+                             ArrayRef<std::string> bitcodeFiles)
+    : toolkitPath(toolkitPath.str()), bitcodeFiles(bitcodeFiles),
+      typeID(typeID) {}
+
+TypeID TargetOptions::getTypeID() const { return typeID; }
+
+StringRef TargetOptions::getToolkitPath() const { return toolkitPath; }
+
+ArrayRef<std::string> TargetOptions::getBitcodeFiles() const {
+  return bitcodeFiles;
+}
+
+MLIR_DEFINE_EXPLICIT_TYPE_ID(::mlir::gpu::TargetOptions)
+
+#include "mlir/Dialect/GPU/IR/CompilationAttrInterfaces.cpp.inc"