From 2945f11c605b059446ac997f62458a6c489e46f7 Mon Sep 17 00:00:00 2001 From: Joseph Huber Date: Thu, 27 Jan 2022 15:50:18 -0500 Subject: [PATCH] [OpenMP] Only generate runtime flags with host input This patch changes the code generation of runtime flags to only occur if a host bitcode file was passed in. This is a cheap way to determine if we are compiling the OpenMP device runtime itself or user code. This is needed because the global flags we generate for the device runtime e.g. __omp_rtl_debug_kind were being generated with default values when we compiled the runtime library. This would then invalidate the ones we want to be able to add in when the user defines it. Reviewed By: jdoerfert Differential Revision: https://reviews.llvm.org/D118399 --- clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp | 3 ++- clang/test/OpenMP/target_globals_codegen.cpp | 5 +++++ 2 files changed, 7 insertions(+), 1 deletion(-) diff --git a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp index e09ea5e..9ca0b20 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp @@ -1198,7 +1198,8 @@ CGOpenMPRuntimeGPU::CGOpenMPRuntimeGPU(CodeGenModule &CGM) llvm_unreachable("OpenMP can only handle device code."); llvm::OpenMPIRBuilder &OMPBuilder = getOMPBuilder(); - if (CGM.getLangOpts().OpenMPTargetNewRuntime) { + if (CGM.getLangOpts().OpenMPTargetNewRuntime && + !CGM.getLangOpts().OMPHostIRFile.empty()) { OMPBuilder.createGlobalFlag(CGM.getLangOpts().OpenMPTargetDebug, "__omp_rtl_debug_kind"); OMPBuilder.createGlobalFlag(CGM.getLangOpts().OpenMPTeamSubscription, diff --git a/clang/test/OpenMP/target_globals_codegen.cpp b/clang/test/OpenMP/target_globals_codegen.cpp index 62e1a9a..7aec922 100644 --- a/clang/test/OpenMP/target_globals_codegen.cpp +++ b/clang/test/OpenMP/target_globals_codegen.cpp @@ -6,6 +6,7 @@ // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-target-new-runtime -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=CHECK-DEFAULT // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-target-new-runtime -fopenmp-assume-threads-oversubscription -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=CHECK-THREADS // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-target-new-runtime -fopenmp-assume-teams-oversubscription -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=CHECK-TEAMS +// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-target-new-runtime -fopenmp-assume-teams-oversubscription -fopenmp-is-device -o - | FileCheck %s --check-prefix=CHECK-RUNTIME // expected-no-diagnostics #ifndef HEADER @@ -32,6 +33,10 @@ // CHECK-TEAMS: @__omp_rtl_assume_teams_oversubscription = weak_odr hidden constant i32 1 // CHECK-TEAMS: @__omp_rtl_assume_threads_oversubscription = weak_odr hidden constant i32 0 //. +// CHECK-RUNTIME-NOT: @__omp_rtl_debug_kind = weak_odr hidden constant i32 0 +// CHECK-RUNTIME-NOT: @__omp_rtl_assume_teams_oversubscription = weak_odr hidden constant i32 1 +// CHECK-RUNTIME-NOT: @__omp_rtl_assume_threads_oversubscription = weak_odr hidden constant i32 0 +//. void foo() { #pragma omp target { } -- 2.7.4