Author: abataev Date: Tue May 1 07:09:46 2018 New Revision: 331261 URL: http://llvm.org/viewvc/llvm-project?rev=331261&view=rev Log: [OPENMP] Emit template instatiation|specialization functions for devices.
If the function is an instantiation|specialization of the template and is used in the device code, the definitions of such functions should be emitted for the device. Modified: cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp cfe/trunk/test/OpenMP/declare_target_codegen.cpp Modified: cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp?rev=331261&r1=331260&r2=331261&view=diff ============================================================================== --- cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp (original) +++ cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp Tue May 1 07:09:46 2018 @@ -907,6 +907,14 @@ isDeclareTargetDeclaration(const ValueDe if (const auto *Attr = D->getAttr<OMPDeclareTargetDeclAttr>()) return Attr->getMapType(); } + if (const auto *V = dyn_cast<VarDecl>(VD)) { + if (const VarDecl *TD = V->getTemplateInstantiationPattern()) + return isDeclareTargetDeclaration(TD); + } else if (const auto *FD = dyn_cast<FunctionDecl>(VD)) { + if (const auto *TD = FD->getTemplateInstantiationPattern()) + return isDeclareTargetDeclaration(TD); + } + return llvm::None; } @@ -7795,7 +7803,8 @@ bool CGOpenMPRuntime::emitTargetFunction scanForTargetRegionsFunctions(FD->getBody(), CGM.getMangledName(GD)); // Do not to emit function if it is not marked as declare target. - return !isDeclareTargetDeclaration(FD); + return !isDeclareTargetDeclaration(FD) && + AlreadyEmittedTargetFunctions.count(FD->getCanonicalDecl()) == 0; } bool CGOpenMPRuntime::emitTargetGlobalVariable(GlobalDecl GD) { Modified: cfe/trunk/test/OpenMP/declare_target_codegen.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/declare_target_codegen.cpp?rev=331261&r1=331260&r2=331261&view=diff ============================================================================== --- cfe/trunk/test/OpenMP/declare_target_codegen.cpp (original) +++ cfe/trunk/test/OpenMP/declare_target_codegen.cpp Tue May 1 07:09:46 2018 @@ -18,7 +18,7 @@ // CHECK-DAG: @d = global i32 0, // CHECK-DAG: @c = external global i32, -// CHECK-DAG: define {{.*}}i32 @{{.*}}{{foo|bar|baz2|baz3}}{{.*}}() +// CHECK-DAG: define {{.*}}i32 @{{.*}}{{foo|bar|baz2|baz3|FA}}{{.*}}() #ifndef HEADER #define HEADER @@ -31,6 +31,11 @@ int baz2(); int baz4() { return 5; } +template <typename T> +T FA() { + return T(); +} + #pragma omp declare target struct S { int a; @@ -54,19 +59,18 @@ int maini1() { { S s(a); static long aaa = 23; - a = foo() + bar() + b + c + d + aa + aaa; + a = foo() + bar() + b + c + d + aa + aaa + FA<int>(); } return baz4(); } -int baz3(); +int baz3() { return 2 + baz2(); } int baz2() { // CHECK-DAG: define void @__omp_offloading_{{.*}}baz2{{.*}}_l[[@LINE+1]](i64 {{.*}}) #pragma omp target ++c; return 2 + baz3(); } -int baz3() { return 2 + baz2(); } // CHECK-NOT: define {{.*}}{{baz1|baz4|maini1}} #endif // HEADER _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits