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 <[email protected]> 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 <[email protected]> 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 <[email protected]> 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_ _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
