jdenny updated this revision to Diff 216420.
jdenny edited the summary of this revision.
jdenny set the repository for this revision to rG LLVM Github Monorepo.
jdenny added a comment.
As requested, backed out the changes related to firstprivate scalars, and
updated tests.
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D65835/new/
https://reviews.llvm.org/D65835
Files:
clang/lib/Sema/SemaOpenMP.cpp
clang/test/OpenMP/target_teams_distribute_parallel_for_lastprivate_messages.cpp
clang/test/OpenMP/target_teams_distribute_parallel_for_simd_lastprivate_messages.cpp
clang/test/OpenMP/target_teams_distribute_simd_firstprivate_messages.cpp
clang/test/OpenMP/target_teams_distribute_simd_lastprivate_messages.cpp
clang/test/OpenMP/target_teams_distribute_simd_private_messages.cpp
clang/test/OpenMP/target_teams_map_codegen.cpp
clang/test/OpenMP/target_teams_map_messages.cpp
Index: clang/test/OpenMP/target_teams_map_messages.cpp
===================================================================
--- clang/test/OpenMP/target_teams_map_messages.cpp
+++ clang/test/OpenMP/target_teams_map_messages.cpp
@@ -1,7 +1,10 @@
-// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=40 -fopenmp -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=45 -fopenmp -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -verify=expected -fopenmp-version=50 -fopenmp -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
-// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
-// RUN: %clang_cc1 -DCCODE -verify -fopenmp -ferror-limit 200 -x c %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -DCCODE -verify=expected,le45 -fopenmp -ferror-limit 200 -x c %s -Wno-openmp-target -Wuninitialized
#ifdef CCODE
void foo(int arg) {
const int n = 0;
@@ -536,10 +539,10 @@
#pragma omp target teams map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
foo();
-#pragma omp target teams private(j) map(j) // expected-error {{private variable cannot be in a map clause in '#pragma omp target teams' directive}} expected-note {{defined as private}}
+#pragma omp target teams private(j) map(j) // le45-error {{private variable cannot be in a map clause in '#pragma omp target teams' directive}} le45-note {{defined as private}}
{}
-#pragma omp target teams firstprivate(j) map(j) // expected-error {{firstprivate variable cannot be in a map clause in '#pragma omp target teams' directive}} expected-note {{defined as firstprivate}}
+#pragma omp target teams firstprivate(j) map(j) // le45-error {{firstprivate variable cannot be in a map clause in '#pragma omp target teams' directive}} le45-note {{defined as firstprivate}}
{}
#pragma omp target teams map(m)
Index: clang/test/OpenMP/target_teams_map_codegen.cpp
===================================================================
--- /dev/null
+++ clang/test/OpenMP/target_teams_map_codegen.cpp
@@ -0,0 +1,172 @@
+// Test host codegen.
+// RUN: %clang_cc1 -DHAS_INT128 -verify -fopenmp -fopenmp-version=50 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s -check-prefixes=CHECK,HOST,INT128,HOST-INT128
+// RUN: %clang_cc1 -DHAS_INT128 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
+// RUN: %clang_cc1 -DHAS_INT128 -fopenmp -fopenmp-version=50 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s -check-prefixes=CHECK,HOST,INT128,HOST-INT128
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=50 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s -check-prefixes=CHECK,HOST
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s -check-prefixes=CHECK,HOST
+
+// Test target codegen - host bc file has to be created first.
+// RUN: %clang_cc1 -DHAS_INT128 -verify -fopenmp -fopenmp-version=50 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1 -DHAS_INT128 -verify -fopenmp -fopenmp-version=50 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s -check-prefixes=CHECK,INT128
+// RUN: %clang_cc1 -DHAS_INT128 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s
+// RUN: %clang_cc1 -DHAS_INT128 -fopenmp -fopenmp-version=50 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s -check-prefixes=CHECK,INT128
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=50 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=50 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
+
+// expected-no-diagnostics
+#ifndef HEADER
+#define HEADER
+
+// HOST: @[[MAPTYPES_FIRSTPRIVATE:.offload_maptypes[0-9.]*]] = private {{.*}}constant [2 x i64] [i64 35, i64 35]
+// HOST: @[[MAPTYPES_REDUCTION:.offload_maptypes[0-9.]*]] = private {{.*}}constant [2 x i64] [i64 35, i64 35]
+// HOST: @[[MAPTYPES_FROM:.offload_maptypes[0-9.]*]] = private {{.*}}constant [1 x i64] [i64 34]
+// HOST: @[[MAPTYPES_TO:.offload_maptypes[0-9.]*]] = private {{.*}}constant [1 x i64] [i64 33]
+// HOST: @[[MAPTYPES_ALLOC:.offload_maptypes[0-9.]*]] = private {{.*}}constant [1 x i64] [i64 32]
+// HOST: @[[MAPTYPES_ARRAY_R0:.offload_maptypes[0-9.]*]] = private {{.*}}constant [2 x i64] [i64 35, i64 35]
+// HOST: @[[MAPTYPES_ARRAY_R1:.offload_maptypes[0-9.]*]] = private {{.*}}constant [2 x i64] [i64 33, i64 33]
+// HOST-INT128: @[[MAPTYPES_INT128_R0:.offload_maptypes[0-9.]*]] = private {{.*}}constant [2 x i64] [i64 35, i64 35]
+// HOST-INT128: @[[MAPTYPES_INT128_R1:.offload_maptypes[0-9.]*]] = private {{.*}}constant [2 x i64] [i64 34, i64 34]
+//
+// CHECK: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_PRIVATE:__omp_offloading_[^"\\]*mapWithPrivate[^"\\]*]]\00"
+// CHECK: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_FIRSTPRIVATE:__omp_offloading_[^"\\]*mapWithFirstprivate[^"\\]*]]\00"
+// CHECK: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_REDUCTION:__omp_offloading_[^"\\]*mapWithReduction[^"\\]*]]\00"
+// CHECK: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_FROM:__omp_offloading_[^"\\]*mapFrom[^"\\]*]]\00"
+// CHECK: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_TO:__omp_offloading_[^"\\]*mapTo[^"\\]*]]\00"
+// CHECK: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_ALLOC:__omp_offloading_[^"\\]*mapAlloc[^"\\]*]]\00"
+// CHECK: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_ARRAY_R0:__omp_offloading_[^"\\]*mapArray[^"\\]*]]\00"
+// CHECK: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_ARRAY_R1:__omp_offloading_[^"\\]*mapArray[^"\\]*]]\00"
+// INT128: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_INT128_R0:__omp_offloading_[^"\\]*mapInt128[^"\\]*]]\00"
+// INT128: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_INT128_R1:__omp_offloading_[^"\\]*mapInt128[^"\\]*]]\00"
+
+// HOST: define {{.*}}mapWithPrivate
+// HOST: call {{.*}} @.[[OFFLOAD_PRIVATE]].region_id
+// HOST-NOT: offload_maptypes
+// HOST-SAME: {{$}}
+//
+// CHECK: define {{.*}} void @[[OFFLOAD_PRIVATE]]()
+// CHECK: call void ({{.*}}@[[OUTLINE_PRIVATE:.omp_outlined.[.0-9]*]]
+//
+// CHECK: define {{.*}} void @[[OUTLINE_PRIVATE]]({{.*}} %.global_tid., {{.*}} %.bound_tid.)
+void mapWithPrivate() {
+ int x, y;
+ #pragma omp target teams private(x) map(x,y) private(y)
+ ;
+}
+
+// HOST: define {{.*}}mapWithFirstprivate
+// HOST: call {{.*}} @.[[OFFLOAD_FIRSTPRIVATE]].region_id{{.*}} @[[MAPTYPES_FIRSTPRIVATE]]
+//
+// CHECK: define {{.*}} void @[[OFFLOAD_FIRSTPRIVATE]](i{{[0-9]*}} %x, i{{[0-9]*}} %y)
+// CHECK: call void ({{.*}}@[[OUTLINE_FIRSTPRIVATE:.omp_outlined.[.0-9]*]]
+//
+// CHECK: define {{.*}} void @[[OUTLINE_FIRSTPRIVATE]]({{.*}} %.global_tid., {{.*}} %.bound_tid., i{{[0-9]*}} %x, i{{[0-9]*}} %y)
+void mapWithFirstprivate() {
+ int x, y;
+ #pragma omp target teams firstprivate(x) map(x,y) firstprivate(y)
+ ;
+}
+
+// HOST: define {{.*}}mapWithReduction
+// HOST: call {{.*}} @.[[OFFLOAD_REDUCTION]].region_id{{.*}} @[[MAPTYPES_REDUCTION]]
+//
+// CHECK: define {{.*}} void @[[OFFLOAD_REDUCTION]](i{{[0-9]*}}* {{[^,]*}}%x, i{{[0-9]*}}* {{[^,]*}}%y)
+// CHECK: call void ({{.*}}@[[OUTLINE_REDUCTION:.omp_outlined.[.0-9]*]]
+//
+// CHECK: define {{.*}} void @[[OUTLINE_REDUCTION]]({{[^)]*}} i{{[0-9]*}}* {{[^,]*}}%x, i{{[0-9]*}}* {{[^,]*}}%y)
+// CHECK: %.omp.reduction.red_list = alloca [2 x i8*]
+void mapWithReduction() {
+ int x, y;
+ #pragma omp target teams reduction(+:x) map(x,y) reduction(+:y)
+ ;
+}
+
+// HOST: define {{.*}}mapFrom
+// HOST: call {{.*}} @.[[OFFLOAD_FROM]].region_id{{.*}} @[[MAPTYPES_FROM]]
+//
+// CHECK: define {{.*}} void @[[OFFLOAD_FROM]](i{{[0-9]*}} %x)
+// CHECK: call void ({{.*}}@[[OUTLINE_FROM:.omp_outlined.[.0-9]*]]
+//
+// CHECK: define {{.*}} void @[[OUTLINE_FROM]]({{.*}} %.global_tid., {{.*}} %.bound_tid., i{{[0-9]*}} %x)
+void mapFrom() {
+ int x;
+ #pragma omp target teams firstprivate(x) map(from:x)
+ ;
+}
+
+// HOST: define {{.*}}mapTo
+// HOST: call {{.*}} @.[[OFFLOAD_TO]].region_id{{.*}} @[[MAPTYPES_TO]]
+//
+// CHECK: define {{.*}} void @[[OFFLOAD_TO]](i{{[0-9]*}} %x)
+// CHECK: call void ({{.*}}@[[OUTLINE_TO:.omp_outlined.[.0-9]*]]
+//
+// CHECK: define {{.*}} void @[[OUTLINE_TO]]({{.*}} %.global_tid., {{.*}} %.bound_tid., i{{[0-9]*}} %x)
+void mapTo() {
+ int x;
+ #pragma omp target teams firstprivate(x) map(to:x)
+ ;
+}
+
+// HOST: define {{.*}}mapAlloc
+// HOST: call {{.*}} @.[[OFFLOAD_ALLOC]].region_id{{.*}} @[[MAPTYPES_ALLOC]]
+//
+// CHECK: define {{.*}} void @[[OFFLOAD_ALLOC]](i{{[0-9]*}} %x)
+// CHECK: call void ({{.*}}@[[OUTLINE_ALLOC:.omp_outlined.[.0-9]*]]
+//
+// CHECK: define {{.*}} void @[[OUTLINE_ALLOC]]({{.*}} %.global_tid., {{.*}} %.bound_tid., i{{[0-9]*}} %x)
+void mapAlloc() {
+ int x;
+ #pragma omp target teams firstprivate(x) map(alloc:x)
+ ;
+}
+
+// HOST: define {{.*}}mapArray
+// HOST: call {{.*}} @.[[OFFLOAD_ARRAY_R0]].region_id{{.*}} @[[MAPTYPES_ARRAY_R0]]
+// HOST: call {{.*}} @.[[OFFLOAD_ARRAY_R1]].region_id{{.*}} @[[MAPTYPES_ARRAY_R1]]
+//
+// CHECK: define {{.*}} void @[[OFFLOAD_ARRAY_R0]]([88 x i{{[0-9]*}}]* {{[^,]*}}%y, [99 x i{{[0-9]*}}]* {{[^,]*}}%z)
+// CHECK: call void ({{.*}}@[[OUTLINE_ARRAY_R0:.omp_outlined.[.0-9]*]]
+//
+// CHECK: define {{.*}} void @[[OUTLINE_ARRAY_R0]]({{.*}} %.global_tid., {{.*}} %.bound_tid., [88 x i{{[0-9]*}}]* {{[^,]*}}%y, [99 x i{{[0-9]*}}]* {{[^,]*}}%z)
+// CHECK: %.omp.reduction.red_list = alloca [1 x i8*]
+//
+// CHECK: define {{.*}} void @[[OFFLOAD_ARRAY_R1]]([88 x i{{[0-9]*}}]* {{[^,]*}}%y, [99 x i{{[0-9]*}}]* {{[^,]*}}%z)
+// CHECK: call void ({{.*}}@[[OUTLINE_ARRAY_R1:.omp_outlined.[.0-9]*]]
+//
+// CHECK: define {{.*}} void @[[OUTLINE_ARRAY_R1]]({{.*}} %.global_tid., {{.*}} %.bound_tid., [88 x i{{[0-9]*}}]* {{[^,]*}}%y, [99 x i{{[0-9]*}}]* {{[^,]*}}%z)
+// CHECK: %.omp.reduction.red_list = alloca [1 x i8*]
+void mapArray() {
+ int x[77], y[88], z[99];
+ #pragma omp target teams private(x) firstprivate(y) reduction(+:z) map(x,y,z)
+ ;
+ #pragma omp target teams private(x) firstprivate(y) reduction(+:z) map(to:x,y,z)
+ ;
+}
+
+# if HAS_INT128
+// HOST-INT128: define {{.*}}mapInt128
+// HOST-INT128: call {{.*}} @.[[OFFLOAD_INT128_R0]].region_id{{.*}} @[[MAPTYPES_INT128_R0]]
+// HOST-INT128: call {{.*}} @.[[OFFLOAD_INT128_R1]].region_id{{.*}} @[[MAPTYPES_INT128_R1]]
+//
+// INT128: define {{.*}} void @[[OFFLOAD_INT128_R0]](i128* {{[^,]*}}%y, i128* {{[^,]*}}%z)
+// INT128: call void ({{.*}}@[[OUTLINE_INT128_R0:.omp_outlined.[.0-9]*]]
+//
+// INT128: define {{.*}} void @[[OUTLINE_INT128_R0]]({{.*}} %.global_tid., {{.*}} %.bound_tid., i128* {{[^,]*}}%y, i128* {{[^,]*}}%z)
+// INT128: %.omp.reduction.red_list = alloca [1 x i8*]
+//
+// INT128: define {{.*}} void @[[OFFLOAD_INT128_R1]](i128* {{[^,]*}}%y, i128* {{[^,]*}}%z)
+// INT128: call void ({{.*}}@[[OUTLINE_INT128_R1:.omp_outlined.[.0-9]*]]
+//
+// INT128: define {{.*}} void @[[OUTLINE_INT128_R1]]({{.*}} %.global_tid., {{.*}} %.bound_tid., i128* {{[^,]*}}%y, i128* {{[^,]*}}%z)
+// INT128: %.omp.reduction.red_list = alloca [1 x i8*]
+void mapInt128() {
+ __int128 x, y, z;
+ #pragma omp target teams private(x) firstprivate(y) reduction(+:z) map(x,y,z)
+ ;
+ #pragma omp target teams private(x) firstprivate(y) reduction(+:z) map(from:x,y,z)
+ ;
+}
+# endif
+#endif
Index: clang/test/OpenMP/target_teams_distribute_simd_private_messages.cpp
===================================================================
--- clang/test/OpenMP/target_teams_distribute_simd_private_messages.cpp
+++ clang/test/OpenMP/target_teams_distribute_simd_private_messages.cpp
@@ -1,6 +1,9 @@
-// RUN: %clang_cc1 -verify -fopenmp %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=40 -fopenmp %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=45 -fopenmp %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected -fopenmp-version=50 -fopenmp %s -Wuninitialized
-// RUN: %clang_cc1 -verify -fopenmp-simd %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd %s -Wuninitialized
typedef void **omp_allocator_handle_t;
extern const omp_allocator_handle_t omp_default_mem_alloc;
@@ -132,7 +135,7 @@
for (int k = 0; k < 10; ++k) {
}
-#pragma omp target teams distribute simd private(j) map(j) // expected-error {{private variable cannot be in a map clause in '#pragma omp target teams distribute simd' directive}} expected-note {{defined as private}}
+#pragma omp target teams distribute simd private(j) map(j) // le45-error {{private variable cannot be in a map clause in '#pragma omp target teams distribute simd' directive}} le45-note {{defined as private}}
for (int k = 0; k < argc; ++k) ++k;
return 0;
Index: clang/test/OpenMP/target_teams_distribute_simd_lastprivate_messages.cpp
===================================================================
--- clang/test/OpenMP/target_teams_distribute_simd_lastprivate_messages.cpp
+++ clang/test/OpenMP/target_teams_distribute_simd_lastprivate_messages.cpp
@@ -1,6 +1,9 @@
-// RUN: %clang_cc1 -verify -fopenmp %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=40 -fopenmp %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=45 -fopenmp %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected -fopenmp-version=50 -fopenmp %s -Wuninitialized
-// RUN: %clang_cc1 -verify -fopenmp-simd %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd %s -Wuninitialized
typedef void **omp_allocator_handle_t;
extern const omp_allocator_handle_t omp_default_mem_alloc;
@@ -240,7 +243,7 @@
#pragma omp target teams distribute simd lastprivate(si) // OK
for (i = 0; i < argc; ++i) si = i + 1;
-#pragma omp target teams distribute simd lastprivate(k) map(k) // expected-error {{lastprivate variable cannot be in a map clause in '#pragma omp target teams distribute simd' directive}} expected-note {{defined as lastprivate}}
+#pragma omp target teams distribute simd lastprivate(k) map(k) // le45-error {{lastprivate variable cannot be in a map clause in '#pragma omp target teams distribute simd' directive}} le45-note {{defined as lastprivate}}
for (i = 0; i < argc; ++i) foo();
return foomain<S4, S5>(argc, argv); // expected-note {{in instantiation of function template specialization 'foomain<S4, S5>' requested here}}
Index: clang/test/OpenMP/target_teams_distribute_simd_firstprivate_messages.cpp
===================================================================
--- clang/test/OpenMP/target_teams_distribute_simd_firstprivate_messages.cpp
+++ clang/test/OpenMP/target_teams_distribute_simd_firstprivate_messages.cpp
@@ -1,6 +1,9 @@
-// RUN: %clang_cc1 -verify -fopenmp %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=40 -fopenmp %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=45 -fopenmp %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected -fopenmp-version=50 -fopenmp %s -Wuninitialized
-// RUN: %clang_cc1 -verify -fopenmp-simd %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd %s -Wuninitialized
typedef void **omp_allocator_handle_t;
extern const omp_allocator_handle_t omp_default_mem_alloc;
@@ -131,8 +134,10 @@
#pragma omp target teams distribute simd firstprivate(h) // expected-error {{threadprivate or thread local variable cannot be firstprivate}}
for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute simd private(i), firstprivate(i) // expected-error {{private variable cannot be firstprivate}} expected-note 2 {{defined as private}}
- for (i = 0; i < argc; ++i) foo(); // expected-error {{loop iteration variable in the associated loop of 'omp target teams distribute simd' directive may not be private, predetermined as linear}}
+#pragma omp target teams distribute simd private(i), firstprivate(i) // expected-error {{private variable cannot be firstprivate}} expected-note {{defined as private}}
+ // le45-note@-1 {{defined as private}}
+ // FIXME: Error is dropped for OpenMP 5.0. Probably shouldn't be.
+ for (i = 0; i < argc; ++i) foo(); // le45-error {{loop iteration variable in the associated loop of 'omp target teams distribute simd' directive may not be private, predetermined as linear}}
#pragma omp target teams distribute simd firstprivate(i)
for (j = 0; j < argc; ++j) foo();
@@ -147,7 +152,7 @@
#pragma omp target teams distribute simd lastprivate(argc), firstprivate(argc)
for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute simd firstprivate(argc) map(argc) // expected-error {{firstprivate variable cannot be in a map clause in '#pragma omp target teams distribute simd' directive}} expected-note {{defined as firstprivate}}
+#pragma omp target teams distribute simd firstprivate(argc) map(argc) // le45-error {{firstprivate variable cannot be in a map clause in '#pragma omp target teams distribute simd' directive}} le45-note {{defined as firstprivate}}
for (i = 0; i < argc; ++i) foo();
return 0;
Index: clang/test/OpenMP/target_teams_distribute_parallel_for_simd_lastprivate_messages.cpp
===================================================================
--- clang/test/OpenMP/target_teams_distribute_parallel_for_simd_lastprivate_messages.cpp
+++ clang/test/OpenMP/target_teams_distribute_parallel_for_simd_lastprivate_messages.cpp
@@ -1,6 +1,9 @@
-// RUN: %clang_cc1 -verify -fopenmp %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=40 -fopenmp %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=45 -fopenmp %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected -fopenmp-version=50 -fopenmp %s -Wuninitialized
-// RUN: %clang_cc1 -verify -fopenmp-simd %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd %s -Wuninitialized
typedef void **omp_allocator_handle_t;
extern const omp_allocator_handle_t omp_default_mem_alloc;
@@ -240,7 +243,7 @@
#pragma omp target teams distribute parallel for simd lastprivate(si) // OK
for (i = 0; i < argc; ++i) si = i + 1;
-#pragma omp target teams distribute parallel for simd lastprivate(k) map(k) // expected-error {{lastprivate variable cannot be in a map clause in '#pragma omp target teams distribute parallel for simd' directive}} expected-note {{defined as lastprivate}}
+#pragma omp target teams distribute parallel for simd lastprivate(k) map(k) // le45-error {{lastprivate variable cannot be in a map clause in '#pragma omp target teams distribute parallel for simd' directive}} le45-note {{defined as lastprivate}}
for (i = 0; i < argc; ++i) foo();
return foomain<S4, S5>(argc, argv); // expected-note {{in instantiation of function template specialization 'foomain<S4, S5>' requested here}}
}
Index: clang/test/OpenMP/target_teams_distribute_parallel_for_lastprivate_messages.cpp
===================================================================
--- clang/test/OpenMP/target_teams_distribute_parallel_for_lastprivate_messages.cpp
+++ clang/test/OpenMP/target_teams_distribute_parallel_for_lastprivate_messages.cpp
@@ -1,6 +1,9 @@
-// RUN: %clang_cc1 -verify -fopenmp %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=40 -fopenmp %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=45 -fopenmp %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected -fopenmp-version=50 -fopenmp %s -Wuninitialized
-// RUN: %clang_cc1 -verify -fopenmp-simd %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd %s -Wuninitialized
typedef void **omp_allocator_handle_t;
extern const omp_allocator_handle_t omp_default_mem_alloc;
@@ -240,7 +243,7 @@
#pragma omp target teams distribute parallel for lastprivate(si) // OK
for (i = 0; i < argc; ++i) si = i + 1;
-#pragma omp target teams distribute parallel for lastprivate(k) map(k) // expected-error {{lastprivate variable cannot be in a map clause in '#pragma omp target teams distribute parallel for' directive}} expected-note {{defined as lastprivate}}
+#pragma omp target teams distribute parallel for lastprivate(k) map(k) // le45-error {{lastprivate variable cannot be in a map clause in '#pragma omp target teams distribute parallel for' directive}} le45-note {{defined as lastprivate}}
for (i = 0; i < argc; ++i) foo();
return foomain<S4, S5>(argc, argv); // expected-note {{in instantiation of function template specialization 'foomain<S4, S5>' requested here}}
Index: clang/lib/Sema/SemaOpenMP.cpp
===================================================================
--- clang/lib/Sema/SemaOpenMP.cpp
+++ clang/lib/Sema/SemaOpenMP.cpp
@@ -10891,7 +10891,13 @@
// OpenMP 4.5 [2.15.5.1, Restrictions, p.3]
// A list item cannot appear in both a map clause and a data-sharing
// attribute clause on the same construct
- if (isOpenMPTargetExecutionDirective(CurrDir)) {
+ //
+ // OpenMP 5.0 [2.19.7.1, Restrictions, p.7]
+ // A list item cannot appear in both a map clause and a data-sharing
+ // attribute clause on the same construct unless the construct is a
+ // combined construct.
+ if ((LangOpts.OpenMP <= 45 && isOpenMPTargetExecutionDirective(CurrDir)) ||
+ CurrDir == OMPD_target) {
OpenMPClauseKind ConflictKind;
if (DSAStack->checkMappableExprComponentListsForDecl(
VD, /*CurrentRegionOnly=*/true,
@@ -11126,7 +11132,14 @@
// OpenMP 4.5 [2.15.5.1, Restrictions, p.3]
// A list item cannot appear in both a map clause and a data-sharing
// attribute clause on the same construct
- if (isOpenMPTargetExecutionDirective(CurrDir)) {
+ //
+ // OpenMP 5.0 [2.19.7.1, Restrictions, p.7]
+ // A list item cannot appear in both a map clause and a data-sharing
+ // attribute clause on the same construct unless the construct is a
+ // combined construct.
+ if ((LangOpts.OpenMP <= 45 &&
+ isOpenMPTargetExecutionDirective(CurrDir)) ||
+ CurrDir == OMPD_target) {
OpenMPClauseKind ConflictKind;
if (DSAStack->checkMappableExprComponentListsForDecl(
VD, /*CurrentRegionOnly=*/true,
@@ -14237,7 +14250,14 @@
// OpenMP 4.5 [2.15.5.1, Restrictions, p.3]
// A list item cannot appear in both a map clause and a data-sharing
// attribute clause on the same construct
- if (VD && isOpenMPTargetExecutionDirective(DKind)) {
+ //
+ // OpenMP 5.0 [2.19.7.1, Restrictions, p.7]
+ // A list item cannot appear in both a map clause and a data-sharing
+ // attribute clause on the same construct unless the construct is a
+ // combined construct.
+ if (VD && ((SemaRef.LangOpts.OpenMP <= 45 &&
+ isOpenMPTargetExecutionDirective(DKind)) ||
+ DKind == OMPD_target)) {
DSAStackTy::DSAVarData DVar = DSAS->getTopDSA(VD, /*FromParent=*/false);
if (isOpenMPPrivate(DVar.CKind)) {
SemaRef.Diag(ELoc, diag::err_omp_variable_in_given_clause_and_dsa)
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
