[CIR][AMDGPU] Add module flags for AMDGPU target using amendOperation of CIRDialectLLVMIRTranslationInterface (#186073)

Add the amendOperation override to handle CIR dialect attributes during
MLIR-to-LLVM IR translation. This dispatches to  amendModule for ModuleOp,
enabling module metadata.

This PR also adds support to emit AMDGPU-specific module flags
amdhsa_code_object_version and amdgpu_printf_kind to match OGCG
behavior.

In CIRGenModule, the flags are stored as CIR module attributes:

cir.amdhsa_code_object_version (integer)
cir.amdgpu_printf_kind (string: "hostcall" or "buffered")
During lowering to LLVM IR (in LowerToLLVMIR.cpp), these attributes are
converted to LLVM module flags.

Upstreaming basic changes from clangIR PRs: 

61e9ebd9f8
https://github.com/llvm/clangir/pull/768
https://github.com/llvm/clangir/pull/773
https://github.com/llvm/clangir/pull/2100

clang/include/clang/CIR/Dialect/IR/CIRDialect.td

@@ -76,6 +76,9 @@ def CIR_Dialect : Dialect {
     static llvm::StringRef getResAttrsAttrName() { return "res_attrs"; }
     static llvm::StringRef getArgAttrsAttrName() { return "arg_attrs"; }
 
+    static llvm::StringRef getAMDGPUCodeObjectVersionAttrName() { return "cir.amdhsa_code_object_version"; }
+    static llvm::StringRef getAMDGPUPrintfKindAttrName() { return "cir.amdgpu_printf_kind"; }
+
     void registerAttributes();
     void registerTypes();
 

clang/lib/CIR/CodeGen/CIRGenAMDGPU.cpp

@@ -0,0 +1,41 @@
+//===- CIRGenAMDGPU.cpp - AMDGPU-specific logic for CIR generation --------===//
+//
+// 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 contains code dealing with AMDGPU-specific logic of CIR generation.
+//
+//===----------------------------------------------------------------------===//
+
+#include "CIRGenModule.h"
+
+#include "clang/Basic/TargetOptions.h"
+#include "clang/CIR/Dialect/IR/CIRDialect.h"
+#include "llvm/TargetParser/Triple.h"
+
+using namespace clang;
+using namespace clang::CIRGen;
+
+void CIRGenModule::emitAMDGPUMetadata() {
+  // Emit code object version module flag.
+  if (target.getTargetOpts().CodeObjectVersion !=
+      llvm::CodeObjectVersionKind::COV_None) {
+    theModule->setAttr(
+        cir::CIRDialect::getAMDGPUCodeObjectVersionAttrName(),
+        builder.getI32IntegerAttr(target.getTargetOpts().CodeObjectVersion));
+  }
+
+  // Emit printf kind module flag for HIP.
+  if (langOpts.HIP) {
+    llvm::StringRef printfKind =
+        target.getTargetOpts().AMDGPUPrintfKindVal ==
+                TargetOptions::AMDGPUPrintfKind::Hostcall
+            ? "hostcall"
+            : "buffered";
+    theModule->setAttr(cir::CIRDialect::getAMDGPUPrintfKindAttrName(),
+                       builder.getStringAttr(printfKind));
+  }
+}

clang/lib/CIR/CodeGen/CIRGenModule.cpp

@@ -3008,6 +3008,10 @@ void CIRGenModule::release() {
   theModule->setAttr(cir::CIRDialect::getModuleLevelAsmAttrName(),
                      builder.getArrayAttr(globalScopeAsm));
 
+  if (getTriple().isAMDGPU() ||
+      (getTriple().isSPIRV() && getTriple().getVendor() == llvm::Triple::AMD))
+    emitAMDGPUMetadata();
+
   // There's a lot of code that is not implemented yet.
   assert(!cir::MissingFeatures::cgmRelease());
 }

clang/lib/CIR/CodeGen/CIRGenModule.h

@@ -781,6 +781,9 @@ public:
   /// Print out an error that codegen doesn't support the specified decl yet.
   void errorUnsupported(const Decl *d, llvm::StringRef type);
 
+  /// Emits AMDGPU specific Metadata.
+  void emitAMDGPUMetadata();
+
 private:
   // An ordered map of canonical GlobalDecls to their mangled names.
   llvm::MapVector<clang::GlobalDecl, llvm::StringRef> mangledDeclNames;

clang/lib/CIR/CodeGen/CMakeLists.txt

@@ -14,6 +14,7 @@ add_clang_library(clangCIR
   CIRGenBuiltin.cpp
   CIRGenBuiltinAArch64.cpp
   CIRGenBuiltinAMDGPU.cpp
+  CIRGenAMDGPU.cpp
   CIRGenBuiltinX86.cpp
   CIRGenCall.cpp
   CIRGenClass.cpp

clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp

@@ -47,6 +47,49 @@ public:
 
     return mlir::success();
   }
+
+  /// Any named attribute in the CIR dialect, i.e, with name started with
+  /// "cir.", will be handled here.
+  virtual mlir::LogicalResult amendOperation(
+      mlir::Operation *op, llvm::ArrayRef<llvm::Instruction *> instructions,
+      mlir::NamedAttribute attribute,
+      mlir::LLVM::ModuleTranslation &moduleTranslation) const override {
+    if (auto mod = dyn_cast<mlir::ModuleOp>(op)) {
+      if (mlir::failed(amendModule(mod, attribute, moduleTranslation)))
+        return mlir::failure();
+    }
+    return mlir::success();
+  }
+
+private:
+  // Translate CIR's module attributes to LLVM's module metadata
+  mlir::LogicalResult
+  amendModule(mlir::ModuleOp mod, mlir::NamedAttribute attribute,
+              mlir::LLVM::ModuleTranslation &moduleTranslation) const {
+    llvm::Module *llvmModule = moduleTranslation.getLLVMModule();
+    llvm::LLVMContext &llvmContext = llvmModule->getContext();
+
+    if (attribute.getName() == "cir.amdhsa_code_object_version") {
+      if (auto intAttr =
+              mlir::dyn_cast<mlir::IntegerAttr>(attribute.getValue())) {
+        llvmModule->addModuleFlag(llvm::Module::Error,
+                                  "amdhsa_code_object_version",
+                                  static_cast<uint32_t>(intAttr.getInt()));
+      }
+    }
+
+    if (attribute.getName() == "cir.amdgpu_printf_kind") {
+      if (auto strAttr =
+              mlir::dyn_cast<mlir::StringAttr>(attribute.getValue())) {
+        llvm::MDString *mdStr =
+            llvm::MDString::get(llvmContext, strAttr.getValue());
+        llvmModule->addModuleFlag(llvm::Module::Error, "amdgpu_printf_kind",
+                                  mdStr);
+      }
+    }
+
+    return mlir::success();
+  }
 };
 
 void registerCIRDialectTranslation(mlir::DialectRegistry &registry) {

clang/test/CIR/CodeGenHIP/amdgpu-module-flags.hip

@@ -0,0 +1,26 @@
+#include "../CodeGenCUDA/Inputs/cuda.h"
+
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip -fclangir \
+// RUN:            -fcuda-is-device -emit-cir %s -o %t.cir
+// RUN: FileCheck --check-prefix=CIR %s --input-file=%t.cir
+
+// RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip -fclangir \
+// RUN:            -fcuda-is-device -emit-llvm %s -o %t.cir.ll
+// RUN: FileCheck --check-prefix=LLVM %s --input-file=%t.cir.ll
+
+// RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip \
+// RUN:            -fcuda-is-device -emit-llvm %s -o %t.ogcg.ll
+// RUN: FileCheck --check-prefix=LLVM %s --input-file=%t.ogcg.ll
+
+// Test that AMDGPU module flags are emitted correctly.
+
+// CIR: module {{.*}} attributes {
+// CIR-SAME: cir.amdgpu_printf_kind = "hostcall"
+// CIR-SAME: cir.amdhsa_code_object_version = 600
+
+// LLVM: !llvm.module.flags = !{
+// LLVM-DAG: !{i32 1, !"amdhsa_code_object_version", i32 600}
+// LLVM-DAG: !{i32 1, !"amdgpu_printf_kind", !"hostcall"}
+
+__global__ void kernel() {}