[clang] a48600c - [OPENMP]Do not emit special virtual function for NVPTX target.

[Alexey Bataev via cfe-commits]

Author: Alexey Bataev
Date: 2020-01-14T16:59:22-05:00
New Revision: a48600c0a653d34f4af760f117755ed1776adf9d

URL: https://github.com/llvm/llvm-project/commit/a48600c0a653d34f4af760f117755ed1776adf9d
DIFF: https://github.com/llvm/llvm-project/commit/a48600c0a653d34f4af760f117755ed1776adf9d.diff

LOG: [OPENMP]Do not emit special virtual function for NVPTX target.

There are no special virtual function handlers (like __cxa_pure_virtual)
defined for NVPTX target, so just emit such functions as null pointers
to prevent issues with linking and unresolved references.

Added: 
    clang/test/OpenMP/nvptx_target_pure_deleted_codegen.cpp

Modified: 
    clang/lib/CodeGen/CGVTables.cpp

Removed: 
    


################################################################################
diff  --git a/clang/lib/CodeGen/CGVTables.cpp b/clang/lib/CodeGen/CGVTables.cpp
index a8dce3c2e859..59631e802373 100644
--- a/clang/lib/CodeGen/CGVTables.cpp
+++ b/clang/lib/CodeGen/CGVTables.cpp
@@ -676,7 +676,12 @@ void CodeGenVTables::addVTableComponent(
       // Method is acceptable, continue processing as usual.
     }
 
-    auto getSpecialVirtualFn = [&](StringRef name) {
+    auto getSpecialVirtualFn = [&](StringRef name) -> llvm::Constant * {
+      // For NVPTX devices in OpenMP emit special functon as null pointers,
+      // otherwise linking ends up with unresolved references.
+      if (CGM.getLangOpts().OpenMP && CGM.getLangOpts().OpenMPIsDevice &&
+          CGM.getTriple().isNVPTX())
+        return llvm::ConstantPointerNull::get(CGM.Int8PtrTy);
       llvm::FunctionType *fnTy =
           llvm::FunctionType::get(CGM.VoidTy, /*isVarArg=*/false);
       llvm::Constant *fn = cast<llvm::Constant>(

diff  --git a/clang/test/OpenMP/nvptx_target_pure_deleted_codegen.cpp b/clang/test/OpenMP/nvptx_target_pure_deleted_codegen.cpp
new file mode 100644
index 000000000000..1e83cf6f8704
--- /dev/null
+++ b/clang/test/OpenMP/nvptx_target_pure_deleted_codegen.cpp
@@ -0,0 +1,34 @@
+// Test target codegen - host bc file has to be created first.
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
+// 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 - -fno-rtti | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - -fno-rtti | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -fexceptions -fcxx-exceptions -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - -fno-rtti | FileCheck %s
+// expected-no-diagnostics
+
+#ifndef HEADER
+#define HEADER
+
+
+// CHECK-DAG: @_ZTV7Derived = linkonce_odr hidden unnamed_addr constant { [3 x i8*] } { [3 x i8*] [i8* null, i8* null, i8* bitcast (void (%class.Derived*)* @_ZN7Derived3fooEv to i8*)] }
+// CHECK-DAG: @_ZTV4Base = linkonce_odr hidden unnamed_addr constant { [3 x i8*] } zeroinitializer
+class Base {
+  public:
+  virtual void foo() = 0;
+};
+
+class Derived : public Base {
+public:
+  void foo() override {}
+  void bar() = delete;
+};
+
+void target() {
+#pragma omp target
+  {
+  Derived D;
+  D.foo();
+  }
+}
+
+#endif