[clang] [CIR][AMDGPU] Add module flags for AMDGPU target using amendOperation of CIRDialectLLVMIRTranslationInterface (PR #186073)
via cfe-commits
cfe-commits at lists.llvm.org
Thu Mar 19 22:50:30 PDT 2026
https://github.com/skc7 updated https://github.com/llvm/llvm-project/pull/186073
>From ef38d27602a43ffe9ef38065c9815eb8d6d70998 Mon Sep 17 00:00:00 2001
From: skc7 <Krishna.Sankisa at amd.com>
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 <Krishna.Sankisa at amd.com>
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 <Krishna.Sankisa at amd.com>
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 <Krishna.Sankisa at amd.com>
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 <Krishna.Sankisa at amd.com>
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,
More information about the cfe-commits
mailing list