[mlir] Default for gpu-binary-annotation option.

[csigg]

Summary

Provide default for gpuBinaryAnnotation so that we don't need to specify it in tests.

The annotation likely only needs to be target specific if we want to lower to e.g. both CUDA and ROCDL.

[csigg]

Move option to GpuToLLVMConversionPass (from Base) to remove the duplicate forward declaration of getDefaultGpuBinaryAnnotation().

[csigg]

Remove triple, target and feature options.

[csigg]

Fix.

[bondhugula]

A couple of minor comments.

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

30

Nit: Triple /// comments please consistent with other places below.

50

It'd be good to document gpuBinaryAnnotation here - that a default annotation is used if an empty string is provided?

[bondhugula]

In D98168#2613946, @csigg wrote:

Remove triple, target and feature options.

Could you include this in the commit summary please?

[csigg]

Address reviewer comments.

[csigg]

In D98168#2614317, @bondhugula wrote:

In D98168#2613946, @csigg wrote:

Remove triple, target and feature options.

Could you include this in the commit summary please?

I reverted part of a previous version. This change does not remove any options.

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

 class Operation;
 class OwningRewritePatternList;
 
 template <typename T>
 class OperationPass;
 
 namespace gpu {
 class GPUModuleOp;
+
+std::string getDefaultGpuBinaryAnnotation();
 } // namespace gpu
 
 namespace LLVM {
 class LLVMDialect;
 } // 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 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 = "");
+createGpuToLLVMConversionPass(StringRef gpuBinaryAnnotation = {});
 
 /// Collect a set of patterns to convert from the GPU dialect to LLVM.
 void populateGpuToLLVMConversionPatterns(LLVMTypeConverter &converter,
                                          OwningRewritePatternList &patterns,
-                                         StringRef gpuBinaryAnnotation);
+                                         StringRef gpuBinaryAnnotation = {});
 
 /// Creates a pass to convert kernel functions into GPU target object blobs.
 ///
 /// This transformation takes the body of each function that is annotated with
 /// the 'gpu.kernel' attribute, copies it to a new LLVM module, compiles the
 /// module with help of the GPU backend to target object and then invokes
 /// the provided blobGenerator to produce a binary blob. Such blob is then
 /// attached as a string attribute to the kernel function.

 /// - gpuBinaryAnnotation : the name of the attribute which contains the blob.
 ///
 /// After the transformation, the body of the kernel function is removed (i.e.,
 /// it is turned into a declaration).
 std::unique_ptr<OperationPass<gpu::GPUModuleOp>>
 createConvertGPUKernelToBlobPass(LoweringCallback loweringCallback,
                                  BlobGenerator blobGenerator, StringRef triple,
                                  StringRef targetChip, StringRef features,
-                                 StringRef gpuBinaryAnnotation);
+                                 StringRef gpuBinaryAnnotation = {});
 
 } // namespace mlir
 
 #endif // MLIR_CONVERSION_GPUCOMMON_GPUCOMMONPASS_H_

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 dependentDialects = ["LLVM::LLVMDialect"];
   let options = [
     Option<"gpuBinaryAnnotation", "gpu-binary-annotation", "std::string",
-           "", "Annotation attribute string for GPU binary">,
+           /*default=*/"mlir::gpu::getDefaultGpuBinaryAnnotation()",
+           "Annotation attribute string for GPU binary">,
   ];
 }
 
 def LowerHostCodeToLLVM : Pass<"lower-host-to-llvm", "ModuleOp"> {
   let summary = "Lowers the host module code and `gpu.launch_func` to LLVM";
   let constructor = "mlir::createLowerHostCodeToLLVMPass()";
   let dependentDialects = ["LLVM::LLVMDialect"];
 }

