2020-03-31 04:06:01 +08:00
|
|
|
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=50 -ast-print %s | FileCheck %s
|
|
|
|
|
// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s
|
|
|
|
|
// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
|
2017-12-30 02:07:07 +08:00
|
|
|
|
2020-03-31 04:06:01 +08:00
|
|
|
// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=50 -ast-print %s | FileCheck %s
|
|
|
|
|
// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s
|
|
|
|
|
// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
|
2020-07-30 00:18:45 +08:00
|
|
|
|
|
|
|
|
// RUN: %clang_cc1 -DOMP51 -verify -fopenmp -fopenmp-version=51 -ast-print %s | FileCheck -check-prefixes=CHECK,OMP51 %s
|
|
|
|
|
// RUN: %clang_cc1 -DOMP51 -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -emit-pch -o %t %s
|
|
|
|
|
// RUN: %clang_cc1 -DOMP51 -fopenmp -fopenmp-version=51 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck -check-prefixes=CHECK,OMP51 %s
|
|
|
|
|
|
|
|
|
|
// RUN: %clang_cc1 -DOMP51 -verify -fopenmp-simd -fopenmp-version=51 -ast-print %s | FileCheck -check-prefixes=CHECK,OMP51 %s
|
|
|
|
|
// RUN: %clang_cc1 -DOMP51 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -emit-pch -o %t %s
|
|
|
|
|
// RUN: %clang_cc1 -DOMP51 -fopenmp-simd -fopenmp-version=51 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck -check-prefixes=CHECK,OMP51 %s
|
2016-05-27 01:39:58 +08:00
|
|
|
// expected-no-diagnostics
|
|
|
|
|
|
|
|
|
|
#ifndef HEADER
|
|
|
|
|
#define HEADER
|
|
|
|
|
|
|
|
|
|
void foo() {}
|
|
|
|
|
|
|
|
|
|
template <class T, class U>
|
|
|
|
|
T foo(T targ, U uarg) {
|
2020-03-31 04:06:01 +08:00
|
|
|
static T a, *p;
|
2016-05-27 01:39:58 +08:00
|
|
|
U b;
|
|
|
|
|
int l;
|
2020-03-31 04:06:01 +08:00
|
|
|
#pragma omp target update to(([a][targ])p, a) if(l>5) device(l) nowait depend(inout:l)
|
2016-05-27 01:49:04 +08:00
|
|
|
|
2020-03-31 04:06:01 +08:00
|
|
|
#pragma omp target update from(b, ([a][targ])p) if(l<5) device(l-1) nowait depend(inout:l)
|
2020-07-10 02:27:32 +08:00
|
|
|
|
[OpenMP5.0] map item can be non-contiguous for target update
In order not to modify the `tgt_target_data_update` information but still be
able to pass the extra information for non-contiguous map item (offset,
count, and stride for each dimension), this patch overload `arg` when
the maptype is set as `OMP_MAP_DESCRIPTOR`. The origin `arg` is for
passing the pointer information, however, the overloaded `arg` is an
array of descriptor_dim:
struct descriptor_dim {
int64_t offset;
int64_t count;
int64_t stride
};
and the array size is the same as dimension size. In addition, since we
have count and stride information in descriptor_dim, we can replace/overload the
`arg_size` parameter by using dimension size.
For supporting `stride` in array section, we use a dummy dimension in
descriptor to store the unit size. The formula for counting the stride
in dimension D_n: `unit size * (D_0 * D_1 ... * D_n-1) * D_n.stride`.
Demonstrate how it works:
```
double arr[3][4][5];
D0: { offset = 0, count = 1, stride = 8 } // offset, count, dimension size always be 0, 1, 1 for this extra dimension, stride is the unit size
D1: { offset = 0, count = 2, stride = 8 * 1 * 2 = 16 } // stride = unit size * (product of dimension size of D0) * D1.stride = 4 * 1 * 2 = 8
D2: { offset = 2, count = 2, stride = 8 * (1 * 5) * 1 = 40 } // stride = unit size * (product of dimension size of D0, D1) * D2.stride = 4 * 5 * 1 = 20
D3: { offset = 0, count = 2, stride = 8 * (1 * 5 * 4) * 2 = 320 } // stride = unit size * (product of dimension size of D0, D1, D2) * D3.stride = 4 * 25 * 2 = 200
// X here means we need to offload this data, therefore, runtime will transfer
// data from offset 80, 96, 120, 136, 400, 416, 440, 456
// Runtime patch: https://reviews.llvm.org/D82245
// OOOOO OOOOO OOOOO
// OOOOO OOOOO OOOOO
// XOXOO OOOOO XOXOO
// XOXOO OOOOO XOXOO
```
Reviewed By: ABataev
Differential Revision: https://reviews.llvm.org/D84192
2020-11-07 11:03:55 +08:00
|
|
|
U marr[10][10][10];
|
|
|
|
|
#pragma omp target update to(marr[2][0:2][0:2])
|
|
|
|
|
|
|
|
|
|
#pragma omp target update from(marr[2][0:2][0:2])
|
|
|
|
|
|
|
|
|
|
#pragma omp target update from(marr[:2][0:2][0:2:1])
|
|
|
|
|
|
|
|
|
|
#pragma omp target update to(marr[:l][:l][l:])
|
|
|
|
|
|
|
|
|
|
#pragma omp target update to(marr[:2][:1][:])
|
|
|
|
|
|
|
|
|
|
#pragma omp target update from(marr[:2][:1][:])
|
|
|
|
|
|
|
|
|
|
#pragma omp target update to(marr[:2][:][:1])
|
|
|
|
|
|
|
|
|
|
#pragma omp target update from(marr[:2][:][:1])
|
|
|
|
|
|
|
|
|
|
#pragma omp target update to(marr[:2][:] [1:])
|
|
|
|
|
|
|
|
|
|
#pragma omp target update from(marr[:2][:][1:])
|
|
|
|
|
|
|
|
|
|
#pragma omp target update to(marr[:1][3:2][:2])
|
|
|
|
|
|
|
|
|
|
#pragma omp target update from(marr[:1][3:2][:2])
|
|
|
|
|
|
|
|
|
|
#pragma omp target update to(marr[:1][:2][0])
|
|
|
|
|
|
|
|
|
|
#pragma omp target update from(marr[:1][:2][0])
|
|
|
|
|
|
2020-07-10 02:27:32 +08:00
|
|
|
int arr[100][100];
|
|
|
|
|
#pragma omp target update to(arr[2][0:1:2])
|
|
|
|
|
|
|
|
|
|
#pragma omp target update from(arr[2][0:1:2])
|
2020-07-30 00:18:45 +08:00
|
|
|
|
|
|
|
|
#ifdef OMP51
|
|
|
|
|
#pragma omp target update to(present: arr[2][0:1:2])
|
|
|
|
|
|
|
|
|
|
#pragma omp target update from(present: arr[2][0:1:2], a)
|
|
|
|
|
#endif
|
|
|
|
|
|
2016-05-27 01:39:58 +08:00
|
|
|
return a + targ + (T)b;
|
|
|
|
|
}
|
2020-03-31 04:06:01 +08:00
|
|
|
// CHECK: static T a, *p;
|
2016-11-10 16:49:37 +08:00
|
|
|
// CHECK-NEXT: U b;
|
2016-05-27 01:39:58 +08:00
|
|
|
// CHECK-NEXT: int l;
|
2020-03-31 04:06:01 +08:00
|
|
|
// CHECK-NEXT: #pragma omp target update to(([a][targ])p,a) if(l > 5) device(l) nowait depend(inout : l){{$}}
|
|
|
|
|
// CHECK-NEXT: #pragma omp target update from(b,([a][targ])p) if(l < 5) device(l - 1) nowait depend(inout : l)
|
|
|
|
|
// CHECK: static int a, *p;
|
2016-05-27 01:39:58 +08:00
|
|
|
// CHECK-NEXT: float b;
|
|
|
|
|
// CHECK-NEXT: int l;
|
2020-03-31 04:06:01 +08:00
|
|
|
// CHECK-NEXT: #pragma omp target update to(([a][targ])p,a) if(l > 5) device(l) nowait depend(inout : l)
|
|
|
|
|
// CHECK-NEXT: #pragma omp target update from(b,([a][targ])p) if(l < 5) device(l - 1) nowait depend(inout : l)
|
|
|
|
|
// CHECK: static char a, *p;
|
2016-11-10 16:49:37 +08:00
|
|
|
// CHECK-NEXT: float b;
|
2016-05-27 01:39:58 +08:00
|
|
|
// CHECK-NEXT: int l;
|
2020-03-31 04:06:01 +08:00
|
|
|
// CHECK-NEXT: #pragma omp target update to(([a][targ])p,a) if(l > 5) device(l) nowait depend(inout : l)
|
|
|
|
|
// CHECK-NEXT: #pragma omp target update from(b,([a][targ])p) if(l < 5) device(l - 1) nowait depend(inout : l)
|
[OpenMP5.0] map item can be non-contiguous for target update
In order not to modify the `tgt_target_data_update` information but still be
able to pass the extra information for non-contiguous map item (offset,
count, and stride for each dimension), this patch overload `arg` when
the maptype is set as `OMP_MAP_DESCRIPTOR`. The origin `arg` is for
passing the pointer information, however, the overloaded `arg` is an
array of descriptor_dim:
struct descriptor_dim {
int64_t offset;
int64_t count;
int64_t stride
};
and the array size is the same as dimension size. In addition, since we
have count and stride information in descriptor_dim, we can replace/overload the
`arg_size` parameter by using dimension size.
For supporting `stride` in array section, we use a dummy dimension in
descriptor to store the unit size. The formula for counting the stride
in dimension D_n: `unit size * (D_0 * D_1 ... * D_n-1) * D_n.stride`.
Demonstrate how it works:
```
double arr[3][4][5];
D0: { offset = 0, count = 1, stride = 8 } // offset, count, dimension size always be 0, 1, 1 for this extra dimension, stride is the unit size
D1: { offset = 0, count = 2, stride = 8 * 1 * 2 = 16 } // stride = unit size * (product of dimension size of D0) * D1.stride = 4 * 1 * 2 = 8
D2: { offset = 2, count = 2, stride = 8 * (1 * 5) * 1 = 40 } // stride = unit size * (product of dimension size of D0, D1) * D2.stride = 4 * 5 * 1 = 20
D3: { offset = 0, count = 2, stride = 8 * (1 * 5 * 4) * 2 = 320 } // stride = unit size * (product of dimension size of D0, D1, D2) * D3.stride = 4 * 25 * 2 = 200
// X here means we need to offload this data, therefore, runtime will transfer
// data from offset 80, 96, 120, 136, 400, 416, 440, 456
// Runtime patch: https://reviews.llvm.org/D82245
// OOOOO OOOOO OOOOO
// OOOOO OOOOO OOOOO
// XOXOO OOOOO XOXOO
// XOXOO OOOOO XOXOO
```
Reviewed By: ABataev
Differential Revision: https://reviews.llvm.org/D84192
2020-11-07 11:03:55 +08:00
|
|
|
// CHECK: marr[10][10][10];
|
|
|
|
|
// CHECK-NEXT: #pragma omp target update to(marr[2][0:2][0:2])
|
|
|
|
|
// CHECK-NEXT: #pragma omp target update from(marr[2][0:2][0:2])
|
|
|
|
|
// CHECK-NEXT: #pragma omp target update from(marr[:2][0:2][0:2:1])
|
|
|
|
|
// CHECK-NEXT: #pragma omp target update to(marr[:l][:l][l:])
|
|
|
|
|
// CHECK-NEXT: #pragma omp target update to(marr[:2][:1][:])
|
|
|
|
|
// CHECK-NEXT: #pragma omp target update from(marr[:2][:1][:])
|
|
|
|
|
// CHECK-NEXT: #pragma omp target update to(marr[:2][:][:1])
|
|
|
|
|
// CHECK-NEXT: #pragma omp target update from(marr[:2][:][:1])
|
|
|
|
|
// CHECK-NEXT: #pragma omp target update to(marr[:2][:][1:])
|
|
|
|
|
// CHECK-NEXT: #pragma omp target update from(marr[:2][:][1:])
|
|
|
|
|
// CHECK-NEXT: #pragma omp target update to(marr[:1][3:2][:2])
|
|
|
|
|
// CHECK-NEXT: #pragma omp target update from(marr[:1][3:2][:2])
|
|
|
|
|
// CHECK-NEXT: #pragma omp target update to(marr[:1][:2][0])
|
|
|
|
|
// CHECK-NEXT: #pragma omp target update from(marr[:1][:2][0])
|
2020-07-10 02:27:32 +08:00
|
|
|
// CHECK: int arr[100][100];
|
|
|
|
|
// CHECK-NEXT: #pragma omp target update to(arr[2][0:1:2])
|
|
|
|
|
// CHECK-NEXT: #pragma omp target update from(arr[2][0:1:2])
|
2020-07-30 00:18:45 +08:00
|
|
|
// OMP5-NEXT: #pragma omp target update to(present: arr[2][0:1:2])
|
|
|
|
|
// OMP5-NEXT: #pragma omp target update from(present: arr[2][0:1:2], a)
|
2016-05-27 01:39:58 +08:00
|
|
|
|
|
|
|
|
int main(int argc, char **argv) {
|
|
|
|
|
static int a;
|
|
|
|
|
int n;
|
|
|
|
|
float f;
|
|
|
|
|
|
|
|
|
|
// CHECK: static int a;
|
|
|
|
|
// CHECK-NEXT: int n;
|
|
|
|
|
// CHECK-NEXT: float f;
|
2016-06-22 11:10:32 +08:00
|
|
|
#pragma omp target update to(a) if(f>0.0) device(n) nowait depend(in:n)
|
|
|
|
|
// CHECK-NEXT: #pragma omp target update to(a) if(f > 0.) device(n) nowait depend(in : n)
|
|
|
|
|
#pragma omp target update from(f) if(f<0.0) device(n+1) nowait depend(in:n)
|
|
|
|
|
// CHECK-NEXT: #pragma omp target update from(f) if(f < 0.) device(n + 1) nowait depend(in : n)
|
2020-07-10 02:27:32 +08:00
|
|
|
#pragma omp target update to(argv[2][0:1:2])
|
|
|
|
|
// CHECK-NEXT: #pragma omp target update to(argv[2][0:1:2])
|
|
|
|
|
#pragma omp target update from(argv[2][0:1:2])
|
|
|
|
|
// CHECK-NEXT: #pragma omp target update from(argv[2][0:1:2])
|
2020-07-30 00:18:45 +08:00
|
|
|
#ifdef OMP51
|
|
|
|
|
#pragma omp target update to(present: argv[2][0:1:2])
|
|
|
|
|
// OMP5-NEXT: #pragma omp target update to(present: arr[2][0:1:2])
|
|
|
|
|
#pragma omp target update from(argv[2][0:1:2], a)
|
|
|
|
|
// OMP5-NEXT: #pragma omp target update from(present: arr[2][0:1:2], a)
|
|
|
|
|
#endif
|
|
|
|
|
|
[OpenMP5.0] map item can be non-contiguous for target update
In order not to modify the `tgt_target_data_update` information but still be
able to pass the extra information for non-contiguous map item (offset,
count, and stride for each dimension), this patch overload `arg` when
the maptype is set as `OMP_MAP_DESCRIPTOR`. The origin `arg` is for
passing the pointer information, however, the overloaded `arg` is an
array of descriptor_dim:
struct descriptor_dim {
int64_t offset;
int64_t count;
int64_t stride
};
and the array size is the same as dimension size. In addition, since we
have count and stride information in descriptor_dim, we can replace/overload the
`arg_size` parameter by using dimension size.
For supporting `stride` in array section, we use a dummy dimension in
descriptor to store the unit size. The formula for counting the stride
in dimension D_n: `unit size * (D_0 * D_1 ... * D_n-1) * D_n.stride`.
Demonstrate how it works:
```
double arr[3][4][5];
D0: { offset = 0, count = 1, stride = 8 } // offset, count, dimension size always be 0, 1, 1 for this extra dimension, stride is the unit size
D1: { offset = 0, count = 2, stride = 8 * 1 * 2 = 16 } // stride = unit size * (product of dimension size of D0) * D1.stride = 4 * 1 * 2 = 8
D2: { offset = 2, count = 2, stride = 8 * (1 * 5) * 1 = 40 } // stride = unit size * (product of dimension size of D0, D1) * D2.stride = 4 * 5 * 1 = 20
D3: { offset = 0, count = 2, stride = 8 * (1 * 5 * 4) * 2 = 320 } // stride = unit size * (product of dimension size of D0, D1, D2) * D3.stride = 4 * 25 * 2 = 200
// X here means we need to offload this data, therefore, runtime will transfer
// data from offset 80, 96, 120, 136, 400, 416, 440, 456
// Runtime patch: https://reviews.llvm.org/D82245
// OOOOO OOOOO OOOOO
// OOOOO OOOOO OOOOO
// XOXOO OOOOO XOXOO
// XOXOO OOOOO XOXOO
```
Reviewed By: ABataev
Differential Revision: https://reviews.llvm.org/D84192
2020-11-07 11:03:55 +08:00
|
|
|
float marr[10][10][10];
|
|
|
|
|
// CHECK: marr[10][10][10];
|
|
|
|
|
#pragma omp target update to(marr[2][0:2][0:2])
|
|
|
|
|
// CHECK-NEXT: #pragma omp target update to(marr[2][0:2][0:2])
|
|
|
|
|
#pragma omp target update from(marr[2][0:2][0:2])
|
|
|
|
|
// CHECK-NEXT: #pragma omp target update from(marr[2][0:2][0:2])
|
|
|
|
|
#pragma omp target update to(marr[:n][:n][n:])
|
|
|
|
|
// CHECK: #pragma omp target update to(marr[:n][:n][n:])
|
|
|
|
|
#pragma omp target update from(marr[:2][:1][:])
|
|
|
|
|
// CHECK-NEXT: #pragma omp target update from(marr[:2][:1][:])
|
|
|
|
|
#pragma omp target update to(marr[:2][:][:1])
|
|
|
|
|
// CHECK-NEXT: #pragma omp target update to(marr[:2][:][:1])
|
|
|
|
|
#pragma omp target update from(marr[:2][:][:1])
|
|
|
|
|
// CHECK-NEXT: #pragma omp target update from(marr[:2][:][:1])
|
|
|
|
|
#pragma omp target update to(marr[:2][:][1:])
|
|
|
|
|
// CHECK-NEXT: #pragma omp target update to(marr[:2][:][1:])
|
|
|
|
|
#pragma omp target update from(marr[:2][:][1:])
|
|
|
|
|
// CHECK-NEXT: #pragma omp target update from(marr[:2][:][1:])
|
|
|
|
|
#pragma omp target update to(marr[:1][3:2][:2])
|
|
|
|
|
// CHECK-NEXT: #pragma omp target update to(marr[:1][3:2][:2])
|
|
|
|
|
#pragma omp target update from(marr[:1][3:2][:2])
|
|
|
|
|
// CHECK-NEXT: #pragma omp target update from(marr[:1][3:2][:2])
|
|
|
|
|
#pragma omp target update to(marr[:1][:2][0])
|
|
|
|
|
// CHECK-NEXT: #pragma omp target update to(marr[:1][:2][0])
|
|
|
|
|
#pragma omp target update from(marr[:1][:2][0])
|
|
|
|
|
// CHECK-NEXT: #pragma omp target update from(marr[:1][:2][0])
|
|
|
|
|
#pragma omp target update to(marr[:2:][0:2][0:2:1])
|
|
|
|
|
// CHECK-NEXT: #pragma omp target update to(marr[:2:][0:2][0:2:1])
|
|
|
|
|
#pragma omp target update from(marr[:2:][0:2][0:2:1])
|
|
|
|
|
// CHECK-NEXT: #pragma omp target update from(marr[:2:][0:2][0:2:1])
|
|
|
|
|
#pragma omp target update to(marr[:2:][:2:][0:2:1])
|
|
|
|
|
// CHECK-NEXT: #pragma omp target update to(marr[:2:][:2:][0:2:1])
|
|
|
|
|
#pragma omp target update from(marr[:2:][:2:][0:2:1])
|
|
|
|
|
// CHECK-NEXT: #pragma omp target update from(marr[:2:][:2:][0:2:1])
|
2020-07-10 02:27:32 +08:00
|
|
|
|
2016-05-27 01:39:58 +08:00
|
|
|
return foo(argc, f) + foo(argv[0][0], f) + a;
|
|
|
|
|
}
|
|
|
|
|
|
|
|
|
|
#endif
|