[clang] [mlir] [CIR][OpenMP] Add host op filtering pass to CIR pipeline (PR #209592)
Jan Leyonberg via cfe-commits
cfe-commits at lists.llvm.org
Wed Jul 15 04:15:46 PDT 2026
https://github.com/jsjodin updated https://github.com/llvm/llvm-project/pull/209592
>From d6d61356abf6768b91fb1c4fe5094417d1dd39b9 Mon Sep 17 00:00:00 2001
From: Jan Leyonberg <jan_sjodin at yahoo.com>
Date: Tue, 14 Jul 2026 14:54:30 -0400
Subject: [PATCH 1/2] [CIR][OpenMP] Add host op filtering pass to CIR pipeline
This patch adds the host op filtering pass which prevents host code being
lowered when compiling for the target device.
---
.../CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp | 1 +
.../CodeGenOpenMP/target-host-op-filtering.c | 35 +++++++++++++++++++
2 files changed, 36 insertions(+)
create mode 100644 clang/test/CIR/CodeGenOpenMP/target-host-op-filtering.c
diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
index 0e412090a16da..67406d261ce1a 100644
--- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
+++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
@@ -5134,6 +5134,7 @@ void populateCIRToLLVMPasses(mlir::OpPassManager &pm) {
mlir::populateCIRPreLoweringPasses(pm);
pm.addPass(mlir::omp::createMarkDeclareTargetPass());
pm.addPass(createConvertCIRToLLVMPass());
+ pm.addPass(mlir::omp::createHostOpFilteringPass());
}
std::unique_ptr<llvm::Module>
diff --git a/clang/test/CIR/CodeGenOpenMP/target-host-op-filtering.c b/clang/test/CIR/CodeGenOpenMP/target-host-op-filtering.c
new file mode 100644
index 0000000000000..ab20b285c79ed
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenMP/target-host-op-filtering.c
@@ -0,0 +1,35 @@
+// Verifies the omp-host-op-filtering pass runs in the CIR-to-LLVM device
+// pipeline.//
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa \
+// RUN: -fclangir -emit-llvm-bc %s -o %t-cir-host.bc
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fopenmp -fopenmp-is-target-device \
+// RUN: -fopenmp-host-ir-file-path %t-cir-host.bc \
+// RUN: -fclangir -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefix=LLVM
+
+void use(int);
+int host_helper(int);
+
+void f(int x) {
+ // Host-only computation that must not leak into the device module.
+ int y = host_helper(x);
+ use(y);
+#pragma omp target map(tofrom : x)
+ {
+ x = x + 1;
+ }
+}
+
+// LLVM-LABEL: define hidden void @f(
+// LLVM-NOT: call i32 @host_helper
+// LLVM-NOT: call void @use
+// LLVM: ret void
+
+// LLVM-LABEL: define weak_odr protected amdgpu_kernel void @__omp_offloading_{{.*}}_f_l
+// LLVM: call i32 @__kmpc_target_init(
+// LLVM: user_code.entry:
+// LLVM: %[[LD:.*]] = load i32, ptr %{{.*}}, align 4
+// LLVM: %[[ADD:.*]] = add nsw i32 %[[LD]], 1
+// LLVM: store i32 %[[ADD]], ptr %{{.*}}, align 4
+// LLVM: call void @__kmpc_target_deinit()
+// LLVM: ret void
>From 9ce9d47344b8b493b0c6077b692cb7eaf967f2c6 Mon Sep 17 00:00:00 2001
From: Jan Leyonberg <jan_sjodin at yahoo.com>
Date: Wed, 15 Jul 2026 07:13:31 -0400
Subject: [PATCH 2/2] Add guard for adding OpenMP passes.
---
clang/include/clang/CIR/LowerToLLVM.h | 2 +-
clang/include/clang/CIR/Passes.h | 5 +++--
clang/lib/CIR/FrontendAction/CIRGenAction.cpp | 9 +++++----
clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp | 12 +++++++-----
clang/tools/cir-opt/cir-opt.cpp | 9 +++++++--
clang/tools/cir-translate/cir-translate.cpp | 7 ++++++-
.../mlir/Dialect/OpenMP/OpenMPOpsInterfaces.td | 9 +++++++++
7 files changed, 38 insertions(+), 15 deletions(-)
diff --git a/clang/include/clang/CIR/LowerToLLVM.h b/clang/include/clang/CIR/LowerToLLVM.h
index df5f9221842ea..305f54017ebe0 100644
--- a/clang/include/clang/CIR/LowerToLLVM.h
+++ b/clang/include/clang/CIR/LowerToLLVM.h
@@ -32,7 +32,7 @@ namespace cir {
namespace direct {
std::unique_ptr<llvm::Module>
lowerDirectlyFromCIRToLLVMIR(mlir::ModuleOp mlirModule,
- llvm::LLVMContext &llvmCtx,
+ llvm::LLVMContext &llvmCtx, bool enableOpenMP,
llvm::StringRef mlirSaveTempsOutFile = {},
llvm::vfs::FileSystem *fs = nullptr);
} // namespace direct
diff --git a/clang/include/clang/CIR/Passes.h b/clang/include/clang/CIR/Passes.h
index 1d202ce5bfb52..a42f20405af24 100644
--- a/clang/include/clang/CIR/Passes.h
+++ b/clang/include/clang/CIR/Passes.h
@@ -22,8 +22,9 @@ namespace direct {
/// Create a pass that fully lowers CIR to the LLVMIR dialect.
std::unique_ptr<mlir::Pass> createConvertCIRToLLVMPass();
-/// Adds passes that fully lower CIR to the LLVMIR dialect.
-void populateCIRToLLVMPasses(mlir::OpPassManager &pm);
+/// Adds passes that fully lower CIR to the LLVMIR dialect. When enableOpenMP
+/// is set (-fopenmp), the OpenMP lowering passes are also added.
+void populateCIRToLLVMPasses(mlir::OpPassManager &pm, bool enableOpenMP);
} // namespace direct
} // end namespace cir
diff --git a/clang/lib/CIR/FrontendAction/CIRGenAction.cpp b/clang/lib/CIR/FrontendAction/CIRGenAction.cpp
index 92c9ef5499693..28229ac8c58fd 100644
--- a/clang/lib/CIR/FrontendAction/CIRGenAction.cpp
+++ b/clang/lib/CIR/FrontendAction/CIRGenAction.cpp
@@ -57,9 +57,10 @@ getBackendActionFromOutputType(CIRGenAction::OutputType Action) {
static std::unique_ptr<llvm::Module>
lowerFromCIRToLLVMIR(mlir::ModuleOp MLIRModule, llvm::LLVMContext &LLVMCtx,
+ bool EnableOpenMP,
llvm::StringRef mlirSaveTempsOutFile = {},
llvm::vfs::FileSystem *fs = nullptr) {
- return direct::lowerDirectlyFromCIRToLLVMIR(MLIRModule, LLVMCtx,
+ return direct::lowerDirectlyFromCIRToLLVMIR(MLIRModule, LLVMCtx, EnableOpenMP,
mlirSaveTempsOutFile, fs);
}
@@ -179,9 +180,9 @@ class CIRGenConsumer : public clang::ASTConsumer {
MlirModule->print(out);
}
- std::unique_ptr<llvm::Module> LLVMModule =
- lowerFromCIRToLLVMIR(MlirModule, LLVMCtx, mlirSaveTempsOutFile,
- &CI.getVirtualFileSystem());
+ std::unique_ptr<llvm::Module> LLVMModule = lowerFromCIRToLLVMIR(
+ MlirModule, LLVMCtx, C.getLangOpts().OpenMP, mlirSaveTempsOutFile,
+ &CI.getVirtualFileSystem());
if (linkInModules(*LLVMModule))
return;
diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
index 67406d261ce1a..ebe74399dda62 100644
--- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
+++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
@@ -5130,23 +5130,25 @@ std::unique_ptr<mlir::Pass> createConvertCIRToLLVMPass() {
return std::make_unique<ConvertCIRToLLVMPass>();
}
-void populateCIRToLLVMPasses(mlir::OpPassManager &pm) {
+void populateCIRToLLVMPasses(mlir::OpPassManager &pm, bool enableOpenMP) {
mlir::populateCIRPreLoweringPasses(pm);
- pm.addPass(mlir::omp::createMarkDeclareTargetPass());
+ if (enableOpenMP)
+ pm.addPass(mlir::omp::createMarkDeclareTargetPass());
pm.addPass(createConvertCIRToLLVMPass());
- pm.addPass(mlir::omp::createHostOpFilteringPass());
+ if (enableOpenMP)
+ pm.addPass(mlir::omp::createHostOpFilteringPass());
}
std::unique_ptr<llvm::Module>
lowerDirectlyFromCIRToLLVMIR(mlir::ModuleOp mlirModule, LLVMContext &llvmCtx,
- StringRef mlirSaveTempsOutFile,
+ bool enableOpenMP, StringRef mlirSaveTempsOutFile,
llvm::vfs::FileSystem *fs) {
llvm::TimeTraceScope scope("lower from CIR to LLVM directly");
mlir::MLIRContext *mlirCtx = mlirModule.getContext();
mlir::PassManager pm(mlirCtx);
- populateCIRToLLVMPasses(pm);
+ populateCIRToLLVMPasses(pm, enableOpenMP);
(void)mlir::applyPassManagerCLOptions(pm);
diff --git a/clang/tools/cir-opt/cir-opt.cpp b/clang/tools/cir-opt/cir-opt.cpp
index 6742985e149e8..a76e4ca6a70e9 100644
--- a/clang/tools/cir-opt/cir-opt.cpp
+++ b/clang/tools/cir-opt/cir-opt.cpp
@@ -26,7 +26,12 @@
#include "clang/CIR/Passes.h"
struct CIRToLLVMPipelineOptions
- : public mlir::PassPipelineOptions<CIRToLLVMPipelineOptions> {};
+ : public mlir::PassPipelineOptions<CIRToLLVMPipelineOptions> {
+ Option<bool> enableOpenMP{
+ *this, "enable-openmp",
+ llvm::cl::desc("Add OpenMP-specific CIR-to-LLVM lowering passes"),
+ llvm::cl::init(false)};
+};
int main(int argc, char **argv) {
// TODO: register needed MLIR passes for CIR?
@@ -44,7 +49,7 @@ int main(int argc, char **argv) {
mlir::PassPipelineRegistration<CIRToLLVMPipelineOptions> pipeline(
"cir-to-llvm", "",
[](mlir::OpPassManager &pm, const CIRToLLVMPipelineOptions &options) {
- cir::direct::populateCIRToLLVMPasses(pm);
+ cir::direct::populateCIRToLLVMPasses(pm, options.enableOpenMP);
});
::mlir::registerPass([]() -> std::unique_ptr<::mlir::Pass> {
diff --git a/clang/tools/cir-translate/cir-translate.cpp b/clang/tools/cir-translate/cir-translate.cpp
index 26adf2cf47d17..9bbbb3939a62f 100644
--- a/clang/tools/cir-translate/cir-translate.cpp
+++ b/clang/tools/cir-translate/cir-translate.cpp
@@ -159,8 +159,13 @@ void registerToLLVMTranslation() {
return mlir::failure();
llvm::LLVMContext llvmContext;
+ bool enableOpenMP = false;
+ if (auto offloadMod = llvm::dyn_cast<mlir::omp::OffloadModuleInterface>(
+ cirModule.getOperation()))
+ enableOpenMP = offloadMod.getIsOpenMP();
std::unique_ptr<llvm::Module> llvmModule =
- cir::direct::lowerDirectlyFromCIRToLLVMIR(cirModule, llvmContext);
+ cir::direct::lowerDirectlyFromCIRToLLVMIR(cirModule, llvmContext,
+ enableOpenMP);
if (!llvmModule)
return mlir::failure();
llvmModule->print(output, nullptr);
diff --git a/mlir/include/mlir/Dialect/OpenMP/OpenMPOpsInterfaces.td b/mlir/include/mlir/Dialect/OpenMP/OpenMPOpsInterfaces.td
index 51f925b17f47e..5ba66c323845a 100644
--- a/mlir/include/mlir/Dialect/OpenMP/OpenMPOpsInterfaces.td
+++ b/mlir/include/mlir/Dialect/OpenMP/OpenMPOpsInterfaces.td
@@ -473,6 +473,15 @@ def OffloadModuleInterface : OpInterface<"OffloadModuleInterface"> {
return ::llvm::dyn_cast<BoolAttr>(isTargetDevice).getValue();
return false;
}]>,
+ InterfaceMethod<
+ /*description=*/[{
+ Return true if the current module was generated with OpenMP enabled.
+ }],
+ /*retTy=*/"bool",
+ /*methodName=*/"getIsOpenMP",
+ (ins), [{}], [{
+ return $_op->hasAttr("omp.is_target_device");
+ }]>,
InterfaceMethod<
/*description=*/[{
Set the attribute on the current module with the specified boolean
More information about the cfe-commits
mailing list