================
@@ -0,0 +1,3823 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
UTC_ARGS: --check-globals all --filter-out-after "getelem.*kernel" --filter-out
"= alloca.*" --include-generated-funcs --replace-value-regex
"__omp_offloading_[0-9a-z]+_[0-9a-z]+" "reduction_size[.].+[.]"
"pl_cond[.].+[.|,]" --prefix-filecheck-ir-name _ --global-value-regex
"\.offload_.*" --global-hex-value-regex ".offload_maptypes.*" --version 6
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=51 -std=c++20 -triple
x86_64-unknown-unknown -emit-llvm %s -o - | FileCheck %s
+
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -std=c++20 -triple
x86_64-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -std=c++20 -triple
x86_64-unknown-unknown -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
+
+// expected-no-diagnostics
+
+#ifndef HEADER
+#define HEADER
+
+struct Point { int x, y; };
+struct Point3D { int x, y, z; };
+
+void test_target_explicit_map() {
+ Point p{1, 2};
+ auto [a, b] = p;
+
+#pragma omp target map(tofrom: p)
+ {
+ a = a + 1;
+ b = b + 1;
+ }
+}
+
+void test_target_implicit_map() {
+ Point p{1, 2};
+ auto [a, b] = p;
+
+#pragma omp target
+ {
+ int sum = a + b;
+ }
+}
+
+void test_target_parallel() {
+ Point p{3, 4};
+ auto [a, b] = p;
+
+#pragma omp target parallel
+ {
+ int sum = a + b;
+ }
+}
+
+void test_target_parallel_for() {
+ Point p{5, 6};
+ auto [a, b] = p;
+
+#pragma omp target parallel for
+ for (int i = 0; i < 10; i++) {
+ int result = a + b + i;
+ }
+}
+
+void test_firstprivate_dsa() {
+ Point p{7, 8};
+ auto [a, b] = p;
+
+#pragma omp parallel firstprivate(p)
+ {
+ int sum = a + b;
+ }
+}
+
+void test_shared_dsa() {
+ Point p{9, 10};
+ auto [a, b] = p;
+
+#pragma omp parallel shared(p)
+ {
+ int sum = a + b;
+ }
+}
+
+void test_reduction_dsa() {
+ Point p{1, 2};
+ auto [a, b] = p;
+ int sum = 0;
+
+#pragma omp parallel for reduction(+:sum)
+ for (int i = 0; i < 10; i++) {
+ sum += a + b;
+ }
+}
+
+void test_parallel_for_() {
+ Point p{11, 12};
+ auto [a, b] = p;
+
+#pragma omp parallel for
+ for (int i = 0; i < 10; i++) {
+ int result = a + b + i;
+ }
+}
+
+void test_parallel_for_simd_() {
+ Point p{13, 14};
+ auto [a, b] = p;
+
+#pragma omp parallel for simd
+ for (int i = 0; i < 10; i++) {
+ int result = a + b + i;
+ }
+}
+
+void test_target_teams_distribute() {
+ Point p{15, 16};
+ auto [a, b] = p;
+
+#pragma omp target teams distribute
+ for (int i = 0; i < 10; i++) {
+ int result = a + b + i;
+ }
+}
+
+void test_task() {
+ Point p{17, 18};
+ auto [a, b] = p;
+
+#pragma omp task
+ {
+ int sum = a + b;
+ }
+}
+
+void test_task_depend() {
+ Point p{19, 20};
+ auto [a, b] = p;
+
+#pragma omp task depend(in: a, b)
+ {
+ int sum = a + b;
+ }
+}
+
+void test_taskloop_() {
+ Point p{21, 22};
+ auto [a, b] = p;
+
+#pragma omp taskloop
+ for (int i = 0; i < 10; i++) {
+ int result = a + b + i;
+ }
+}
+
+template<typename T>
+int test_template_bas(T p) {
+ auto [a, b] = p;
+ int result = 0;
+
+#pragma omp parallel reduction(+:result)
+ {
+ result = a + b;
+ }
+ return result;
+}
+
+template<typename T>
+int test_template_target(T p) {
+ auto [a, b] = p;
+ int result = 0;
+
+#pragma omp target map(tofrom: result)
+ {
+ result = a + b;
+ }
+ return result;
+}
+
+template<typename T>
+int test_template_task(T p) {
+ auto [a, b] = p;
+ int result = 0;
+
+#pragma omp task shared(result)
+ {
+ result = a + b;
+ }
+#pragma omp taskwait
+ return result;
+}
+
+template<typename T>
+int test_template_3_bindings(T p) {
+ auto [x, y, z] = p;
+ int result = 0;
+
+#pragma omp parallel reduction(+:result)
+ {
+ result = x + y + z;
+ }
+ return result;
+}
+
+void instantiate_templates() {
+ Point p2{1, 2};
+ Point3D p3{1, 2, 3};
+
+ test_template_bas(p2);
+ test_template_bas(Point{3, 4});
+ test_template_target(p2);
+ test_template_task(p2);
+ test_template_3_bindings(p3);
+}
+
+void test_static_binding() {
+ static Point p{23, 24};
+ static auto [a, b] = p;
+
+#pragma omp parallel
+ {
+ int sum = a + b;
+ }
+}
+
+void test_static_binding_shared() {
+ static Point p{25, 26};
+ static auto [a, b] = p;
+
+#pragma omp parallel shared(p)
+ {
+ a = a + 1;
+ b = b + 1;
+ }
+}
+
+void test_array_target() {
+ int arr[2] = {27, 28};
+ auto [a, b] = arr;
+
+#pragma omp target
+ {
+ int sum = a + b;
+ }
+}
+
+void test_array_task() {
+ int arr[2] = {29, 30};
+ auto [a, b] = arr;
+
+#pragma omp task
+ {
+ int sum = a + b;
+ }
+}
+
+void test_nested() {
+ Point p{31, 32};
+ auto [a, b] = p;
+
+#pragma omp parallel
+ {
+#pragma omp critical
+ {
+ int sum = a + b;
+ }
+
+#pragma omp task
+ {
+ int product = a * b;
+ }
+ }
+}
+
+void test_reference_binding() {
+ Point p{31, 32};
+ auto& [a, b] = p;
+
+#pragma omp parallel
+ {
+ int sum = a + b;
+ }
+}
+
+void test_const_binding() {
+ const Point p{33, 34};
+ const auto [a, b] = p;
+
+#pragma omp parallel
+ {
+ int sum = a + b;
+ }
+}
+
+void test_multiple_bindings() {
+ Point p1{33, 34};
+ Point p2{35, 36};
+ auto [a, b] = p1;
+ auto [c, d] = p2;
+
+#pragma omp parallel
+ {
+ int sum = a + b + c + d;
+ }
+}
+
+void test_multiple_bindings_mixed_dsa() {
+ Point p1{37, 38};
+ Point p2{39, 40};
+ auto [a, b] = p1;
+ auto [c, d] = p2;
+
+#pragma omp parallel firstprivate(p1) shared(p2)
+ {
+ int result = a + b + c + d;
+ }
+}
+
+void test_array_3_elements() {
+ int arr[3] = {35, 36, 37};
+ auto [a, b, c] = arr;
+
+#pragma omp parallel
+ {
+ int sum = a + b + c;
+ }
+}
+
+void test_single() {
+ Point p{38, 39};
+ auto [a, b] = p;
+
+#pragma omp parallel
+#pragma omp single
+ {
+ int sum = a + b;
+ }
+}
+
+void test_sections() {
+ Point p{40, 41};
+ auto [a, b] = p;
+
+#pragma omp parallel sections
+ {
+#pragma omp section
+ { int sum = a + b; }
+#pragma omp section
+ { int diff = a - b; }
+ }
+}
+
+void test_nested_parallel() {
+ Point p{42, 43};
+ auto [a, b] = p;
+
+#pragma omp parallel
+ {
+#pragma omp parallel
+ {
+ int sum = a + b;
+ }
+ }
+}
+
+void test_simd_() {
+ Point p{44, 45};
+ auto [a, b] = p;
+
+#pragma omp simd
+ for (int i = 0; i < 10; i++) {
+ int result = a + b + i;
+ }
+}
+
+void test_private_individual_bindings() {
+ Point p{1, 2};
+ auto [a, b] = p;
+#pragma omp parallel private(a)
+ {
+ a = 2;
+ }
+}
+
+void test_firstprivate_individual_bindings() {
+ Point p{1, 2};
+ auto [a, b] = p;
+#pragma omp parallel firstprivate(b)
+ {
+ b = b + 10;
+ }
+}
+
+void test_static_bindings() {
+ static auto [a, b] = Point{1, 2};
+#pragma omp parallel
+ {
+ (void)(a + b);
+ }
+}
+
+void use(int a);
+
+void test_shadowing() {
+ auto [a, b] = Point{1, 2};
+#pragma omp parallel private(a)
+ {
+ use(a);
+ {
+ auto [a, b] = Point{10, 20};
+ use(a);
+ }
+ use(a);
+ }
+ use(a);
+}
+
+void test_simd_private_then_parallel() {
+ Point p{1,2};
+ auto [a,b] = p;
+#pragma omp simd private(a)
+ for (int i=0;i<10;++i) a += i;
+#pragma omp parallel
+ {
+ use(a);
+ }
+}
+
+void test_linear_binding() {
+ Point p{1, 2};
+ auto [a, b] = p;
+
+#pragma omp simd linear(a:1)
+ for (int i = 0; i < 10; ++i) {
+ a += 1;
+ use(a);
+ }
+}
+
+void test_reduction_binding_sum() {
+ Point p{0, 0};
+ auto [a, b] = p;
+
+#pragma omp parallel for reduction(+:a)
+ for (int i = 0; i < 100; ++i) {
+ a += i;
+ }
+ use(a);
+}
+
+void test_reduction_binding_operators() {
+ Point p{1, 100};
+ auto [a, b] = p;
+
+#pragma omp parallel for reduction(*:a) reduction(min:b)
+ for (int i = 1; i <= 10; ++i) {
+ a *= 2;
+ if (i < b) b = i;
+ }
+ use(a);
+ use(b);
+}
+
+void test_reduction_binding_nontrivial() {
+ struct NonTrivial {
+ int value;
+ NonTrivial() : value(0) {}
+ NonTrivial(int v) : value(v) {}
+ ~NonTrivial() {}
+ NonTrivial& operator+=(int x) { value += x; return *this; }
+ NonTrivial& operator+=(const NonTrivial& other) { value += other.value;
return *this; }
+ };
+ struct PairNonTrivial {
+ NonTrivial a;
+ NonTrivial b;
+ };
+
+ PairNonTrivial p{NonTrivial(0), NonTrivial(0)};
+ auto [a, b] = p;
+
+#pragma omp parallel for reduction(+:a)
+ for (int i = 0; i < 10; ++i) {
+ a += i;
+ }
+ use(a.value);
+}
+
+void test_lastprivate_binding() {
+ Point p{1, 2};
+ auto [a, b] = p;
+
+#pragma omp parallel for lastprivate(a)
+ for (int i = 0; i < 10; ++i) {
+ a = i * 10;
+ }
+ use(a);
+}
+
+void test_reduction_binding_max() {
+ Point p{-100, -100};
+ auto [a, b] = p;
+
+#pragma omp parallel for reduction(max:a,b)
+ for (int i = 0; i < 100; ++i) {
+ if (i > a) a = i;
+ if (i > b) b = i;
+ }
+ use(a);
+ use(b);
+}
+
+void test_binding_name_collision() {
+ Point p1{1, 2};
+ auto [a, b] = p1;
+
+#pragma omp parallel reduction(+:a)
+ {
+ Point p2{10, 20};
+ auto [a, b] = p2;
+ a += 1;
+ use(a);
+ }
+}
+
+struct NonTrivialCopy {
+ int value;
+ int copy_count;
+
+ NonTrivialCopy(int v) : value(v), copy_count(0) {}
+ NonTrivialCopy() : value(0), copy_count(0) {}
+ NonTrivialCopy(const NonTrivialCopy &other)
+ : value(other.value), copy_count(other.copy_count + 1) {}
+
+ ~NonTrivialCopy() {}
+};
+
+struct Pair {
+ NonTrivialCopy x;
+ NonTrivialCopy y;
+};
+
+void test_firstprivate_nontrivial_sb() {
+ Pair p{NonTrivialCopy(10), NonTrivialCopy(20)};
+ auto [a, b] = p;
+
+#pragma omp parallel firstprivate(a)
+ {
+ (void)a.value;
+ }
+}
+
+void test_firstprivate_ref_binding_sb() {
+ Pair p{NonTrivialCopy(10), NonTrivialCopy(20)};
+ auto& [a, b] = p;
+
+#pragma omp parallel firstprivate(a)
+ {
+ (void)a.value;
+ }
+}
+
+void test_firstprivate_const_ref_binding_sb() {
+ Pair p{NonTrivialCopy(10), NonTrivialCopy(20)};
+ const auto& [a, b] = p;
+
+#pragma omp parallel firstprivate(a)
+ {
+ (void)a.value;
+ }
+}
+
+void test_firstprivate_multiple_bindings_sb() {
+ Pair p{NonTrivialCopy(10), NonTrivialCopy(20)};
+ auto [a, b] = p;
+
+#pragma omp parallel firstprivate(a, b)
+ {
+ (void)(a.value + b.value);
+ }
+}
+
+struct WithDtor {
+ int *ptr;
+ WithDtor(int v) : ptr(new int(v)) {}
+ WithDtor(const WithDtor &other) : ptr(new int(*other.ptr)) {}
+ ~WithDtor() { delete ptr; }
+};
+
+struct PairWithDtor {
+ WithDtor x;
+ WithDtor y;
+};
+
+void test_firstprivate_with_destructor_sb() {
+ PairWithDtor p{WithDtor(100), WithDtor(200)};
+ auto [a, b] = p;
+
+#pragma omp parallel firstprivate(a)
+ {
+ (void)(*a.ptr);
+ }
+}
+
+void test_firstprivate_array_bindings_sb() {
+ NonTrivialCopy arr[3] = {NonTrivialCopy(1), NonTrivialCopy(2),
NonTrivialCopy(3)};
+ auto [a, b, c] = arr;
+
+
+#pragma omp parallel firstprivate(a, b, c)
+ {
+ (void)(a.value + b.value + c.value);
+ }
+}
+
+struct Inner {
+ int val;
+ Inner(int v) : val(v) {}
+ Inner(const Inner &o) : val(o.val * 2) {}
+ ~Inner() {}
+};
+
+struct Outer {
+ Inner i1;
+ Inner i2;
+};
+
+void test_firstprivate_nested_struct_sb() {
+ Outer o{Inner(5), Inner(10)};
+ auto [x, y] = o;
+
+#pragma omp parallel firstprivate(x)
+ {
+ (void)x.val;
+ }
+}
+
+void test_firstprivate_ref_binding_both_bindings_sb() {
+ Pair p{NonTrivialCopy(10), NonTrivialCopy(20)};
+ auto& [a, b] = p;
+
+#pragma omp parallel firstprivate(a, b)
+ {
+ (void)(a.value + b.value);
+ }
+}
+
+
+void test_firstprivate_const_ref_binding_both_bindings_sb() {
+ Pair p{NonTrivialCopy(10), NonTrivialCopy(20)};
+ const auto& [a, b] = p;
+
+#pragma omp parallel firstprivate(a, b)
+ {
+ (void)(a.value + b.value);
+ }
+}
+
+void test_lambda_capture_binding_by_value() {
+ Point p{10, 20};
+ auto [a, b] = p;
+ auto lambda = [a]() {
+ return a + a;
+ };
+ int result = lambda();
+}
+
+
+void test_lambda_capture_binding_by_ref() {
+ Point p{10, 20};
+ auto [a, b] = p;
+ auto lambda = [&a]() {
+ a = 100;
+ };
+ lambda();
+}
+
+
+void test_lambda_capture_multiple_bindings() {
+ Point p{10, 20};
+ auto [a, b] = p;
+ auto lambda = [a, &b]() {
+ b = a + 50;
+ };
+ lambda();
+}
+
+
+void test_lambda_implicit_capture() {
+ Point p{10, 20};
+ auto [a, b] = p;
+ auto lambda = [=]() {
+ return a + b;
+ };
+ lambda();
+}
+#endif
+// CHECK-LABEL: define dso_local void @_Z24test_target_explicit_mapv(
+// CHECK-SAME: ) #[[ATTR0:[0-9]+]] {
+// CHECK: [[ENTRY:.*:]]
+// CHECK: call void @llvm.memcpy.p0.p0.i64(ptr align 4 [[P:%.*]], ptr align
4 @__const._Z24test_target_explicit_mapv.p, i64 8, i1 false)
+// CHECK: call void @llvm.memcpy.p0.p0.i64(ptr align 4 [[TMP0:%.*]], ptr
align 4 [[P]], i64 8, i1 false)
+// CHECK: [[TMP1:%.*]] = load [[STRUCT_POINT:%.*]], ptr [[TMP0]], align 4
+// CHECK: store [[STRUCT_POINT]] [[TMP1]], ptr [[DOTCASTED:%.*]], align 4
+// CHECK: [[TMP2:%.*]] = load i64, ptr [[DOTCASTED]], align 8
+// CHECK: call void
@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z24test_target_explicit_mapv_l19(i64
[[TMP2]], ptr null) #[[ATTR3:[0-9]+]]
+// CHECK: ret void
+//
+//
+// CHECK-LABEL: define internal void
@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z24test_target_explicit_mapv_l19(
+// CHECK-SAME: i64 noundef [[TMP0:%.*]], ptr noalias noundef [[DYN_PTR:%.*]])
#[[ATTR2:[0-9]+]] {
+// CHECK: [[ENTRY:.*:]]
+// CHECK: store i64 [[TMP0]], ptr [[DOTADDR:%.*]], align 8
+// CHECK: store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR:%.*]], align 8
+// CHECK: [[X:%.*]] = getelementptr inbounds nuw [[STRUCT_POINT:%.*]], ptr
[[DOTADDR]], i32 0, i32 0
+// CHECK: [[TMP1:%.*]] = load i32, ptr [[X]], align 4
+// CHECK: [[ADD:%.*]] = add nsw i32 [[TMP1]], 1
+// CHECK: [[X1:%.*]] = getelementptr inbounds nuw [[STRUCT_POINT]], ptr
[[DOTADDR]], i32 0, i32 0
+// CHECK: store i32 [[ADD]], ptr [[X1]], align 4
+// CHECK: [[Y:%.*]] = getelementptr inbounds nuw [[STRUCT_POINT]], ptr
[[DOTADDR]], i32 0, i32 1
+// CHECK: [[TMP2:%.*]] = load i32, ptr [[Y]], align 4
+// CHECK: [[ADD2:%.*]] = add nsw i32 [[TMP2]], 1
+// CHECK: [[Y3:%.*]] = getelementptr inbounds nuw [[STRUCT_POINT]], ptr
[[DOTADDR]], i32 0, i32 1
+// CHECK: store i32 [[ADD2]], ptr [[Y3]], align 4
+// CHECK: ret void
+//
+//
+// CHECK-LABEL: define dso_local void @_Z24test_target_implicit_mapv(
+// CHECK-SAME: ) #[[ATTR0]] {
+// CHECK: [[ENTRY:.*:]]
+// CHECK: call void @llvm.memcpy.p0.p0.i64(ptr align 4 [[P:%.*]], ptr align
4 @__const._Z24test_target_implicit_mapv.p, i64 8, i1 false)
+// CHECK: call void @llvm.memcpy.p0.p0.i64(ptr align 4 [[TMP0:%.*]], ptr
align 4 [[P]], i64 8, i1 false)
+// CHECK: ret void
+//
+//
+// CHECK-LABEL: define dso_local void @_Z20test_target_parallelv(
+// CHECK-SAME: ) #[[ATTR0]] {
+// CHECK: [[ENTRY:.*:]]
+// CHECK: call void @llvm.memcpy.p0.p0.i64(ptr align 4 [[P:%.*]], ptr align
4 @__const._Z20test_target_parallelv.p, i64 8, i1 false)
+// CHECK: call void @llvm.memcpy.p0.p0.i64(ptr align 4 [[TMP0:%.*]], ptr
align 4 [[P]], i64 8, i1 false)
+// CHECK: ret void
+//
+//
+// CHECK-LABEL: define dso_local void @_Z24test_target_parallel_forv(
+// CHECK-SAME: ) #[[ATTR0]] {
+// CHECK: [[ENTRY:.*:]]
+// CHECK: call void @llvm.memcpy.p0.p0.i64(ptr align 4 [[P:%.*]], ptr align
4 @__const._Z24test_target_parallel_forv.p, i64 8, i1 false)
+// CHECK: call void @llvm.memcpy.p0.p0.i64(ptr align 4 [[TMP0:%.*]], ptr
align 4 [[P]], i64 8, i1 false)
+// CHECK: ret void
----------------
alexey-bataev wrote:
Same question, where is the target kernel?
https://github.com/llvm/llvm-project/pull/190832
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
