From 8d6e5d9647a0ad395326e3359a9b590c509b1ab3 Mon Sep 17 00:00:00 2001 From: Gheorghe-Teodor Bercea Date: Mon, 7 Aug 2017 20:57:59 +0000 Subject: [PATCH] [OpenMP] Prevent emission of exception handling code when using OpenMP to offload to NVIDIA devices. Summary: For the OpenMP toolchain which offloads to NVIDIA GPUs make sure that no exception handling code is emitted. Reviewers: arpith-jacob, sfantao, caomhin, carlo.bertolli, ABataev, Hahnfeld, hfinkel, tstellar Reviewed By: ABataev, Hahnfeld Subscribers: rengolin, Hahnfeld, cfe-commits Differential Revision: https://reviews.llvm.org/D29904 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@310306 91177308-0d34-0410-b5e6-96231b3b80d8 --- lib/Frontend/CompilerInvocation.cpp | 7 +++++++ test/OpenMP/target_parallel_no_exceptions.cpp | 18 ++++++++++++++++++ 2 files changed, 25 insertions(+) create mode 100644 test/OpenMP/target_parallel_no_exceptions.cpp diff --git a/lib/Frontend/CompilerInvocation.cpp b/lib/Frontend/CompilerInvocation.cpp index 003ea55471..eb73e28954 100644 --- a/lib/Frontend/CompilerInvocation.cpp +++ b/lib/Frontend/CompilerInvocation.cpp @@ -2308,6 +2308,13 @@ static void ParseLangArgs(LangOptions &Opts, ArgList &Args, InputKind IK, } } + // Set the flag to prevent the implementation from emitting device exception + // handling code for those requiring so. + if (Opts.OpenMPIsDevice && T.isNVPTX()) { + Opts.Exceptions = 0; + Opts.CXXExceptions = 0; + } + // Get the OpenMP target triples if any. if (Arg *A = Args.getLastArg(options::OPT_fopenmp_targets_EQ)) { diff --git a/test/OpenMP/target_parallel_no_exceptions.cpp b/test/OpenMP/target_parallel_no_exceptions.cpp new file mode 100644 index 0000000000..95189a3ade --- /dev/null +++ b/test/OpenMP/target_parallel_no_exceptions.cpp @@ -0,0 +1,18 @@ +/// Make sure no exception messages are inclided in the llvm output. +// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc +// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix CHK-EXCEPTION + +void test_increment() { +#pragma omp target + { + []() { return; }(); + } +} + +int main() { + test_increment(); + return 0; +} + +//CHK-EXCEPTION-NOT: invoke + -- 2.40.0