Skip to content

Commit 979966f

Browse files
alexey-bataevalexey-bataev
committed
[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. llvm-svn: 303768
1 parent d20066c commit 979966f

File tree

2 files changed

+27
-21
lines changed

2 files changed

+27
-21
lines changed

‎clang/lib/Sema/SemaDecl.cpp

Lines changed: 6 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -6516,7 +6516,7 @@ NamedDecl *Sema::ActOnVariableDeclarator(
65166516
diag::err_thread_non_global)
65176517
<< DeclSpec::getSpecifierName(TSCS);
65186518
else if (!Context.getTargetInfo().isTLSSupported()) {
6519-
if (getLangOpts().CUDA) {
6519+
if (getLangOpts().CUDA || getLangOpts().OpenMPIsDevice) {
65206520
// Postpone error emission until we've collected attributes required to
65216521
// figure out whether it's a host or device variable and whether the
65226522
// error should be ignored.
@@ -6578,8 +6578,11 @@ NamedDecl *Sema::ActOnVariableDeclarator(
65786578
// Handle attributes prior to checking for duplicates in MergeVarDecl
65796579
ProcessDeclAttributes(S, NewVD, D);
65806580

6581-
if (getLangOpts().CUDA) {
6582-
if (EmitTLSUnsupportedError && DeclAttrsMatchCUDAMode(getLangOpts(), NewVD))
6581+
if (getLangOpts().CUDA || getLangOpts().OpenMPIsDevice) {
6582+
if (EmitTLSUnsupportedError &&
6583+
((getLangOpts().CUDA && DeclAttrsMatchCUDAMode(getLangOpts(), NewVD)) ||
6584+
(getLangOpts().OpenMPIsDevice &&
6585+
NewVD->hasAttr<OMPDeclareTargetDeclAttr>())))
65836586
Diag(D.getDeclSpec().getThreadStorageClassSpecLoc(),
65846587
diag::err_thread_unsupported);
65856588
// CUDA B.2.5: "__shared__ and __constant__ variables have implied static

‎clang/test/OpenMP/nvptx_target_codegen.cpp

Lines changed: 21 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -9,12 +9,14 @@
99
#define HEADER
1010

1111
// Check that the execution mode of all 6 target regions is set to Generic Mode.
12-
// CHECK-DAG: {{@__omp_offloading_.+l98}}_exec_mode = weak constant i8 1
13-
// CHECK-DAG: {{@__omp_offloading_.+l175}}_exec_mode = weak constant i8 1
14-
// CHECK-DAG: {{@__omp_offloading_.+l284}}_exec_mode = weak constant i8 1
15-
// CHECK-DAG: {{@__omp_offloading_.+l321}}_exec_mode = weak constant i8 1
16-
// CHECK-DAG: {{@__omp_offloading_.+l339}}_exec_mode = weak constant i8 1
17-
// CHECK-DAG: {{@__omp_offloading_.+l304}}_exec_mode = weak constant i8 1
12+
// CHECK-DAG: {{@__omp_offloading_.+l100}}_exec_mode = weak constant i8 1
13+
// CHECK-DAG: {{@__omp_offloading_.+l177}}_exec_mode = weak constant i8 1
14+
// CHECK-DAG: {{@__omp_offloading_.+l287}}_exec_mode = weak constant i8 1
15+
// CHECK-DAG: {{@__omp_offloading_.+l324}}_exec_mode = weak constant i8 1
16+
// CHECK-DAG: {{@__omp_offloading_.+l342}}_exec_mode = weak constant i8 1
17+
// CHECK-DAG: {{@__omp_offloading_.+l307}}_exec_mode = weak constant i8 1
18+
19+
__thread int id;
1820

1921
template<typename tx, typename ty>
2022
struct TT{
@@ -31,7 +33,7 @@ int foo(int n) {
3133
double cn[5][n];
3234
TT<long long, char> d;
3335

34-
// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+foo.+l98}}_worker()
36+
// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+foo.+l100}}_worker()
3537
// CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8,
3638
// CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*,
3739
// CHECK: store i8* null, i8** [[OMP_WORK_FN]],
@@ -62,7 +64,7 @@ int foo(int n) {
6264
// CHECK: [[EXIT]]
6365
// CHECK: ret void
6466

65-
// CHECK: define {{.*}}void [[T1:@__omp_offloading_.+foo.+l98]]()
67+
// CHECK: define {{.*}}void [[T1:@__omp_offloading_.+foo.+l100]]()
6668
// CHECK-DAG: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
6769
// CHECK-DAG: [[NTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
6870
// CHECK-DAG: [[WS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
@@ -104,7 +106,7 @@ int foo(int n) {
104106
{
105107
}
106108

107-
// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+foo.+l175}}_worker()
109+
// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+foo.+l177}}_worker()
108110
// CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8,
109111
// CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*,
110112
// CHECK: store i8* null, i8** [[OMP_WORK_FN]],
@@ -135,7 +137,7 @@ int foo(int n) {
135137
// CHECK: [[EXIT]]
136138
// CHECK: ret void
137139

138-
// CHECK: define {{.*}}void [[T2:@__omp_offloading_.+foo.+l175]](i[[SZ:32|64]] [[ARG1:%[a-zA-Z_]+]])
140+
// CHECK: define {{.*}}void [[T2:@__omp_offloading_.+foo.+l177]](i[[SZ:32|64]] [[ARG1:%[a-zA-Z_]+]], i[[SZ:32|64]] [[ID:%[a-zA-Z_]+]])
139141
// CHECK: [[AA_ADDR:%.+]] = alloca i[[SZ]],
140142
// CHECK: store i[[SZ]] [[ARG1]], i[[SZ]]* [[AA_ADDR]],
141143
// CHECK: [[AA_CADDR:%.+]] = bitcast i[[SZ]]* [[AA_ADDR]] to i16*
@@ -175,9 +177,10 @@ int foo(int n) {
175177
#pragma omp target if(1)
176178
{
177179
aa += 1;
180+
id = aa;
178181
}
179182

180-
// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+foo.+l284}}_worker()
183+
// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+foo.+l287}}_worker()
181184
// CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8,
182185
// CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*,
183186
// CHECK: store i8* null, i8** [[OMP_WORK_FN]],
@@ -208,7 +211,7 @@ int foo(int n) {
208211
// CHECK: [[EXIT]]
209212
// CHECK: ret void
210213

211-
// CHECK: define {{.*}}void [[T3:@__omp_offloading_.+foo.+l284]](i[[SZ]]
214+
// CHECK: define {{.*}}void [[T3:@__omp_offloading_.+foo.+l287]](i[[SZ]]
212215
// Create local storage for each capture.
213216
// CHECK: [[LOCAL_A:%.+]] = alloca i[[SZ]]
214217
// CHECK: [[LOCAL_B:%.+]] = alloca [10 x float]*
@@ -361,7 +364,7 @@ int bar(int n){
361364
return a;
362365
}
363366

364-
// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+static.+321}}_worker()
367+
// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+static.+324}}_worker()
365368
// CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8,
366369
// CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*,
367370
// CHECK: store i8* null, i8** [[OMP_WORK_FN]],
@@ -392,7 +395,7 @@ int bar(int n){
392395
// CHECK: [[EXIT]]
393396
// CHECK: ret void
394397

395-
// CHECK: define {{.*}}void [[T4:@__omp_offloading_.+static.+l321]](i[[SZ]]
398+
// CHECK: define {{.*}}void [[T4:@__omp_offloading_.+static.+l324]](i[[SZ]]
396399
// Create local storage for each capture.
397400
// CHECK: [[LOCAL_A:%.+]] = alloca i[[SZ]]
398401
// CHECK: [[LOCAL_AA:%.+]] = alloca i[[SZ]]
@@ -447,7 +450,7 @@ int bar(int n){
447450

448451

449452

450-
// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+S1.+l339}}_worker()
453+
// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+S1.+l342}}_worker()
451454
// CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8,
452455
// CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*,
453456
// CHECK: store i8* null, i8** [[OMP_WORK_FN]],
@@ -478,7 +481,7 @@ int bar(int n){
478481
// CHECK: [[EXIT]]
479482
// CHECK: ret void
480483

481-
// CHECK: define {{.*}}void [[T5:@__omp_offloading_.+S1.+l339]](
484+
// CHECK: define {{.*}}void [[T5:@__omp_offloading_.+S1.+l342]](
482485
// Create local storage for each capture.
483486
// CHECK: [[LOCAL_THIS:%.+]] = alloca [[S1:%struct.*]]*
484487
// CHECK: [[LOCAL_B:%.+]] = alloca i[[SZ]]
@@ -537,7 +540,7 @@ int bar(int n){
537540

538541

539542

540-
// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l304}}_worker()
543+
// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l307}}_worker()
541544
// CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8,
542545
// CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*,
543546
// CHECK: store i8* null, i8** [[OMP_WORK_FN]],
@@ -568,7 +571,7 @@ int bar(int n){
568571
// CHECK: [[EXIT]]
569572
// CHECK: ret void
570573

571-
// CHECK: define {{.*}}void [[T6:@__omp_offloading_.+template.+l304]](i[[SZ]]
574+
// CHECK: define {{.*}}void [[T6:@__omp_offloading_.+template.+l307]](i[[SZ]]
572575
// Create local storage for each capture.
573576
// CHECK: [[LOCAL_A:%.+]] = alloca i[[SZ]]
574577
// CHECK: [[LOCAL_AA:%.+]] = alloca i[[SZ]]

0 commit comments

Comments
 (0)