https://github.com/skc7 updated https://github.com/llvm/llvm-project/pull/186073
>From ef38d27602a43ffe9ef38065c9815eb8d6d70998 Mon Sep 17 00:00:00 2001 From: skc7 <[email protected]> Date: Thu, 12 Mar 2026 15:28:39 +0530 Subject: [PATCH 1/5] [CIR][NFC] Add amendOperation to CIRDialectLLVMIRTranslationInterface --- .../Lowering/DirectToLLVM/LowerToLLVMIR.cpp | 29 +++++++++++++++++++ 1 file changed, 29 insertions(+) diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp index 30b9eaaca2d37..8de63bfb169a6 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp +++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp @@ -47,6 +47,35 @@ class CIRDialectLLVMIRTranslationInterface 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 func = dyn_cast<mlir::LLVM::LLVMFuncOp>(op)) { + amendFunction(func, instructions, attribute, moduleTranslation); + } else if (auto mod = dyn_cast<mlir::ModuleOp>(op)) { + amendModule(mod, attribute, moduleTranslation); + } + return mlir::success(); + } + +private: + // Translate CIR's extra function attributes to LLVM's function attributes. + void amendFunction(mlir::LLVM::LLVMFuncOp func, + llvm::ArrayRef<llvm::Instruction *> instructions, + mlir::NamedAttribute attribute, + mlir::LLVM::ModuleTranslation &moduleTranslation) const { + // TODO(cir): Implement this + } + + // Translate CIR's module attributes to LLVM's module metadata + void amendModule(mlir::ModuleOp mod, mlir::NamedAttribute attribute, + mlir::LLVM::ModuleTranslation &moduleTranslation) const { + // TODO(cir): Implement this + } }; void registerCIRDialectTranslation(mlir::DialectRegistry ®istry) { >From 863e058e6e5d7790cee7a8f2de2647eac763c889 Mon Sep 17 00:00:00 2001 From: Chaitanya <[email protected]> Date: Fri, 13 Mar 2026 10:14:53 +0530 Subject: [PATCH 2/5] [CIR][AMDGPU] Add module flags for AMDGPU target (#186081) Upstreaming clangIR PR: https://github.com/llvm/clangir/pull/2100 This PR 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. --- .../clang/CIR/Dialect/IR/CIRDialect.td | 3 ++ clang/lib/CIR/CodeGen/CIRGenAMDGPU.cpp | 41 +++++++++++++++++++ clang/lib/CIR/CodeGen/CIRGenModule.cpp | 3 ++ clang/lib/CIR/CodeGen/CIRGenModule.h | 3 ++ clang/lib/CIR/CodeGen/CMakeLists.txt | 1 + .../Lowering/DirectToLLVM/LowerToLLVMIR.cpp | 23 ++++++++++- .../CIR/CodeGenHIP/amdgpu-module-flags.hip | 30 ++++++++++++++ 7 files changed, 103 insertions(+), 1 deletion(-) create mode 100644 clang/lib/CIR/CodeGen/CIRGenAMDGPU.cpp create mode 100644 clang/test/CIR/CodeGenHIP/amdgpu-module-flags.hip diff --git a/clang/include/clang/CIR/Dialect/IR/CIRDialect.td b/clang/include/clang/CIR/Dialect/IR/CIRDialect.td index 3056179f08264..f1f94c868e5b0 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIRDialect.td +++ b/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(); diff --git a/clang/lib/CIR/CodeGen/CIRGenAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenAMDGPU.cpp new file mode 100644 index 0000000000000..896e74e548c61 --- /dev/null +++ b/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)); + } +} diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp index cb931f969a41d..fd08cdae37881 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp @@ -2878,6 +2878,9 @@ void CIRGenModule::release() { theModule->setAttr(cir::CIRDialect::getModuleLevelAsmAttrName(), builder.getArrayAttr(globalScopeAsm)); + if (getTriple().isAMDGPU()) + emitAMDGPUMetadata(); + // There's a lot of code that is not implemented yet. assert(!cir::MissingFeatures::cgmRelease()); } diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.h b/clang/lib/CIR/CodeGen/CIRGenModule.h index bef154955b9b6..baaf7db20dd31 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.h +++ b/clang/lib/CIR/CodeGen/CIRGenModule.h @@ -778,6 +778,9 @@ class CIRGenModule : public CIRGenTypeCache { /// 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; diff --git a/clang/lib/CIR/CodeGen/CMakeLists.txt b/clang/lib/CIR/CodeGen/CMakeLists.txt index f982fcf5b1b8a..8548cc8424527 100644 --- a/clang/lib/CIR/CodeGen/CMakeLists.txt +++ b/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 diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp index 8de63bfb169a6..7daeb88ec0900 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp +++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp @@ -74,7 +74,28 @@ class CIRDialectLLVMIRTranslationInterface // Translate CIR's module attributes to LLVM's module metadata void amendModule(mlir::ModuleOp mod, mlir::NamedAttribute attribute, mlir::LLVM::ModuleTranslation &moduleTranslation) const { - // TODO(cir): Implement this + llvm::Module *llvmModule = moduleTranslation.getLLVMModule(); + llvm::LLVMContext &llvmContext = llvmModule->getContext(); + + // AMDGPU module flags + 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); + } + } } }; diff --git a/clang/test/CIR/CodeGenHIP/amdgpu-module-flags.hip b/clang/test/CIR/CodeGenHIP/amdgpu-module-flags.hip new file mode 100644 index 0000000000000..5d1f48291658c --- /dev/null +++ b/clang/test/CIR/CodeGenHIP/amdgpu-module-flags.hip @@ -0,0 +1,30 @@ +#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=OGCG %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"} + +// OGCG: !llvm.module.flags = !{ +// OGCG-DAG: !{i32 1, !"amdhsa_code_object_version", i32 600} +// OGCG-DAG: !{i32 1, !"amdgpu_printf_kind", !"hostcall"} + +__global__ void kernel() {} >From 9d4302927fca3fec5ec9ac7c5aee91263e5bd785 Mon Sep 17 00:00:00 2001 From: skc7 <[email protected]> Date: Wed, 18 Mar 2026 10:53:15 +0530 Subject: [PATCH 3/5] Fix amendFunction and amendModule returns --- .../Lowering/DirectToLLVM/LowerToLLVMIR.cpp | 29 ++++++++++++------- .../CIR/CodeGenHIP/amdgpu-module-flags.hip | 6 +--- 2 files changed, 20 insertions(+), 15 deletions(-) diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp index 7daeb88ec0900..88aeea33fd2c4 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp +++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp @@ -55,35 +55,41 @@ class CIRDialectLLVMIRTranslationInterface mlir::NamedAttribute attribute, mlir::LLVM::ModuleTranslation &moduleTranslation) const override { if (auto func = dyn_cast<mlir::LLVM::LLVMFuncOp>(op)) { - amendFunction(func, instructions, attribute, moduleTranslation); + if (mlir::failed( + amendFunction(func, instructions, attribute, moduleTranslation))) + return mlir::failure(); } else if (auto mod = dyn_cast<mlir::ModuleOp>(op)) { - amendModule(mod, attribute, moduleTranslation); + if (mlir::failed(amendModule(mod, attribute, moduleTranslation))) + return mlir::failure(); } return mlir::success(); } private: // Translate CIR's extra function attributes to LLVM's function attributes. - void amendFunction(mlir::LLVM::LLVMFuncOp func, - llvm::ArrayRef<llvm::Instruction *> instructions, - mlir::NamedAttribute attribute, - mlir::LLVM::ModuleTranslation &moduleTranslation) const { - // TODO(cir): Implement this + mlir::LogicalResult + amendFunction(mlir::LLVM::LLVMFuncOp func, + llvm::ArrayRef<llvm::Instruction *> instructions, + mlir::NamedAttribute attribute, + mlir::LLVM::ModuleTranslation &moduleTranslation) const { + // TODO(CIR): process extra function attributes. + return mlir::success(); } // Translate CIR's module attributes to LLVM's module metadata - void amendModule(mlir::ModuleOp mod, mlir::NamedAttribute attribute, - mlir::LLVM::ModuleTranslation &moduleTranslation) const { + mlir::LogicalResult + amendModule(mlir::ModuleOp mod, mlir::NamedAttribute attribute, + mlir::LLVM::ModuleTranslation &moduleTranslation) const { llvm::Module *llvmModule = moduleTranslation.getLLVMModule(); llvm::LLVMContext &llvmContext = llvmModule->getContext(); - // AMDGPU module flags 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())); + return mlir::success(); } } @@ -94,8 +100,11 @@ class CIRDialectLLVMIRTranslationInterface llvm::MDString::get(llvmContext, strAttr.getValue()); llvmModule->addModuleFlag(llvm::Module::Error, "amdgpu_printf_kind", mdStr); + return mlir::success(); } } + + return mlir::success(); } }; diff --git a/clang/test/CIR/CodeGenHIP/amdgpu-module-flags.hip b/clang/test/CIR/CodeGenHIP/amdgpu-module-flags.hip index 5d1f48291658c..7a597bca64bb9 100644 --- a/clang/test/CIR/CodeGenHIP/amdgpu-module-flags.hip +++ b/clang/test/CIR/CodeGenHIP/amdgpu-module-flags.hip @@ -11,7 +11,7 @@ // RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip \ // RUN: -fcuda-is-device -emit-llvm %s -o %t.ogcg.ll -// RUN: FileCheck --check-prefix=OGCG %s --input-file=%t.ogcg.ll +// RUN: FileCheck --check-prefix=LLVM %s --input-file=%t.ogcg.ll // Test that AMDGPU module flags are emitted correctly. @@ -23,8 +23,4 @@ // LLVM-DAG: !{i32 1, !"amdhsa_code_object_version", i32 600} // LLVM-DAG: !{i32 1, !"amdgpu_printf_kind", !"hostcall"} -// OGCG: !llvm.module.flags = !{ -// OGCG-DAG: !{i32 1, !"amdhsa_code_object_version", i32 600} -// OGCG-DAG: !{i32 1, !"amdgpu_printf_kind", !"hostcall"} - __global__ void kernel() {} >From 72a827368b658df9d652cdd2892d137b748da6ca Mon Sep 17 00:00:00 2001 From: skc7 <[email protected]> Date: Wed, 18 Mar 2026 12:04:57 +0530 Subject: [PATCH 4/5] remove redundant returns --- clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp | 2 -- 1 file changed, 2 deletions(-) diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp index 88aeea33fd2c4..3fc13168e742e 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp +++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp @@ -89,7 +89,6 @@ class CIRDialectLLVMIRTranslationInterface llvmModule->addModuleFlag(llvm::Module::Error, "amdhsa_code_object_version", static_cast<uint32_t>(intAttr.getInt())); - return mlir::success(); } } @@ -100,7 +99,6 @@ class CIRDialectLLVMIRTranslationInterface llvm::MDString::get(llvmContext, strAttr.getValue()); llvmModule->addModuleFlag(llvm::Module::Error, "amdgpu_printf_kind", mdStr); - return mlir::success(); } } >From cb012328e3c94d2e028c3d212d801f56c8c98f2f Mon Sep 17 00:00:00 2001 From: skc7 <[email protected]> Date: Fri, 20 Mar 2026 11:19:23 +0530 Subject: [PATCH 5/5] remove amendFunction --- clang/lib/CIR/CodeGen/CIRGenModule.cpp | 3 ++- .../CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp | 16 +--------------- 2 files changed, 3 insertions(+), 16 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp index fd08cdae37881..d55788ffa2f2a 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp @@ -2878,7 +2878,8 @@ void CIRGenModule::release() { theModule->setAttr(cir::CIRDialect::getModuleLevelAsmAttrName(), builder.getArrayAttr(globalScopeAsm)); - if (getTriple().isAMDGPU()) + if (getTriple().isAMDGPU() || + (getTriple().isSPIRV() && getTriple().getVendor() == llvm::Triple::AMD)) emitAMDGPUMetadata(); // There's a lot of code that is not implemented yet. diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp index 3fc13168e742e..2a95cfb9371b1 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp +++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp @@ -54,11 +54,7 @@ class CIRDialectLLVMIRTranslationInterface mlir::Operation *op, llvm::ArrayRef<llvm::Instruction *> instructions, mlir::NamedAttribute attribute, mlir::LLVM::ModuleTranslation &moduleTranslation) const override { - if (auto func = dyn_cast<mlir::LLVM::LLVMFuncOp>(op)) { - if (mlir::failed( - amendFunction(func, instructions, attribute, moduleTranslation))) - return mlir::failure(); - } else if (auto mod = dyn_cast<mlir::ModuleOp>(op)) { + if (auto mod = dyn_cast<mlir::ModuleOp>(op)) { if (mlir::failed(amendModule(mod, attribute, moduleTranslation))) return mlir::failure(); } @@ -66,16 +62,6 @@ class CIRDialectLLVMIRTranslationInterface } private: - // Translate CIR's extra function attributes to LLVM's function attributes. - mlir::LogicalResult - amendFunction(mlir::LLVM::LLVMFuncOp func, - llvm::ArrayRef<llvm::Instruction *> instructions, - mlir::NamedAttribute attribute, - mlir::LLVM::ModuleTranslation &moduleTranslation) const { - // TODO(CIR): process extra function attributes. - return mlir::success(); - } - // Translate CIR's module attributes to LLVM's module metadata mlir::LogicalResult amendModule(mlir::ModuleOp mod, mlir::NamedAttribute attribute, _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