mlir/lib/Conversion/GPUCommon/ConvertKernelFuncToBlob.cpp

 /// The function body is erased.
 class GpuKernelToBlobPass
     : public PassWrapper<GpuKernelToBlobPass, OperationPass<gpu::GPUModuleOp>> {
 public:
   GpuKernelToBlobPass(LoweringCallback loweringCallback,
                       BlobGenerator blobGenerator, StringRef triple,
                       StringRef targetChip, StringRef features,
                       StringRef gpuBinaryAnnotation)
-      : loweringCallback(loweringCallback), blobGenerator(blobGenerator),
-        triple(triple), targetChip(targetChip), features(features),
-        blobAnnotation(gpuBinaryAnnotation) {}
+      : loweringCallback(loweringCallback), blobGenerator(blobGenerator) {
+    if (!triple.empty())
+      this->triple = llvm::Triple(triple);
+    if (!targetChip.empty())
+      this->targetChip = targetChip.str();
+    if (!features.empty())
+      this->features = features.str();
+    if (!gpuBinaryAnnotation.empty())
+      this->gpuBinaryAnnotation = gpuBinaryAnnotation.str();
+  }
+
+  GpuKernelToBlobPass(const GpuKernelToBlobPass &other)
+      : loweringCallback(other.loweringCallback),
+        blobGenerator(other.blobGenerator) {}
 
   void runOnOperation() override {
     gpu::GPUModuleOp module = getOperation();
 
     // Lower the module to an LLVM IR module using a separate context to enable
     // multi-threaded processing.
     llvm::LLVMContext llvmContext;
     std::unique_ptr<llvm::Module> llvmModule =
         loweringCallback(module, llvmContext, "LLVMDialectModule");
     if (!llvmModule)
       return signalPassFailure();
 
     // Translate the llvm module to a target blob and attach the result as
     // attribute to the module.
     if (auto blobAttr = translateGPUModuleToBinaryAnnotation(
             *llvmModule, module.getLoc(), module.getName()))
-      module->setAttr(blobAnnotation, blobAttr);
+      module->setAttr(gpuBinaryAnnotation, blobAttr);
     else
       signalPassFailure();
   }
 
 private:
   std::string translateModuleToISA(llvm::Module &module,
                                    llvm::TargetMachine &targetMachine);
 
   /// Converts llvmModule to a blob with target instructions using the
   /// user-provided generator. Location is used for error reporting and name is
   /// forwarded to the blob generator to use in its logging mechanisms.
   OwnedBlob convertModuleToBlob(llvm::Module &llvmModule, Location loc,
                                 StringRef name);
 
   /// Translates llvmModule to a blob with target instructions and returns the
   /// result as attribute.
   StringAttr translateGPUModuleToBinaryAnnotation(llvm::Module &llvmModule,
                                                   Location loc, StringRef name);
 
   LoweringCallback loweringCallback;
   BlobGenerator blobGenerator;
-  llvm::Triple triple;
-  std::string targetChip;
-  std::string features;
-  std::string blobAnnotation;
+
+  Option<llvm::Triple> triple{*this, "triple",
+                              ::llvm::cl::desc("Target triple")};
+
+  Option<std::string> targetChip{*this, "target-chip",
+                                 ::llvm::cl::desc("Target architecture")};
+
+  Option<std::string> features{*this, "features",
+                               ::llvm::cl::desc("Target features")};
+
+  Option<std::string> gpuBinaryAnnotation{
+      *this, "gpu-binary-annotation",
+      ::llvm::cl::desc("Annotation attribute string for GPU binary"),
+      ::llvm::cl::init(gpu::getDefaultGpuBinaryAnnotation())};
 };
 
 } // anonymous namespace
 
+std::string gpu::getDefaultGpuBinaryAnnotation() { return "gpu.binary"; }
+
 std::string
 GpuKernelToBlobPass::translateModuleToISA(llvm::Module &module,
                                           llvm::TargetMachine &targetMachine) {
   std::string targetISA;
   {
     llvm::raw_string_ostream stream(targetISA);
     llvm::buffer_ostream pstream(stream);
     llvm::legacy::PassManager codegenPasses;

mlir/lib/Conversion/PassDetail.h

 
 namespace complex {
 class ComplexDialect;
 } // end namespace complex
 
 namespace gpu {
 class GPUDialect;
 class GPUModuleOp;
+
+std::string getDefaultGpuBinaryAnnotation();
 } // end namespace gpu
 
 namespace LLVM {
 class LLVMArmNeonDialect;
 class LLVMArmSVEDialect;
 class LLVMAVX512Dialect;
 class LLVMDialect;
 } // end namespace LLVM

mlir/tools/mlir-cuda-runner/mlir-cuda-runner.cpp

OwnedBlob compilePtxToCubin(const std::string ptx, Location loc,
 
   return result;
 }
 
 struct GpuToCubinPipelineOptions
     : public mlir::PassPipelineOptions<GpuToCubinPipelineOptions> {
   Option<std::string> gpuBinaryAnnotation{
       *this, "gpu-binary-annotation",
-      llvm::cl::desc("Annotation attribute string for GPU binary")};
+      llvm::cl::desc("Annotation attribute string for GPU binary"),
+      llvm::cl::init(gpu::getDefaultGpuBinaryAnnotation())};
 };
 
 // Register cuda-runner specific passes.
 static void registerCudaRunnerPasses() {
   PassPipelineRegistration<GpuToCubinPipelineOptions> registerGpuToCubin(
       "gpu-to-cubin", "Generate CUBIN from gpu.launch regions",
       [&](OpPassManager &pm, const GpuToCubinPipelineOptions &options) {
         pm.addPass(createGpuKernelOutliningPass());