r326590 - [OPENMP] Treat local variables in CUDA mode as thread local.
Alexey Bataev via cfe-commits
cfe-commits at lists.llvm.org
Fri Mar 2 09:17:12 PST 2018
Author: abataev
Date: Fri Mar 2 09:17:12 2018
New Revision: 326590
URL: http://llvm.org/viewvc/llvm-project?rev=326590&view=rev
Log:
[OPENMP] Treat local variables in CUDA mode as thread local.
In CUDA mode all local variables are actually thread
local|threadprivate, not private, and, thus, they cannot be shared
between threads|lanes.
Added:
cfe/trunk/test/OpenMP/nvptx_target_cuda_mode_messages.cpp
Modified:
cfe/trunk/include/clang/Driver/Options.td
cfe/trunk/lib/Sema/SemaOpenMP.cpp
Modified: cfe/trunk/include/clang/Driver/Options.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Driver/Options.td?rev=326590&r1=326589&r2=326590&view=diff
==============================================================================
--- cfe/trunk/include/clang/Driver/Options.td (original)
+++ cfe/trunk/include/clang/Driver/Options.td Fri Mar 2 09:17:12 2018
@@ -1427,7 +1427,7 @@ def fopenmp_simd : Flag<["-"], "fopenmp-
HelpText<"Emit OpenMP code only for SIMD-based constructs.">;
def fno_openmp_simd : Flag<["-"], "fno-openmp-simd">, Group<f_Group>, Flags<[CC1Option, NoArgumentUnused]>;
def fopenmp_cuda_mode : Flag<["-"], "fopenmp-cuda-mode">, Group<f_Group>, Flags<[CC1Option, NoArgumentUnused]>;
-def fno_openmp_cuda_mode : Flag<["-"], "fno-openmp-cuda-mode">, Group<f_Group>, Flags<[CC1Option, NoArgumentUnused]>;
+def fno_openmp_cuda_mode : Flag<["-"], "fno-openmp-cuda-mode">, Group<f_Group>, Flags<[NoArgumentUnused]>;
def fno_optimize_sibling_calls : Flag<["-"], "fno-optimize-sibling-calls">, Group<f_Group>;
def foptimize_sibling_calls : Flag<["-"], "foptimize-sibling-calls">, Group<f_Group>;
def fno_escaping_block_tail_calls : Flag<["-"], "fno-escaping-block-tail-calls">, Group<f_Group>, Flags<[CC1Option]>;
Modified: cfe/trunk/lib/Sema/SemaOpenMP.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaOpenMP.cpp?rev=326590&r1=326589&r2=326590&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaOpenMP.cpp (original)
+++ cfe/trunk/lib/Sema/SemaOpenMP.cpp Fri Mar 2 09:17:12 2018
@@ -936,10 +936,11 @@ DSAStackTy::getTopMostTaskgroupReduction
bool DSAStackTy::isOpenMPLocal(VarDecl *D, StackTy::reverse_iterator Iter) {
D = D->getCanonicalDecl();
- if (!isStackEmpty() && Stack.back().first.size() > 1) {
+ if (!isStackEmpty()) {
reverse_iterator I = Iter, E = Stack.back().first.rend();
Scope *TopScope = nullptr;
- while (I != E && !isParallelOrTaskRegion(I->Directive))
+ while (I != E && !isParallelOrTaskRegion(I->Directive) &&
+ !isOpenMPTargetExecutionDirective(I->Directive))
++I;
if (I == E)
return false;
@@ -956,20 +957,7 @@ DSAStackTy::DSAVarData DSAStackTy::getTo
D = getCanonicalDecl(D);
DSAVarData DVar;
- // OpenMP [2.9.1.1, Data-sharing Attribute Rules for Variables Referenced
- // in a Construct, C/C++, predetermined, p.1]
- // Variables appearing in threadprivate directives are threadprivate.
auto *VD = dyn_cast<VarDecl>(D);
- if ((VD && VD->getTLSKind() != VarDecl::TLS_None &&
- !(VD->hasAttr<OMPThreadPrivateDeclAttr>() &&
- SemaRef.getLangOpts().OpenMPUseTLS &&
- SemaRef.getASTContext().getTargetInfo().isTLSSupported())) ||
- (VD && VD->getStorageClass() == SC_Register &&
- VD->hasAttr<AsmLabelAttr>() && !VD->isLocalVarDecl())) {
- addDSA(D, buildDeclRefExpr(SemaRef, VD, D->getType().getNonReferenceType(),
- D->getLocation()),
- OMPC_threadprivate);
- }
auto TI = Threadprivates.find(D);
if (TI != Threadprivates.end()) {
DVar.RefExpr = TI->getSecond().RefExpr.getPointer();
@@ -981,6 +969,62 @@ DSAStackTy::DSAVarData DSAStackTy::getTo
VD->getAttr<OMPThreadPrivateDeclAttr>()->getLocation());
DVar.CKind = OMPC_threadprivate;
addDSA(D, DVar.RefExpr, OMPC_threadprivate);
+ return DVar;
+ }
+ // OpenMP [2.9.1.1, Data-sharing Attribute Rules for Variables Referenced
+ // in a Construct, C/C++, predetermined, p.1]
+ // Variables appearing in threadprivate directives are threadprivate.
+ if ((VD && VD->getTLSKind() != VarDecl::TLS_None &&
+ !(VD->hasAttr<OMPThreadPrivateDeclAttr>() &&
+ SemaRef.getLangOpts().OpenMPUseTLS &&
+ SemaRef.getASTContext().getTargetInfo().isTLSSupported())) ||
+ (VD && VD->getStorageClass() == SC_Register &&
+ VD->hasAttr<AsmLabelAttr>() && !VD->isLocalVarDecl())) {
+ DVar.RefExpr = buildDeclRefExpr(
+ SemaRef, VD, D->getType().getNonReferenceType(), D->getLocation());
+ DVar.CKind = OMPC_threadprivate;
+ addDSA(D, DVar.RefExpr, OMPC_threadprivate);
+ return DVar;
+ }
+ if (SemaRef.getLangOpts().OpenMPCUDAMode && VD &&
+ VD->isLocalVarDeclOrParm() && !isStackEmpty() &&
+ !isLoopControlVariable(D).first) {
+ auto IterTarget =
+ std::find_if(Stack.back().first.rbegin(), Stack.back().first.rend(),
+ [](const SharingMapTy &Data) {
+ return isOpenMPTargetExecutionDirective(Data.Directive);
+ });
+ if (IterTarget != Stack.back().first.rend()) {
+ auto ParentIterTarget = std::next(IterTarget, 1);
+ auto Iter = Stack.back().first.rbegin();
+ while (Iter != ParentIterTarget) {
+ if (isOpenMPLocal(VD, Iter)) {
+ DVar.RefExpr =
+ buildDeclRefExpr(SemaRef, VD, D->getType().getNonReferenceType(),
+ D->getLocation());
+ DVar.CKind = OMPC_threadprivate;
+ return DVar;
+ }
+ std::advance(Iter, 1);
+ }
+ if (!isClauseParsingMode() || IterTarget != Stack.back().first.rbegin()) {
+ auto DSAIter = IterTarget->SharingMap.find(D);
+ if (DSAIter != IterTarget->SharingMap.end() &&
+ isOpenMPPrivate(DSAIter->getSecond().Attributes)) {
+ DVar.RefExpr = DSAIter->getSecond().RefExpr.getPointer();
+ DVar.CKind = OMPC_threadprivate;
+ return DVar;
+ } else if (!SemaRef.IsOpenMPCapturedByRef(
+ D, std::distance(ParentIterTarget,
+ Stack.back().first.rend()))) {
+ DVar.RefExpr =
+ buildDeclRefExpr(SemaRef, VD, D->getType().getNonReferenceType(),
+ IterTarget->ConstructLoc);
+ DVar.CKind = OMPC_threadprivate;
+ return DVar;
+ }
+ }
+ }
}
if (isStackEmpty())
Added: cfe/trunk/test/OpenMP/nvptx_target_cuda_mode_messages.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/nvptx_target_cuda_mode_messages.cpp?rev=326590&view=auto
==============================================================================
--- cfe/trunk/test/OpenMP/nvptx_target_cuda_mode_messages.cpp (added)
+++ cfe/trunk/test/OpenMP/nvptx_target_cuda_mode_messages.cpp Fri Mar 2 09:17:12 2018
@@ -0,0 +1,108 @@
+// RUN: %clang_cc1 -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-cuda-mode -fopenmp-host-ir-file-path %t-ppc-host.bc -o -
+
+template <typename tx, typename ty>
+struct TT {
+ tx X;
+ ty Y;
+};
+
+int foo(int n, double *ptr) {
+ int a = 0;
+ short aa = 0;
+ float b[10];
+ double c[5][10];
+ TT<long long, char> d;
+
+#pragma omp target firstprivate(a) map(tofrom: b) // expected-note 2 {{defined as threadprivate or thread local}}
+ {
+ int c; // expected-note {{defined as threadprivate or thread local}}
+#pragma omp parallel shared(a, b, c, aa) // expected-error 3 {{threadprivate or thread local variable cannot be shared}}
+ b[a] = a;
+#pragma omp parallel for
+ for (int i = 0; i < 10; ++i) // expected-note {{defined as threadprivate or thread local}}
+#pragma omp parallel shared(i) // expected-error {{threadprivate or thread local variable cannot be shared}}
+ ++i;
+ }
+
+#pragma omp target map(aa, b, c, d)
+ {
+ int e; // expected-note {{defined as threadprivate or thread local}}
+#pragma omp parallel private(b, e) // expected-error {{threadprivate or thread local variable cannot be private}}
+ {
+ aa += 1;
+ b[2] = 1.0;
+ c[1][2] = 1.0;
+ d.X = 1;
+ d.Y = 1;
+ }
+ }
+
+#pragma omp target private(ptr)
+ {
+ ptr[0]++;
+ }
+
+ return a;
+}
+
+template <typename tx>
+tx ftemplate(int n) {
+ tx a = 0;
+ tx b[10];
+
+#pragma omp target reduction(+ \
+ : a, b) // expected-note {{defined as threadprivate or thread local}}
+ {
+ int e; // expected-note {{defined as threadprivate or thread local}}
+#pragma omp parallel shared(a, e) // expected-error 2 {{threadprivate or thread local variable cannot be shared}}
+ a += 1;
+ b[2] += 1;
+ }
+
+ return a;
+}
+
+static int fstatic(int n) {
+ int a = 0;
+ char aaa = 0;
+ int b[10];
+
+#pragma omp target firstprivate(a, aaa, b)
+ {
+ a += 1;
+ aaa += 1;
+ b[2] += 1;
+ }
+
+ return a;
+}
+
+struct S1 {
+ double a;
+
+ int r1(int n) {
+ int b = n + 1;
+
+#pragma omp target firstprivate(b) // expected-note {{defined as threadprivate or thread local}}
+ {
+ int c; // expected-note {{defined as threadprivate or thread local}}
+#pragma omp parallel shared(b, c) // expected-error 2 {{threadprivate or thread local variable cannot be shared}}
+ this->a = (double)b + 1.5;
+ }
+
+ return (int)b;
+ }
+};
+
+int bar(int n, double *ptr) {
+ int a = 0;
+ a += foo(n, ptr);
+ S1 S;
+ a += S.r1(n);
+ a += fstatic(n);
+ a += ftemplate<int>(n); // expected-note {{in instantiation of function template specialization 'ftemplate<int>' requested here}}
+
+ return a;
+}
+
More information about the cfe-commits
mailing list