[mlir][gpu] Pass GPU module to TargetAttrInterface::createObject.

This patch adds an argument to `gpu::TargetAttrInterface::createObject` to pass
the GPU module. This is useful as `gpu::ObjectAttr` contains a property dict
for metadata, and passing the module allows extracting things like the symbol
table and adding it to the property dict.

Author: fabianmcg

mlir/include/mlir/Dialect/GPU/IR/CompilationAttrInterfaces.td

@@ -47,7 +47,8 @@ def GPUTargetAttrInterface : AttrInterface<"TargetAttrInterface"> {
         meant to be used for passing additional options that are not in the
         attribute.
       }], "::mlir::Attribute", "createObject",
-        (ins "const ::llvm::SmallVector<char, 0>&":$object,
+        (ins "::mlir::Operation*":$module,
+             "const ::llvm::SmallVector<char, 0>&":$object,
              "const ::mlir::gpu::TargetOptions&":$options)>
   ];
 }

mlir/lib/Target/LLVM/NVVM/Target.cpp

@@ -49,7 +49,7 @@ class NVVMTargetAttrImpl
   serializeToObject(Attribute attribute, Operation *module,
                     const gpu::TargetOptions &options) const;
 
-  Attribute createObject(Attribute attribute,
+  Attribute createObject(Attribute attribute, Operation *module,
                          const SmallVector<char, 0> &object,
                          const gpu::TargetOptions &options) const;
 };
@@ -591,7 +591,7 @@ NVVMTargetAttrImpl::serializeToObject(Attribute attribute, Operation *module,
 }
 
 Attribute
-NVVMTargetAttrImpl::createObject(Attribute attribute,
+NVVMTargetAttrImpl::createObject(Attribute attribute, Operation *module,
                                  const SmallVector<char, 0> &object,
                                  const gpu::TargetOptions &options) const {
   auto target = cast<NVVMTargetAttr>(attribute);

mlir/lib/Target/LLVM/ROCDL/Target.cpp

@@ -59,7 +59,7 @@ class ROCDLTargetAttrImpl
   serializeToObject(Attribute attribute, Operation *module,
                     const gpu::TargetOptions &options) const;
 
-  Attribute createObject(Attribute attribute,
+  Attribute createObject(Attribute attribute, Operation *module,
                          const SmallVector<char, 0> &object,
                          const gpu::TargetOptions &options) const;
 };
@@ -500,7 +500,7 @@ std::optional<SmallVector<char, 0>> ROCDLTargetAttrImpl::serializeToObject(
 }
 
 Attribute
-ROCDLTargetAttrImpl::createObject(Attribute attribute,
+ROCDLTargetAttrImpl::createObject(Attribute attribute, Operation *module,
                                   const SmallVector<char, 0> &object,
                                   const gpu::TargetOptions &options) const {
   gpu::CompilationTarget format = options.getCompilationTarget();

mlir/lib/Target/SPIRV/Target.cpp

@@ -34,7 +34,7 @@ class SPIRVTargetAttrImpl
   serializeToObject(Attribute attribute, Operation *module,
                     const gpu::TargetOptions &options) const;
 
-  Attribute createObject(Attribute attribute,
+  Attribute createObject(Attribute attribute, Operation *module,
                          const SmallVector<char, 0> &object,
                          const gpu::TargetOptions &options) const;
 };
@@ -89,7 +89,7 @@ std::optional<SmallVector<char, 0>> SPIRVTargetAttrImpl::serializeToObject(
 
 // Prepare Attribute for gpu.binary with serialized kernel object
 Attribute
-SPIRVTargetAttrImpl::createObject(Attribute attribute,
+SPIRVTargetAttrImpl::createObject(Attribute attribute, Operation *module,
                                   const SmallVector<char, 0> &object,
                                   const gpu::TargetOptions &options) const {
   gpu::CompilationTarget format = options.getCompilationTarget();