[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
Thu Jul 16 09:04:29 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/3] [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/3] 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
>From 3eb5ebe97171062ec310da246edb65bb8489ef76 Mon Sep 17 00:00:00 2001
From: Jan Leyonberg <jan_sjodin at yahoo.com>
Date: Thu, 16 Jul 2026 12:03:44 -0400
Subject: [PATCH 3/3] Change how OpenMP is checked in the module.
---
clang/lib/CIR/CodeGen/CIRGenModule.cpp | 3 ++-
clang/test/CIR/CodeGenOpenMP/omp-module-attrs.c | 1 +
clang/tools/cir-translate/cir-translate.cpp | 6 ++----
mlir/include/mlir/Dialect/OpenMP/OpenMPDialect.h | 2 +-
.../mlir/Dialect/OpenMP/OpenMPOpsInterfaces.td | 9 ---------
.../OpenMP/{OpenMPOffloadUtils.h => OpenMPUtils.h} | 13 +++++++++----
6 files changed, 15 insertions(+), 19 deletions(-)
rename mlir/include/mlir/Dialect/OpenMP/{OpenMPOffloadUtils.h => OpenMPUtils.h} (91%)
diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
index 66a9dc9ea605c..de322c261a8cc 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
@@ -16,7 +16,7 @@
#include "CIRGenConstantEmitter.h"
#include "CIRGenFunction.h"
-#include "mlir/Dialect/OpenMP/OpenMPOffloadUtils.h"
+#include "mlir/Dialect/OpenMP/OpenMPUtils.h"
#include "mlir/IR/SymbolTable.h"
#include "clang/AST/ASTContext.h"
#include "clang/AST/ASTLambda.h"
@@ -151,6 +151,7 @@ CIRGenModule::CIRGenModule(mlir::MLIRContext &mlirContext,
getTriple().isGPU(), langOpts.OpenMPForceUSM, langOpts.OpenMP,
langOpts.OMPHostIRFile, langOpts.OMPTargetTriples, langOpts.NoGPULib);
mlir::omp::setOffloadModuleInterfaceAttributes(theModule, ompOpts);
+ mlir::omp::setOpenMPVersionAttribute(theModule, langOpts.OpenMP);
}
if (langOpts.CUDA)
diff --git a/clang/test/CIR/CodeGenOpenMP/omp-module-attrs.c b/clang/test/CIR/CodeGenOpenMP/omp-module-attrs.c
index 833268b02d312..7481975a6df7e 100644
--- a/clang/test/CIR/CodeGenOpenMP/omp-module-attrs.c
+++ b/clang/test/CIR/CodeGenOpenMP/omp-module-attrs.c
@@ -7,6 +7,7 @@
// HOST: module {{.*}} attributes {
// HOST-SAME: omp.is_gpu = false
// HOST-SAME: omp.is_target_device = false
+// HOST-SAME: omp.version = #omp.version<version = {{[0-9]+}}>
// Host with target triples
// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fopenmp -fclangir -emit-cir \
diff --git a/clang/tools/cir-translate/cir-translate.cpp b/clang/tools/cir-translate/cir-translate.cpp
index 9bbbb3939a62f..cc3b7dec291d1 100644
--- a/clang/tools/cir-translate/cir-translate.cpp
+++ b/clang/tools/cir-translate/cir-translate.cpp
@@ -14,6 +14,7 @@
#include "mlir/Dialect/Func/IR/FuncOps.h"
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
#include "mlir/Dialect/OpenMP/OpenMPDialect.h"
+#include "mlir/Dialect/OpenMP/OpenMPUtils.h"
#include "mlir/IR/BuiltinOps.h"
#include "mlir/IR/MLIRContext.h"
#include "mlir/InitAllTranslations.h"
@@ -159,10 +160,7 @@ 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();
+ const bool enableOpenMP = mlir::omp::isOpenMPModule(cirModule);
std::unique_ptr<llvm::Module> llvmModule =
cir::direct::lowerDirectlyFromCIRToLLVMIR(cirModule, llvmContext,
enableOpenMP);
diff --git a/mlir/include/mlir/Dialect/OpenMP/OpenMPDialect.h b/mlir/include/mlir/Dialect/OpenMP/OpenMPDialect.h
index c61e9415327f0..2507a5a36b514 100644
--- a/mlir/include/mlir/Dialect/OpenMP/OpenMPDialect.h
+++ b/mlir/include/mlir/Dialect/OpenMP/OpenMPDialect.h
@@ -17,7 +17,7 @@
#include "mlir/Dialect/OpenACCMPCommon/Interfaces/AtomicInterfaces.h"
#include "mlir/Dialect/OpenACCMPCommon/Interfaces/OpenACCMPOpsInterfaces.h"
#include "mlir/Dialect/OpenMP/OpenMPInterfaces.h"
-#include "mlir/Dialect/OpenMP/OpenMPOffloadUtils.h"
+#include "mlir/Dialect/OpenMP/OpenMPUtils.h"
#include "mlir/IR/Dialect.h"
#include "mlir/IR/OpDefinition.h"
#include "mlir/IR/PatternMatch.h"
diff --git a/mlir/include/mlir/Dialect/OpenMP/OpenMPOpsInterfaces.td b/mlir/include/mlir/Dialect/OpenMP/OpenMPOpsInterfaces.td
index 5ba66c323845a..51f925b17f47e 100644
--- a/mlir/include/mlir/Dialect/OpenMP/OpenMPOpsInterfaces.td
+++ b/mlir/include/mlir/Dialect/OpenMP/OpenMPOpsInterfaces.td
@@ -473,15 +473,6 @@ 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
diff --git a/mlir/include/mlir/Dialect/OpenMP/OpenMPOffloadUtils.h b/mlir/include/mlir/Dialect/OpenMP/OpenMPUtils.h
similarity index 91%
rename from mlir/include/mlir/Dialect/OpenMP/OpenMPOffloadUtils.h
rename to mlir/include/mlir/Dialect/OpenMP/OpenMPUtils.h
index f1f2950d2325b..002377ae6d24a 100644
--- a/mlir/include/mlir/Dialect/OpenMP/OpenMPOffloadUtils.h
+++ b/mlir/include/mlir/Dialect/OpenMP/OpenMPUtils.h
@@ -1,4 +1,4 @@
-//===- OpenMPOffloadUtils.h - OpenMP offload utilities ----------*- C++ -*-===//
+//===- OpenMPUtils.h - OpenMP utilities -------------------------*- C++ -*-===//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
@@ -12,8 +12,8 @@
//
//===----------------------------------------------------------------------===//
-#ifndef MLIR_DIALECT_OPENMP_OPENMPOFFLOADUTILS_H_
-#define MLIR_DIALECT_OPENMP_OPENMPOFFLOADUTILS_H_
+#ifndef MLIR_DIALECT_OPENMP_OPENMPUTILS_H_
+#define MLIR_DIALECT_OPENMP_OPENMPUTILS_H_
#include "mlir/Dialect/OpenMP/OpenMPInterfaces.h"
#include "mlir/Dialect/OpenMP/OpenMPOpsAttributes.h"
@@ -99,6 +99,11 @@ getOpenMPVersionAttribute(ModuleOp module, int64_t fallback = -1) {
return fallback;
}
+/// Returns true if the given module contains OpenMP code.
+[[maybe_unused]] static bool isOpenMPModule(ModuleOp module) {
+ return module->hasAttr("omp.version");
+}
+
} // namespace mlir::omp
-#endif // MLIR_DIALECT_OPENMP_OPENMPOFFLOADUTILS_H_
+#endif // MLIR_DIALECT_OPENMP_OPENMPUTILS_H_
More information about the cfe-commits
mailing list