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