From 33c137bf0c92d071fc52b31445174e27355ab27d Mon Sep 17 00:00:00 2001
From: Alexey Bataev <a.bataev@hotmail.com>
Date: Wed, 5 Sep 2018 17:10:30 +0000
Subject: [PATCH] [OPENMP][NVPTX] Disable runtime-type info for CUDA devices.

RTTI is not supported by the NVPTX target.

llvm-svn: 341483
---
 clang/lib/Frontend/CompilerInvocation.cpp     |  5 ++
 .../OpenMP/nvptx_target_rtti_messages.cpp     | 68 +++++++++++++++++++
 2 files changed, 73 insertions(+)
 create mode 100644 clang/test/OpenMP/nvptx_target_rtti_messages.cpp

diff --git a/clang/lib/Frontend/CompilerInvocation.cpp b/clang/lib/Frontend/CompilerInvocation.cpp
index 5210314493de..7a0e78fba638 100644
--- a/clang/lib/Frontend/CompilerInvocation.cpp
+++ b/clang/lib/Frontend/CompilerInvocation.cpp
@@ -2650,6 +2650,11 @@ static void ParseLangArgs(LangOptions &Opts, ArgList &Args, InputKind IK,
     Opts.Exceptions = 0;
     Opts.CXXExceptions = 0;
   }
+  // NVPTX does not support RTTI.
+  if (Opts.OpenMPIsDevice && T.isNVPTX()) {
+    Opts.RTTI = 0;
+    Opts.RTTIData = 0;
+  }
 
   // Get the OpenMP target triples if any.
   if (Arg *A = Args.getLastArg(options::OPT_fopenmp_targets_EQ)) {
diff --git a/clang/test/OpenMP/nvptx_target_rtti_messages.cpp b/clang/test/OpenMP/nvptx_target_rtti_messages.cpp
new file mode 100644
index 000000000000..8e2819e48981
--- /dev/null
+++ b/clang/test/OpenMP/nvptx_target_rtti_messages.cpp
@@ -0,0 +1,68 @@
+// RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc -fexceptions -fcxx-exceptions
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -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 - -ferror-limit 100
+
+#ifndef HEADER
+#define HEADER
+
+namespace std {
+  class type_info;
+}
+
+template <typename T>
+class TemplateClass {
+  T a;
+public:
+  TemplateClass() { (void)typeid(int); } // expected-error {{use of typeid requires -frtti}}
+  T f_method() const { return a; }
+};
+
+int foo();
+
+int baz1();
+
+int baz2();
+
+int baz4() { return 5; }
+
+template <typename T>
+T FA() {
+  TemplateClass<T> s;
+  return s.f_method();
+}
+
+#pragma omp declare target
+struct S {
+  int a;
+  S(int a) : a(a) { (void)typeid(int); } // expected-error {{use of typeid requires -frtti}}
+};
+
+int foo() { return 0; }
+int b = 15;
+int d;
+#pragma omp end declare target
+int c;
+
+int bar() { return 1 + foo() + bar() + baz1() + baz2(); }
+
+int maini1() {
+  int a;
+  static long aa = 32;
+#pragma omp target map(tofrom \
+                       : a, b)
+  {
+    S s(a);
+    static long aaa = 23;
+    a = foo() + bar() + b + c + d + aa + aaa + FA<int>();
+    (void)typeid(int); // expected-error {{use of typeid requires -frtti}}
+  }
+  return baz4();
+}
+
+int baz3() { return 2 + baz2(); }
+int baz2() {
+#pragma omp target
+  (void)typeid(int); // expected-error {{use of typeid requires -frtti}}
+  return 2 + baz3();
+}
+
+#endif // HEADER