Author: abataev
Date: Wed May 24 11:00:02 2017
New Revision: 303768
URL: http://llvm.org/viewvc/llvm-project?rev=303768&view=rev
Log:
[OPENMP] Allow value of thread local variables in target regions.
If the variable is marked as TLS variable and target device does not
support TLS, the error is emitted for the variable even if it is not
used in target regions. Patch fixes this and allows to use the values of
the TLS variables in target regions.
Modified:
cfe/trunk/lib/Sema/SemaDecl.cpp
cfe/trunk/test/OpenMP/nvptx_target_codegen.cpp
Modified: cfe/trunk/lib/Sema/SemaDecl.cpp
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDecl.cpp?rev=303768&r1=303767&r2=303768&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaDecl.cpp (original)
+++ cfe/trunk/lib/Sema/SemaDecl.cpp Wed May 24 11:00:02 2017
@@ -6516,7 +6516,7 @@ NamedDecl *Sema::ActOnVariableDeclarator
diag::err_thread_non_global)
<< DeclSpec::getSpecifierName(TSCS);
else if (!Context.getTargetInfo().isTLSSupported()) {
- if (getLangOpts().CUDA) {
+ if (getLangOpts().CUDA || getLangOpts().OpenMPIsDevice) {
// Postpone error emission until we've collected attributes required to
// figure out whether it's a host or device variable and whether the
// error should be ignored.
@@ -6578,8 +6578,11 @@ NamedDecl *Sema::ActOnVariableDeclarator
// Handle attributes prior to checking for duplicates in MergeVarDecl
ProcessDeclAttributes(S, NewVD, D);
- if (getLangOpts().CUDA) {
- if (EmitTLSUnsupportedError && DeclAttrsMatchCUDAMode(getLangOpts(),
NewVD))
+ if (getLangOpts().CUDA || getLangOpts().OpenMPIsDevice) {
+ if (EmitTLSUnsupportedError &&
+ ((getLangOpts().CUDA && DeclAttrsMatchCUDAMode(getLangOpts(), NewVD))
||
+ (getLangOpts().OpenMPIsDevice &&
+ NewVD->hasAttr<OMPDeclareTargetDeclAttr>())))
Diag(D.getDeclSpec().getThreadStorageClassSpecLoc(),
diag::err_thread_unsupported);
// CUDA B.2.5: "__shared__ and __constant__ variables have implied static
Modified: cfe/trunk/test/OpenMP/nvptx_target_codegen.cpp
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/nvptx_target_codegen.cpp?rev=303768&r1=303767&r2=303768&view=diff
==============================================================================
--- cfe/trunk/test/OpenMP/nvptx_target_codegen.cpp (original)
+++ cfe/trunk/test/OpenMP/nvptx_target_codegen.cpp Wed May 24 11:00:02 2017
@@ -9,12 +9,14 @@
#define HEADER
// Check that the execution mode of all 6 target regions is set to Generic
Mode.
-// CHECK-DAG: {{@__omp_offloading_.+l98}}_exec_mode = weak constant i8 1
-// CHECK-DAG: {{@__omp_offloading_.+l175}}_exec_mode = weak constant i8 1
-// CHECK-DAG: {{@__omp_offloading_.+l284}}_exec_mode = weak constant i8 1
-// CHECK-DAG: {{@__omp_offloading_.+l321}}_exec_mode = weak constant i8 1
-// CHECK-DAG: {{@__omp_offloading_.+l339}}_exec_mode = weak constant i8 1
-// CHECK-DAG: {{@__omp_offloading_.+l304}}_exec_mode = weak constant i8 1
+// CHECK-DAG: {{@__omp_offloading_.+l100}}_exec_mode = weak constant i8 1
+// CHECK-DAG: {{@__omp_offloading_.+l177}}_exec_mode = weak constant i8 1
+// CHECK-DAG: {{@__omp_offloading_.+l287}}_exec_mode = weak constant i8 1
+// CHECK-DAG: {{@__omp_offloading_.+l324}}_exec_mode = weak constant i8 1
+// CHECK-DAG: {{@__omp_offloading_.+l342}}_exec_mode = weak constant i8 1
+// CHECK-DAG: {{@__omp_offloading_.+l307}}_exec_mode = weak constant i8 1
+
+__thread int id;
template<typename tx, typename ty>
struct TT{
@@ -31,7 +33,7 @@ int foo(int n) {
double cn[5][n];
TT<long long, char> d;
- // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+foo.+l98}}_worker()
+ // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+foo.+l100}}_worker()
// CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8,
// CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*,
// CHECK: store i8* null, i8** [[OMP_WORK_FN]],
@@ -62,7 +64,7 @@ int foo(int n) {
// CHECK: [[EXIT]]
// CHECK: ret void
- // CHECK: define {{.*}}void [[T1:@__omp_offloading_.+foo.+l98]]()
+ // CHECK: define {{.*}}void [[T1:@__omp_offloading_.+foo.+l100]]()
// CHECK-DAG: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
// CHECK-DAG: [[NTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
// CHECK-DAG: [[WS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
@@ -104,7 +106,7 @@ int foo(int n) {
{
}
- // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+foo.+l175}}_worker()
+ // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+foo.+l177}}_worker()
// CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8,
// CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*,
// CHECK: store i8* null, i8** [[OMP_WORK_FN]],
@@ -135,7 +137,7 @@ int foo(int n) {
// CHECK: [[EXIT]]
// CHECK: ret void
- // CHECK: define {{.*}}void
[[T2:@__omp_offloading_.+foo.+l175]](i[[SZ:32|64]] [[ARG1:%[a-zA-Z_]+]])
+ // CHECK: define {{.*}}void
[[T2:@__omp_offloading_.+foo.+l177]](i[[SZ:32|64]] [[ARG1:%[a-zA-Z_]+]],
i[[SZ:32|64]] [[ID:%[a-zA-Z_]+]])
// CHECK: [[AA_ADDR:%.+]] = alloca i[[SZ]],
// CHECK: store i[[SZ]] [[ARG1]], i[[SZ]]* [[AA_ADDR]],
// CHECK: [[AA_CADDR:%.+]] = bitcast i[[SZ]]* [[AA_ADDR]] to i16*
@@ -175,9 +177,10 @@ int foo(int n) {
#pragma omp target if(1)
{
aa += 1;
+ id = aa;
}
- // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+foo.+l284}}_worker()
+ // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+foo.+l287}}_worker()
// CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8,
// CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*,
// CHECK: store i8* null, i8** [[OMP_WORK_FN]],
@@ -208,7 +211,7 @@ int foo(int n) {
// CHECK: [[EXIT]]
// CHECK: ret void
- // CHECK: define {{.*}}void [[T3:@__omp_offloading_.+foo.+l284]](i[[SZ]]
+ // CHECK: define {{.*}}void [[T3:@__omp_offloading_.+foo.+l287]](i[[SZ]]
// Create local storage for each capture.
// CHECK: [[LOCAL_A:%.+]] = alloca i[[SZ]]
// CHECK: [[LOCAL_B:%.+]] = alloca [10 x float]*
@@ -361,7 +364,7 @@ int bar(int n){
return a;
}
- // CHECK-LABEL: define {{.*}}void
{{@__omp_offloading_.+static.+321}}_worker()
+ // CHECK-LABEL: define {{.*}}void
{{@__omp_offloading_.+static.+324}}_worker()
// CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8,
// CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*,
// CHECK: store i8* null, i8** [[OMP_WORK_FN]],
@@ -392,7 +395,7 @@ int bar(int n){
// CHECK: [[EXIT]]
// CHECK: ret void
- // CHECK: define {{.*}}void [[T4:@__omp_offloading_.+static.+l321]](i[[SZ]]
+ // CHECK: define {{.*}}void [[T4:@__omp_offloading_.+static.+l324]](i[[SZ]]
// Create local storage for each capture.
// CHECK: [[LOCAL_A:%.+]] = alloca i[[SZ]]
// CHECK: [[LOCAL_AA:%.+]] = alloca i[[SZ]]
@@ -447,7 +450,7 @@ int bar(int n){
- // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+S1.+l339}}_worker()
+ // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+S1.+l342}}_worker()
// CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8,
// CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*,
// CHECK: store i8* null, i8** [[OMP_WORK_FN]],
@@ -478,7 +481,7 @@ int bar(int n){
// CHECK: [[EXIT]]
// CHECK: ret void
- // CHECK: define {{.*}}void [[T5:@__omp_offloading_.+S1.+l339]](
+ // CHECK: define {{.*}}void [[T5:@__omp_offloading_.+S1.+l342]](
// Create local storage for each capture.
// CHECK: [[LOCAL_THIS:%.+]] = alloca [[S1:%struct.*]]*
// CHECK: [[LOCAL_B:%.+]] = alloca i[[SZ]]
@@ -537,7 +540,7 @@ int bar(int n){
- // CHECK-LABEL: define {{.*}}void
{{@__omp_offloading_.+template.+l304}}_worker()
+ // CHECK-LABEL: define {{.*}}void
{{@__omp_offloading_.+template.+l307}}_worker()
// CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8,
// CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*,
// CHECK: store i8* null, i8** [[OMP_WORK_FN]],
@@ -568,7 +571,7 @@ int bar(int n){
// CHECK: [[EXIT]]
// CHECK: ret void
- // CHECK: define {{.*}}void [[T6:@__omp_offloading_.+template.+l304]](i[[SZ]]
+ // CHECK: define {{.*}}void [[T6:@__omp_offloading_.+template.+l307]](i[[SZ]]
// Create local storage for each capture.
// CHECK: [[LOCAL_A:%.+]] = alloca i[[SZ]]
// CHECK: [[LOCAL_AA:%.+]] = alloca i[[SZ]]
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
