maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

[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.
This commit is contained in:
Alexey Bataev 2020-01-14 16:42:23 -05:00
parent 1bd14ce392
commit a48600c0a6
2 changed files with 40 additions and 1 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

7
clang/lib/CodeGen/CGVTables.cpp
View File

@ -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>(

34
clang/test/OpenMP/nvptx_target_pure_deleted_codegen.cpp Normal file
View File

@ -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