// RUN: %libomptarget-compile-run-and-check-generic
-#include <stdio.h>
#include <omp.h>
+#include <stdio.h>
-int test_omp_get_device_num()
-{
+int test_omp_get_device_num() {
/* checks that omp_get_device_num() == omp_get_num_devices() in the host */
int device_num = omp_get_device_num();
printf("device_num = %d\n", device_num);
- #pragma omp target
+#pragma omp target
{}
return (device_num == omp_get_num_devices());
}
-int main()
-{
+int main() {
int i;
- int failed=0;
+ int failed = 0;
if (!test_omp_get_device_num()) {
failed++;
// RUN: %libomptarget-compile-run-and-check-generic
-#include <stdio.h>
#include <omp.h>
+#include <stdio.h>
-int test_omp_get_num_devices()
-{
+int test_omp_get_num_devices() {
/* checks that omp_get_num_devices() > 0 */
int num_devices = omp_get_num_devices();
printf("num_devices = %d\n", num_devices);
- #pragma omp target
+#pragma omp target
{}
return (num_devices > 0);
}
-int main()
-{
+int main() {
int i;
- int failed=0;
+ int failed = 0;
if (!test_omp_get_num_devices()) {
failed++;
for (int i = 0; i < N; ++i)
hst_ptr[i] = 2;
-#pragma omp target teams distribute parallel for device(device) \
- map(tofrom:hst_ptr[0 : N])
+#pragma omp target teams distribute parallel for device(device) \
+ map(tofrom : hst_ptr[0 : N])
for (int i = 0; i < N; ++i)
hst_ptr[i] -= 1;
llvm_omp_target_free_host(hst_ptr, device);
// CHECK: PASS
if (sum == N)
- printf ("PASS\n");
+ printf("PASS\n");
}
for (int i = 0; i < N; ++i)
hst_ptr[i] = 2;
-#pragma omp target teams distribute parallel for map(tofrom : hst_ptr [0:N])
+#pragma omp target teams distribute parallel for map(tofrom : hst_ptr[0 : N])
for (int i = 0; i < N; ++i)
hst_ptr[i] -= 1;
// RUN: %libomptarget-compile-generic && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-generic 2>&1 | %fcheck-generic
// REQUIRES: libomptarget-debug
-#include <stdlib.h>
#include <stdio.h>
+#include <stdlib.h>
int *allocate(size_t n) {
int *ptr = malloc(sizeof(int) * n);
-#pragma omp target enter data map(to : ptr[:n])
+#pragma omp target enter data map(to : ptr[ : n])
return ptr;
}
void deallocate(int *ptr, size_t n) {
-#pragma omp target exit data map(delete : ptr[:n])
+#pragma omp target exit data map(delete : ptr[ : n])
free(ptr);
}
#pragma omp declare target
int *cnt;
-void foo() {
- ++(*cnt);
-}
+void foo() { ++(*cnt); }
#pragma omp end declare target
int main(void) {
int *V = allocate(10);
deallocate(A, 10);
deallocate(V, 10);
-// CHECK-NOT: RefCount=2
+ // CHECK-NOT: RefCount=2
cnt = malloc(sizeof(int));
*cnt = 0;
-#pragma omp target map(cnt[:1])
+#pragma omp target map(cnt[ : 1])
foo();
printf("Cnt = %d.\n", *cnt);
-// CHECK: Cnt = 1.
+ // CHECK: Cnt = 1.
*cnt = 0;
-#pragma omp target data map(cnt[:1])
+#pragma omp target data map(cnt[ : 1])
#pragma omp target
foo();
printf("Cnt = %d.\n", *cnt);
-// CHECK: Cnt = 1.
+ // CHECK: Cnt = 1.
free(cnt);
return 0;
}
-
// DEBUG: Libomptarget
// NDEBUG-NOT: Libomptarget
// NDEBUG-NOT: Target
-
int main() {
int arr[4] = {0, 1, 2, 3};
-#pragma omp target data map(alloc: arr[0:2])
-#pragma omp target data map(alloc: arr[1:2])
+#pragma omp target data map(alloc : arr[0 : 2])
+#pragma omp target data map(alloc : arr[1 : 2])
;
return 0;
}
B[i] = 2 * i;
}
-#pragma omp target enter data map(to : A [FROM:LENGTH], B [FROM:LENGTH])
-#pragma omp target enter data map(alloc : C [FROM:LENGTH])
+#pragma omp target enter data map(to : A[FROM : LENGTH], B[FROM : LENGTH])
+#pragma omp target enter data map(alloc : C[FROM : LENGTH])
// A, B and C have been mapped starting at index FROM, but inside the kernel
// they are captured implicitly so the library must look them up using their
}
}
-#pragma omp target exit data map(from : C [FROM:LENGTH])
-#pragma omp target exit data map(delete : A [FROM:LENGTH], B [FROM:LENGTH])
+#pragma omp target exit data map(from : C[FROM : LENGTH])
+#pragma omp target exit data map(delete : A[FROM : LENGTH], B[FROM : LENGTH])
int errors = 0;
for (int i = FROM; i < FROM + LENGTH; i++)
int main() {
float *A = (float *)malloc(N * sizeof(float));
-#pragma omp target enter data map(to : A [FROM:LENGTH])
+#pragma omp target enter data map(to : A[FROM : LENGTH])
// A, has been mapped starting at index FROM, but inside the use_device_ptr
// clause it is captured by base so the library must look it up using the
float *A_dev = NULL;
#pragma omp target data use_device_ptr(A)
{ A_dev = A; }
-#pragma omp target exit data map(delete : A [FROM:LENGTH])
+#pragma omp target exit data map(delete : A[FROM : LENGTH])
// CHECK: Success
if (A_dev == NULL || A_dev == A)
int main(void) {
int f = 5, r = 6, d = 7, af = 8;
- // Check exit from omp target data.
- // CHECK: f = 5, af = 8
- #pragma omp target data map(from: f) map(always, from: af)
+// Check exit from omp target data.
+// CHECK: f = 5, af = 8
+#pragma omp target data map(from : f) map(always, from : af)
{
- #pragma omp target exit data map(delete: f, af)
- }
- printf("f = %d, af = %d\n", f, af);
+#pragma omp target exit data map(delete : f, af)
+ } printf("f = %d, af = %d\n", f, af);
- // Check omp target exit data.
- // CHECK: f = 5, r = 6, d = 7, af = 8
- #pragma omp target exit data map(from: f) map(release: r) map(delete: d) \
- map(always, from: af)
+// Check omp target exit data.
+// CHECK: f = 5, r = 6, d = 7, af = 8
+#pragma omp target exit data map(from : f) map(release : r) map(delete : d) \
+ map(always, from : af)
printf("f = %d, r = %d, d = %d, af = %d\n", f, r, d, af);
return 0;
// RUN: %libomptarget-compilexx-run-and-check-generic
+#include <cinttypes>
#include <cstdio>
#include <cstdlib>
#include <vector>
-#include <cinttypes>
// Data structure definitions copied from OpenMP RTL.
struct MapComponentInfoTy {
int64_t Type;
void *Name;
MapComponentInfoTy() = default;
- MapComponentInfoTy(void *Base, void *Begin, int64_t Size, int64_t Type, void *Name)
+ MapComponentInfoTy(void *Base, void *Begin, int64_t Size, int64_t Type,
+ void *Name)
: Base(Base), Begin(Begin), Size(Size), Type(Type), Name(Name) {}
};
#endif
int64_t __tgt_mapper_num_components(void *rt_mapper_handle);
void __tgt_push_mapper_component(void *rt_mapper_handle, void *base,
- void *begin, int64_t size, int64_t type,
+ void *begin, int64_t size, int64_t type,
void *name);
#ifdef __cplusplus
}
int a;
double *b;
} C1;
-#pragma omp declare mapper(C1 s) map(to : s.a) map(from : s.b [0:2])
+#pragma omp declare mapper(C1 s) map(to : s.a) map(from : s.b[0 : 2])
typedef struct {
int a;
double *b;
C1 c;
} C;
-#pragma omp declare mapper(C s) map(to : s.a, s.c) map(from : s.b [0:2])
+#pragma omp declare mapper(C s) map(to : s.a, s.c) map(from : s.b[0 : 2])
typedef struct {
int e;
int a;
double *b;
} C1;
-#pragma omp declare mapper(C1 s) map(to : s.a) map(from : s.b [0:2])
+#pragma omp declare mapper(C1 s) map(to : s.a) map(from : s.b[0 : 2])
typedef struct {
int a;
double *b;
C1 c;
} C;
-#pragma omp declare mapper(C s) map(to : s.a, s.c) map(from : s.b [0:2])
+#pragma omp declare mapper(C s) map(to : s.a, s.c) map(from : s.b[0 : 2])
typedef struct {
int e;
int a;
double *b;
} C1;
-#pragma omp declare mapper(C1 s) map(to : s.a) map(from : s.b [0:2])
+#pragma omp declare mapper(C1 s) map(to : s.a) map(from : s.b[0 : 2])
typedef struct {
int a;
double *b;
C1 c;
} C;
-#pragma omp declare mapper(C s) map(to : s.a, s.c) map(from : s.b [0:2])
+#pragma omp declare mapper(C s) map(to : s.a, s.c) map(from : s.b[0 : 2])
typedef struct {
int e;
MyObjectA *arr;
int len;
};
-#pragma omp declare mapper(MyObjectB obj) map(obj, obj.arr[:obj.len])
+#pragma omp declare mapper(MyObjectB obj) map(obj, obj.arr[ : obj.len])
class MyObjectC {
public:
MyObjectB *arr;
int len;
};
-#pragma omp declare mapper(MyObjectC obj) map(obj, obj.arr[:obj.len])
+#pragma omp declare mapper(MyObjectC obj) map(obj, obj.arr[ : obj.len])
int main(void) {
MyObjectC *outer = new MyObjectC[N];
outer[i].show();
printf("Sending data to device...\n");
-#pragma omp target enter data map(to : outer[:N])
+#pragma omp target enter data map(to : outer[ : N])
printf("Calling foo()...\n");
#pragma omp target teams distribute parallel for
printf("foo() complete!\n");
printf("Sending data back to host...\n");
-#pragma omp target exit data map(from : outer[:N])
+#pragma omp target exit data map(from : outer[ : N])
printf("Modified Data Hierarchy:\n");
for (int i = 0; i < N; i++)
int a;
double *b;
} C1;
-#pragma omp declare mapper(C1 s) map(to : s.a) map(from : s.b [0:2])
+#pragma omp declare mapper(C1 s) map(to : s.a) map(from : s.b[0 : 2])
typedef struct {
int a;
double *b;
C1 c;
} C;
-#pragma omp declare mapper(C s) map(to : s.a, s.c) map(from : s.b [0:2])
+#pragma omp declare mapper(C s) map(to : s.a, s.c) map(from : s.b[0 : 2])
typedef struct {
int e;
int a;
double *b;
} C1;
-#pragma omp declare mapper(C1 s) map(to : s.a) map(from : s.b [0:2])
+#pragma omp declare mapper(C1 s) map(to : s.a) map(from : s.b[0 : 2])
typedef struct {
int a;
double *b;
C1 c;
} C;
-#pragma omp declare mapper(C s) map(to : s.a, s.c) map(from : s.b [0:2])
+#pragma omp declare mapper(C s) map(to : s.a, s.c) map(from : s.b[0 : 2])
typedef struct {
int e;
int a;
double *b;
} C;
-#pragma omp declare mapper(id1 : C s) map(to : s.a) map(from : s.b [0:2])
+#pragma omp declare mapper(id1 : C s) map(to : s.a) map(from : s.b[0 : 2])
typedef struct {
int e;
int h;
short *g;
} D;
-#pragma omp declare mapper(default \
- : D r) map(from \
- : r.e) map(mapper(id1), tofrom \
- : r.f) map(tofrom \
- : r.g [0:r.h])
+#pragma omp declare mapper(default : D r) map(from : r.e) \
+ map(mapper(id1), tofrom : r.f) map(tofrom : r.g[0 : r.h])
int main() {
constexpr int N = 10;
int *a;
};
-#pragma omp declare mapper(id: C s) map(s.a[0:NUM])
+#pragma omp declare mapper(id : C s) map(s.a[0 : NUM])
int main() {
C c;
- c.a = (int*) malloc(sizeof(int)*NUM);
+ c.a = (int *)malloc(sizeof(int) * NUM);
for (int i = 0; i < NUM; i++) {
c.a[i] = 1;
}
- #pragma omp target teams distribute parallel for map(mapper(id),tofrom: c)
+#pragma omp target teams distribute parallel for map(mapper(id), tofrom : c)
for (int i = 0; i < NUM; i++) {
++c.a[i];
}
printf("Sum = %d\n", sum);
return 0;
}
-
int *a;
};
-#pragma omp declare mapper(id: C s) map(s.a[0:NUM])
+#pragma omp declare mapper(id : C s) map(s.a[0 : NUM])
int main() {
C c;
- c.a = (int*) malloc(sizeof(int)*NUM);
+ c.a = (int *)malloc(sizeof(int) * NUM);
for (int i = 0; i < NUM; i++) {
c.a[i] = 1;
}
- #pragma omp target data map(mapper(id),tofrom: c)
+#pragma omp target data map(mapper(id), tofrom : c)
{
- #pragma omp target teams distribute parallel for
+#pragma omp target teams distribute parallel for
for (int i = 0; i < NUM; i++) {
++c.a[i];
}
printf("Sum = %d\n", sum);
return 0;
}
-
int *a;
};
-#pragma omp declare mapper(id: C s) map(s.a[0:NUM])
+#pragma omp declare mapper(id : C s) map(s.a[0 : NUM])
int main() {
C c;
- c.a = (int*) malloc(sizeof(int)*NUM);
+ c.a = (int *)malloc(sizeof(int) * NUM);
for (int i = 0; i < NUM; i++) {
c.a[i] = 1;
}
- #pragma omp target enter data map(mapper(id),to: c)
- #pragma omp target teams distribute parallel for
+#pragma omp target enter data map(mapper(id), to : c)
+#pragma omp target teams distribute parallel for
for (int i = 0; i < NUM; i++) {
++c.a[i];
}
- #pragma omp target exit data map(mapper(id),from: c)
+#pragma omp target exit data map(mapper(id), from : c)
int sum = 0;
for (int i = 0; i < NUM; i++) {
sum += c.a[i];
printf("Sum = %d\n", sum);
return 0;
}
-
int *a;
};
-#pragma omp declare mapper(id: C s) map(s.a[0:NUM])
+#pragma omp declare mapper(id : C s) map(s.a[0 : NUM])
int main() {
C c;
int sum = 0;
- c.a = (int*) malloc(sizeof(int)*NUM);
+ c.a = (int *)malloc(sizeof(int) * NUM);
for (int i = 0; i < NUM; i++) {
c.a[i] = 1;
}
- #pragma omp target enter data map(mapper(id),alloc: c)
- #pragma omp target teams distribute parallel for
+#pragma omp target enter data map(mapper(id), alloc : c)
+#pragma omp target teams distribute parallel for
for (int i = 0; i < NUM; i++) {
c.a[i] = 0;
}
- #pragma omp target update from(mapper(id): c)
+#pragma omp target update from(mapper(id) : c)
for (int i = 0; i < NUM; i++) {
sum += c.a[i];
}
for (int i = 0; i < NUM; i++) {
c.a[i] = 1;
}
- #pragma omp target update to(mapper(id): c)
- #pragma omp target teams distribute parallel for
+#pragma omp target update to(mapper(id) : c)
+#pragma omp target teams distribute parallel for
for (int i = 0; i < NUM; i++) {
++c.a[i];
}
}
// CHECK: Sum (after update to) = 1024
printf("Sum (after update to) = %d\n", sum);
- #pragma omp target update from(mapper(id): c)
+#pragma omp target update from(mapper(id) : c)
sum = 0;
for (int i = 0; i < NUM; i++) {
sum += c.a[i];
}
// CHECK: Sum (after second update from) = 2048
printf("Sum (after second update from) = %d\n", sum);
- #pragma omp target exit data map(mapper(id),delete: c)
+#pragma omp target exit data map(mapper(id), delete : c)
return 0;
}
-
// RUN: %libomptarget-compile-run-and-check-generic
-#include <stdio.h>
#include <omp.h>
+#include <stdio.h>
#pragma omp declare target
int isHost;
int main(void) {
isHost = -1;
-#pragma omp target enter data map(to: isHost)
+#pragma omp target enter data map(to : isHost)
#pragma omp target
{ isHost = omp_is_initial_device(); }
printf("Runtime error, isHost=%d\n", isHost);
}
-#pragma omp target exit data map(delete: isHost)
+#pragma omp target exit data map(delete : isHost)
// CHECK: Target region executed on the device
printf("Target region executed on the %s\n", isHost ? "host" : "device");
s1.p = A;
- // DEBUG: Update pointer ([[DEV_PTR:0x[^ ]+]]) -> {{\[}}[[DEV_OBJ_A:0x[^ ]+]]{{\]}}
- #pragma omp target enter data map(alloc : s1.p [0:10])
+// DEBUG: Update pointer ([[DEV_PTR:0x[^ ]+]]) -> {{\[}}[[DEV_OBJ_A:0x[^ ]+]]{{\]}}
+#pragma omp target enter data map(alloc : s1.p[0 : 10])
- // DEBUG-NOT: Update pointer ([[DEV_PTR]]) -> {{\[}}[[DEV_OBJ_A]]{{\]}}
- #pragma omp target map(alloc : s1.p [0:10])
+// DEBUG-NOT: Update pointer ([[DEV_PTR]]) -> {{\[}}[[DEV_OBJ_A]]{{\]}}
+#pragma omp target map(alloc : s1.p[0 : 10])
{
for (int i = 0; i < 10; ++i)
s1.p[i] = i;
}
- #pragma omp target exit data map(from : s1.p [0:10])
+#pragma omp target exit data map(from : s1.p[0 : 10])
int fail_A = 0;
for (int i = 0; i < 10; ++i) {
// XFAIL: amdgcn-amd-amdhsa
// XFAIL: amdgcn-amd-amdhsa-LTO
-#include <stdio.h>
#include <stdint.h>
+#include <stdio.h>
// CHECK: before: [[V1:111]] [[V2:222]] [[PX:0x[^ ]+]] [[PY:0x[^ ]+]]
// CHECK: lambda: [[V1]] [[V2]] [[PX_TGT:0x[^ ]+]] 0x{{.*}}
printf("before: %d %ld %p %p\n", x[1], y[1], &x[0], &y[0]);
- intptr_t xp = (intptr_t) &x[0];
+ intptr_t xp = (intptr_t)&x[0];
#pragma omp target firstprivate(xp)
{
lambda();
- printf("tgt : %d %p %d\n", x[1], &x[0], (&x[0] != (int*) xp));
+ printf("tgt : %d %p %d\n", x[1], &x[0], (&x[0] != (int *)xp));
}
printf("out : %d %ld %p %p\n", x[1], y[1], &x[0], &y[0]);
return 0;
}
-
C[I] = -9;
}
-#pragma omp target data map(tofrom : C [0:N]) map(to : A [0:N], B [0:N])
+#pragma omp target data map(tofrom : C[0 : N]) map(to : A[0 : N], B[0 : N])
{
forall(0, N, [&](int I) { C[I] += A[I] + B[I]; });
}
// RUN: %not %libomptarget-run-generic 1 2>&1 | %fcheck-generic
// RUN: %not %libomptarget-run-generic inf 2>&1 | %fcheck-generic
+#include <limits.h>
#include <omp.h>
#include <stdio.h>
-#include <limits.h>
#include <string.h>
int main(int argc, char *argv[]) {
}
} else {
for (int I = 0; I < DynRef; ++I) {
- #pragma omp target enter data map(alloc: X)
+#pragma omp target enter data map(alloc : X)
}
}
// Disassociate while hold reference count > 0.
int Status = 0;
- #pragma omp target data map(ompx_hold,alloc: X)
+#pragma omp target data map(ompx_hold, alloc : X)
#if HOLD_MORE
- #pragma omp target data map(ompx_hold,alloc: X)
- #pragma omp target data map(ompx_hold,alloc: X)
+#pragma omp target data map(ompx_hold, alloc : X)
+#pragma omp target data map(ompx_hold, alloc : X)
#endif
{
// CHECK: Libomptarget error: Trying to disassociate a pointer with a
#include <stdio.h>
#define CHECK_PRESENCE(Var1, Var2, Var3) \
- printf(" presence of %s, %s, %s: %d, %d, %d\n", \
- #Var1, #Var2, #Var3, \
+ printf(" presence of %s, %s, %s: %d, %d, %d\n", #Var1, #Var2, #Var3, \
omp_target_is_present(&(Var1), omp_get_default_device()), \
omp_target_is_present(&(Var2), omp_get_default_device()), \
omp_target_is_present(&(Var3), omp_get_default_device()))
#define CHECK_VALUES(Var1, Var2) \
- printf(" values of %s, %s: %d, %d\n", \
- #Var1, #Var2, (Var1), (Var2))
+ printf(" values of %s, %s: %d, %d\n", #Var1, #Var2, (Var1), (Var2))
int main() {
- struct S { int i; int j; } s;
+ struct S {
+ int i;
+ int j;
+ } s;
// CHECK: presence of s, s.i, s.j: 0, 0, 0
CHECK_PRESENCE(s, s.i, s.j);
printf("check: ompx_hold only on first member\n");
s.i = 20;
s.j = 30;
- #pragma omp target data map(tofrom: s) map(ompx_hold,tofrom: s.i) \
- map(tofrom: s.j)
+#pragma omp target data map(tofrom : s) map(ompx_hold, tofrom : s.i) \
+ map(tofrom : s.j)
{
// CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
CHECK_PRESENCE(s, s.i, s.j);
- #pragma omp target map(tofrom: s)
+#pragma omp target map(tofrom : s)
{
s.i = 21;
s.j = 31;
}
- #pragma omp target exit data map(delete: s, s.i)
+#pragma omp target exit data map(delete : s, s.i)
// ompx_hold on s.i applies to all of s.
// CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
// CHECK-NEXT: values of s.i, s.j: 20, 30
printf("check: ompx_hold only on last member\n");
s.i = 20;
s.j = 30;
- #pragma omp target data map(tofrom: s) map(tofrom: s.i) \
- map(ompx_hold,tofrom: s.j)
+#pragma omp target data map(tofrom : s) map(tofrom : s.i) \
+ map(ompx_hold, tofrom : s.j)
{
// CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
CHECK_PRESENCE(s, s.i, s.j);
- #pragma omp target map(tofrom: s)
+#pragma omp target map(tofrom : s)
{
s.i = 21;
s.j = 31;
}
- #pragma omp target exit data map(delete: s, s.i)
+#pragma omp target exit data map(delete : s, s.i)
// ompx_hold on s.j applies to all of s.
// CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
// CHECK-NEXT: values of s.i, s.j: 20, 30
printf("check: ompx_hold only on struct\n");
s.i = 20;
s.j = 30;
- #pragma omp target data map(ompx_hold,tofrom: s) map(tofrom: s.i) \
- map(tofrom: s.j)
+#pragma omp target data map(ompx_hold, tofrom : s) map(tofrom : s.i) \
+ map(tofrom : s.j)
{
// CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
CHECK_PRESENCE(s, s.i, s.j);
- #pragma omp target map(tofrom: s)
+#pragma omp target map(tofrom : s)
{
s.i = 21;
s.j = 31;
}
- #pragma omp target exit data map(delete: s, s.i)
+#pragma omp target exit data map(delete : s, s.i)
// CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
// CHECK-NEXT: values of s.i, s.j: 20, 30
CHECK_PRESENCE(s, s.i, s.j);
printf("check: parent DynRefCount=1 is not sufficient for transfer\n");
s.i = 20;
s.j = 30;
- #pragma omp target data map(ompx_hold, tofrom: s)
- #pragma omp target data map(ompx_hold, tofrom: s)
+#pragma omp target data map(ompx_hold, tofrom : s)
+#pragma omp target data map(ompx_hold, tofrom : s)
{
// CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
CHECK_PRESENCE(s, s.i, s.j);
- #pragma omp target map(from: s.i, s.j)
+#pragma omp target map(from : s.i, s.j)
{
s.i = 21;
s.j = 31;
// CHECK-NEXT: values of s.i, s.j: 20, 30
CHECK_PRESENCE(s, s.i, s.j);
CHECK_VALUES(s.i, s.j);
- #pragma omp target map(to: s.i, s.j)
+#pragma omp target map(to : s.i, s.j)
{ // No transfer here even though parent's DynRefCount=1.
// CHECK-NEXT: values of s.i, s.j: 21, 31
CHECK_VALUES(s.i, s.j);
printf("check: parent HoldRefCount=1 is not sufficient for transfer\n");
s.i = 20;
s.j = 30;
- #pragma omp target data map(tofrom: s)
- #pragma omp target data map(tofrom: s)
+#pragma omp target data map(tofrom : s)
+#pragma omp target data map(tofrom : s)
{
// CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
CHECK_PRESENCE(s, s.i, s.j);
- #pragma omp target map(ompx_hold, from: s.i, s.j)
+#pragma omp target map(ompx_hold, from : s.i, s.j)
{
s.i = 21;
s.j = 31;
// CHECK-NEXT: values of s.i, s.j: 20, 30
CHECK_PRESENCE(s, s.i, s.j);
CHECK_VALUES(s.i, s.j);
- #pragma omp target map(ompx_hold, to: s.i, s.j)
+#pragma omp target map(ompx_hold, to : s.i, s.j)
{ // No transfer here even though parent's HoldRefCount=1.
// CHECK-NEXT: values of s.i, s.j: 21, 31
CHECK_VALUES(s.i, s.j);
printf("check: parent TotalRefCount=1 is not sufficient for transfer\n");
s.i = 20;
s.j = 30;
- #pragma omp target data map(ompx_hold, tofrom: s)
+#pragma omp target data map(ompx_hold, tofrom : s)
{
// CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
CHECK_PRESENCE(s, s.i, s.j);
- #pragma omp target map(ompx_hold, tofrom: s.i, s.j)
+#pragma omp target map(ompx_hold, tofrom : s.i, s.j)
{
s.i = 21;
s.j = 31;
}
- #pragma omp target exit data map(from: s.i, s.j)
+#pragma omp target exit data map(from : s.i, s.j)
// No transfer here even though parent's TotalRefCount=1.
// CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
// CHECK-NEXT: values of s.i, s.j: 20, 30
#include <stdio.h>
#define CHECK_PRESENCE(Var1, Var2, Var3) \
- printf(" presence of %s, %s, %s: %d, %d, %d\n", \
- #Var1, #Var2, #Var3, \
+ printf(" presence of %s, %s, %s: %d, %d, %d\n", #Var1, #Var2, #Var3, \
omp_target_is_present(&Var1, omp_get_default_device()), \
omp_target_is_present(&Var2, omp_get_default_device()), \
omp_target_is_present(&Var3, omp_get_default_device()))
// CHECK-NEXT: structured{{.*}}
printf(" structured dec of dyn\n");
- #pragma omp target data map(tofrom: m) map(alloc: r, d)
+#pragma omp target data map(tofrom : m) map(alloc : r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
- #pragma omp target data map(tofrom: m) map(alloc: r, d)
+#pragma omp target data map(tofrom : m) map(alloc : r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
// CHECK-NEXT: dynamic{{.*}}
printf(" dynamic dec/reset of dyn\n");
- #pragma omp target data map(tofrom: m) map(alloc: r, d)
+#pragma omp target data map(tofrom : m) map(alloc : r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
- #pragma omp target data map(tofrom: m) map(alloc: r, d)
+#pragma omp target data map(tofrom : m) map(alloc : r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
- #pragma omp target exit data map(from: m) map(release: r)
+#pragma omp target exit data map(from : m) map(release : r)
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
- #pragma omp target exit data map(from: m) map(release: r) map(delete: d)
+#pragma omp target exit data map(from : m) map(release : r) map(delete : d)
// CHECK-NEXT: presence of m, r, d: 0, 0, 0
CHECK_PRESENCE(m, r, d);
}
// CHECK-NEXT: presence of m, r, d: 0, 0, 0
CHECK_PRESENCE(m, r, d);
- #pragma omp target exit data map(from: m) map(release: r) map(delete: d)
+#pragma omp target exit data map(from : m) map(release : r) map(delete : d)
// CHECK-NEXT: presence of m, r, d: 0, 0, 0
CHECK_PRESENCE(m, r, d);
}
// CHECK-NEXT: dynamic{{.*}}
printf(" dynamic dec/reset of dyn\n");
- #pragma omp target data map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
+#pragma omp target data map(ompx_hold, tofrom : m) map(ompx_hold, alloc : r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
- #pragma omp target data map(ompx_hold, tofrom: m) \
- map(ompx_hold, alloc: r, d)
+#pragma omp target data map(ompx_hold, tofrom : m) map(ompx_hold, alloc : r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
- #pragma omp target exit data map(from: m) map(release: r)
+#pragma omp target exit data map(from : m) map(release : r)
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
- #pragma omp target exit data map(from: m) map(release: r) map(delete: d)
+#pragma omp target exit data map(from : m) map(release : r) map(delete : d)
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
}
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
- #pragma omp target exit data map(from: m) map(release: r) map(delete: d)
+#pragma omp target exit data map(from : m) map(release : r) map(delete : d)
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
}
// CHECK-NEXT: structured{{.*}}
printf(" structured dec of dyn\n");
- #pragma omp target data map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
+#pragma omp target data map(ompx_hold, tofrom : m) map(ompx_hold, alloc : r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
- #pragma omp target data map(ompx_hold, tofrom: m) \
- map(ompx_hold, alloc: r, d)
+#pragma omp target data map(ompx_hold, tofrom : m) map(ompx_hold, alloc : r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
- #pragma omp target data map(tofrom: m) map(alloc: r, d)
+#pragma omp target data map(tofrom : m) map(alloc : r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
- #pragma omp target data map(tofrom: m) map(alloc: r, d)
+#pragma omp target data map(tofrom : m) map(alloc : r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
// CHECK-NEXT: dynamic{{.*}}
printf(" dynamic dec/reset of dyn\n");
- #pragma omp target enter data map(to: m) map(alloc: r, d)
+#pragma omp target enter data map(to : m) map(alloc : r, d)
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
- #pragma omp target enter data map(to: m) map(alloc: r, d)
+#pragma omp target enter data map(to : m) map(alloc : r, d)
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
- #pragma omp target data map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
+#pragma omp target data map(ompx_hold, tofrom : m) map(ompx_hold, alloc : r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
- #pragma omp target data map(ompx_hold, tofrom: m) \
- map(ompx_hold, alloc: r, d)
+#pragma omp target data map(ompx_hold, tofrom : m) map(ompx_hold, alloc : r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
- #pragma omp target exit data map(from: m) map(release: r)
+#pragma omp target exit data map(from : m) map(release : r)
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
- #pragma omp target exit data map(from: m) map(release: r) map(delete: d)
+#pragma omp target exit data map(from : m) map(release : r) map(delete : d)
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
}
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
- #pragma omp target exit data map(from: m) map(release: r) map(delete: d)
+#pragma omp target exit data map(from : m) map(release : r) map(delete : d)
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
}
// CHECK-NEXT: structured{{.*}}
printf(" structured dec of dyn\n");
- #pragma omp target data map(tofrom: m) map(alloc: r, d)
+#pragma omp target data map(tofrom : m) map(alloc : r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
- #pragma omp target data map(tofrom: m) map(alloc: r, d)
+#pragma omp target data map(tofrom : m) map(alloc : r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
- #pragma omp target data map(ompx_hold, tofrom: m) \
- map(ompx_hold, alloc: r, d)
+#pragma omp target data map(ompx_hold, tofrom : m) map(ompx_hold, alloc : r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
- #pragma omp target data map(ompx_hold, tofrom: m) \
- map(ompx_hold, alloc: r, d)
+#pragma omp target data map(ompx_hold, tofrom : m) map(ompx_hold, alloc : r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
// CHECK-NEXT: dynamic{{.*}}
printf(" dynamic dec/reset of dyn\n");
- #pragma omp target enter data map(to: m) map(alloc: r, d)
+#pragma omp target enter data map(to : m) map(alloc : r, d)
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
- #pragma omp target enter data map(to: m) map(alloc: r, d)
+#pragma omp target enter data map(to : m) map(alloc : r, d)
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
- #pragma omp target data map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
+#pragma omp target data map(ompx_hold, tofrom : m) map(ompx_hold, alloc : r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
- #pragma omp target data map(ompx_hold, tofrom: m) \
- map(ompx_hold, alloc: r, d)
+#pragma omp target data map(ompx_hold, tofrom : m) map(ompx_hold, alloc : r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
}
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
- #pragma omp target exit data map(from: m) map(release: r)
+#pragma omp target exit data map(from : m) map(release : r)
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
- #pragma omp target exit data map(from: m) map(release: r) map(delete: d)
+#pragma omp target exit data map(from : m) map(release : r) map(delete : d)
// CHECK-NEXT: presence of m, r, d: 0, 0, 0
CHECK_PRESENCE(m, r, d);
#include <stdio.h>
#define CHECK_PRESENCE(Var1, Var2, Var3) \
- printf(" presence of %s, %s, %s: %d, %d, %d\n", \
- #Var1, #Var2, #Var3, \
+ printf(" presence of %s, %s, %s: %d, %d, %d\n", #Var1, #Var2, #Var3, \
omp_target_is_present(&Var1, omp_get_default_device()), \
omp_target_is_present(&Var2, omp_get_default_device()), \
omp_target_is_present(&Var3, omp_get_default_device()))
// CHECK-NEXT: once
printf(" once\n");
- #pragma omp target map(tofrom: m) map(alloc: r, d)
+#pragma omp target map(tofrom : m) map(alloc : r, d)
;
// CHECK-NEXT: presence of m, r, d: 0, 0, 0
CHECK_PRESENCE(m, r, d);
// CHECK-NEXT: twice
printf(" twice\n");
- #pragma omp target data map(tofrom: m) map(alloc: r, d)
+#pragma omp target data map(tofrom : m) map(alloc : r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
- #pragma omp target map(tofrom: m) map(alloc: r, d)
+#pragma omp target map(tofrom : m) map(alloc : r, d)
;
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
// CHECK-NEXT: once
printf(" once\n");
- #pragma omp target map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
+#pragma omp target map(ompx_hold, tofrom : m) map(ompx_hold, alloc : r, d)
;
// CHECK-NEXT: presence of m, r, d: 0, 0, 0
CHECK_PRESENCE(m, r, d);
// CHECK-NEXT: twice
printf(" twice\n");
- #pragma omp target data map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
+#pragma omp target data map(ompx_hold, tofrom : m) map(ompx_hold, alloc : r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
- #pragma omp target map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
+#pragma omp target map(ompx_hold, tofrom : m) map(ompx_hold, alloc : r, d)
;
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
// CHECK-NEXT: once each
printf(" once each\n");
- #pragma omp target data map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
+#pragma omp target data map(ompx_hold, tofrom : m) map(ompx_hold, alloc : r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
- #pragma omp target map(tofrom: m) map(alloc: r, d)
+#pragma omp target map(tofrom : m) map(alloc : r, d)
;
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
// CHECK-NEXT: twice each
printf(" twice each\n");
- #pragma omp target data map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
+#pragma omp target data map(ompx_hold, tofrom : m) map(ompx_hold, alloc : r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
- #pragma omp target data map(ompx_hold, tofrom: m) \
- map(ompx_hold, alloc: r, d)
+#pragma omp target data map(ompx_hold, tofrom : m) map(ompx_hold, alloc : r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
- #pragma omp target data map(tofrom: m) map(alloc: r, d)
+#pragma omp target data map(tofrom : m) map(alloc : r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
- #pragma omp target map(tofrom: m) map(alloc: r, d)
+#pragma omp target map(tofrom : m) map(alloc : r, d)
;
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
// CHECK-NEXT: once each
printf(" once each\n");
- #pragma omp target data map(tofrom: m) map(alloc: r, d)
+#pragma omp target data map(tofrom : m) map(alloc : r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
- #pragma omp target map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
+#pragma omp target map(ompx_hold, tofrom : m) map(ompx_hold, alloc : r, d)
;
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
// CHECK-NEXT: twice each
printf(" twice each\n");
- #pragma omp target data map(tofrom: m) map(alloc: r, d)
+#pragma omp target data map(tofrom : m) map(alloc : r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
- #pragma omp target data map(tofrom: m) map(alloc: r, d)
+#pragma omp target data map(tofrom : m) map(alloc : r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
- #pragma omp target data map(ompx_hold, tofrom: m) \
- map(ompx_hold, alloc: r, d)
+#pragma omp target data map(ompx_hold, tofrom : m) map(ompx_hold, alloc : r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
- #pragma omp target map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
+#pragma omp target map(ompx_hold, tofrom : m) map(ompx_hold, alloc : r, d)
;
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
s2.ptr1 = malloc(sizeof(int));
s2.ptr2 = malloc(2 * sizeof(int));
-#pragma omp target enter data map(to: s2.ptr2[0:1])
-#pragma omp target map(s.ptr1[0:1], s.ptr2[0:2])
+#pragma omp target enter data map(to : s2.ptr2[0 : 1])
+#pragma omp target map(s.ptr1[0 : 1], s.ptr2[0 : 2])
{
s.ptr1[0] = 1;
s.ptr2[0] = 2;
s.ptr2[1] = 3;
}
-#pragma omp target exit data map(from: s2.ptr1[0:1], s2.ptr2[0:1])
+#pragma omp target exit data map(from : s2.ptr1[0 : 1], s2.ptr2[0 : 1])
// CHECK: s.ptr1[0] = 1
// CHECK: s.ptr2[0] = 2
// RUN: %libomptarget-run-fail-generic 2>&1 \
// RUN: | %fcheck-generic
-
#include <stdio.h>
int main() {
fprintf(stderr, "addr=%p, size=%ld\n", &i, sizeof i);
// CHECK-NOT: Libomptarget
-#pragma omp target data map(alloc: i)
-#pragma omp target map(present, alloc: i)
+#pragma omp target data map(alloc : i)
+#pragma omp target map(present, alloc : i)
;
// CHECK: i is present
// CHECK: Libomptarget error: Call to targetDataBegin failed, abort target.
// CHECK: Libomptarget error: Failed to process data before launching the kernel.
// CHECK: Libomptarget fatal error 1: failure of target construct while offloading is mandatory
-#pragma omp target map(present, alloc: i)
+#pragma omp target map(present, alloc : i)
;
// CHECK-NOT: i is present
// RUN: %libomptarget-run-fail-generic 2>&1 \
// RUN: | %fcheck-generic
-
// END.
#include <stdio.h>
#define BEFORE 0
-#define AFTER 1
+#define AFTER 1
#define SIZE 100
#if EXTENDS == BEFORE
-# define SMALL_BEG (SIZE-2)
-# define SMALL_END SIZE
-# define LARGE_BEG 0
-# define LARGE_END SIZE
+#define SMALL_BEG (SIZE - 2)
+#define SMALL_END SIZE
+#define LARGE_BEG 0
+#define LARGE_END SIZE
#elif EXTENDS == AFTER
-# define SMALL_BEG 0
-# define SMALL_END 2
-# define LARGE_BEG 0
-# define LARGE_END SIZE
+#define SMALL_BEG 0
+#define SMALL_END 2
+#define LARGE_BEG 0
+#define LARGE_END SIZE
#else
-# error EXTENDS undefined
+#error EXTENDS undefined
#endif
-#define SMALL_SIZE (SMALL_END-SMALL_BEG)
-#define LARGE_SIZE (LARGE_END-LARGE_BEG)
+#define SMALL_SIZE (SMALL_END - SMALL_BEG)
+#define LARGE_SIZE (LARGE_END - LARGE_BEG)
-#define SMALL SMALL_BEG:SMALL_SIZE
-#define LARGE LARGE_BEG:LARGE_SIZE
+#define SMALL \
+ SMALL_BEG: \
+ SMALL_SIZE
+#define LARGE \
+ LARGE_BEG: \
+ LARGE_SIZE
int main() {
int arr[SIZE];
LARGE_SIZE * sizeof arr[0]);
// CHECK-NOT: Libomptarget
-#pragma omp target data map(alloc: arr[LARGE])
+#pragma omp target data map(alloc : arr[LARGE])
{
-#pragma omp target map(present, tofrom: arr[SMALL])
+#pragma omp target map(present, tofrom : arr[SMALL])
;
}
// CHECK: Libomptarget error: Call to targetDataBegin failed, abort target.
// CHECK: Libomptarget error: Failed to process data before launching the kernel.
// CHECK: Libomptarget fatal error 1: failure of target construct while offloading is mandatory
-#pragma omp target data map(alloc: arr[SMALL])
+#pragma omp target data map(alloc : arr[SMALL])
{
-#pragma omp target map(present, tofrom: arr[LARGE])
+#pragma omp target map(present, tofrom : arr[LARGE])
;
}
// RUN: %libomptarget-run-fail-generic 2>&1 \
// RUN: | %fcheck-generic
-
#include <stdio.h>
int main() {
fprintf(stderr, "addr=%p, size=%ld\n", &i, sizeof i);
// CHECK-NOT: Libomptarget
-#pragma omp target data map(alloc: i)
-#pragma omp target data map(present, alloc: i)
+#pragma omp target data map(alloc : i)
+#pragma omp target data map(present, alloc : i)
;
// CHECK: i is present
// CHECK: Libomptarget message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] ([[#SIZE]] bytes)
// CHECK: Libomptarget fatal error 1: failure of target construct while offloading is mandatory
-#pragma omp target data map(present, alloc: i)
+#pragma omp target data map(present, alloc : i)
;
// CHECK-NOT: i is present
// RUN: %libomptarget-run-fail-generic 2>&1 \
// RUN: | %fcheck-generic
-
// END.
#include <stdio.h>
#define BEFORE 0
-#define AFTER 1
+#define AFTER 1
#define SIZE 100
#if EXTENDS == BEFORE
-# define SMALL_BEG (SIZE-2)
-# define SMALL_END SIZE
-# define LARGE_BEG 0
-# define LARGE_END SIZE
+#define SMALL_BEG (SIZE - 2)
+#define SMALL_END SIZE
+#define LARGE_BEG 0
+#define LARGE_END SIZE
#elif EXTENDS == AFTER
-# define SMALL_BEG 0
-# define SMALL_END 2
-# define LARGE_BEG 0
-# define LARGE_END SIZE
+#define SMALL_BEG 0
+#define SMALL_END 2
+#define LARGE_BEG 0
+#define LARGE_END SIZE
#else
-# error EXTENDS undefined
+#error EXTENDS undefined
#endif
-#define SMALL_SIZE (SMALL_END-SMALL_BEG)
-#define LARGE_SIZE (LARGE_END-LARGE_BEG)
+#define SMALL_SIZE (SMALL_END - SMALL_BEG)
+#define LARGE_SIZE (LARGE_END - LARGE_BEG)
-#define SMALL SMALL_BEG:SMALL_SIZE
-#define LARGE LARGE_BEG:LARGE_SIZE
+#define SMALL \
+ SMALL_BEG: \
+ SMALL_SIZE
+#define LARGE \
+ LARGE_BEG: \
+ LARGE_SIZE
int main() {
int arr[SIZE];
LARGE_SIZE * sizeof arr[0]);
// CHECK-NOT: Libomptarget
-#pragma omp target data map(alloc: arr[LARGE])
+#pragma omp target data map(alloc : arr[LARGE])
{
-#pragma omp target data map(present, tofrom: arr[SMALL])
+#pragma omp target data map(present, tofrom : arr[SMALL])
;
}
// CHECK: Libomptarget message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0*}}[[#LARGE_ADDR]] ([[#LARGE_BYTES]] bytes)
// CHECK: Libomptarget error: Call to getTargetPointer returned null pointer ('present' map type modifier).
// CHECK: Libomptarget fatal error 1: failure of target construct while offloading is mandatory
-#pragma omp target data map(alloc: arr[SMALL])
+#pragma omp target data map(alloc : arr[SMALL])
{
-#pragma omp target data map(present, tofrom: arr[LARGE])
+#pragma omp target data map(present, tofrom : arr[LARGE])
;
}
// RUN: %libomptarget-run-generic 2>&1 \
// RUN: | %fcheck-generic
-
#include <stdio.h>
int main() {
int i;
-#pragma omp target enter data map(alloc:i)
+#pragma omp target enter data map(alloc : i)
// i isn't present at the end of the target data region, but the "present"
// modifier is only checked at the beginning of a region.
-#pragma omp target data map(present, alloc: i)
+#pragma omp target data map(present, alloc : i)
{
-#pragma omp target exit data map(delete:i)
+#pragma omp target exit data map(delete : i)
}
// CHECK-NOT: Libomptarget
// RUN: %libomptarget-run-fail-generic 2>&1 \
// RUN: | %fcheck-generic
-
#include <stdio.h>
int main() {
fprintf(stderr, "addr=%p, size=%ld\n", &i, sizeof i);
// CHECK-NOT: Libomptarget
-#pragma omp target enter data map(alloc: i)
-#pragma omp target enter data map(present, alloc: i)
-#pragma omp target exit data map(delete: i)
+#pragma omp target enter data map(alloc : i)
+#pragma omp target enter data map(present, alloc : i)
+#pragma omp target exit data map(delete : i)
// CHECK: i is present
fprintf(stderr, "i is present\n");
// CHECK: Libomptarget message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] ([[#SIZE]] bytes)
// CHECK: Libomptarget error: Call to getTargetPointer returned null pointer ('present' map type modifier).
// CHECK: Libomptarget fatal error 1: failure of target construct while offloading is mandatory
-#pragma omp target enter data map(present, alloc: i)
+#pragma omp target enter data map(present, alloc : i)
// CHECK-NOT: i is present
fprintf(stderr, "i is present\n");
// RUN: %libomptarget-run-fail-generic 2>&1 \
// RUN: | %fcheck-generic
-
#include <stdio.h>
int main() {
// CHECK: addr=0x[[#%x,HOST_ADDR:]], size=[[#%u,SIZE:]]
fprintf(stderr, "addr=%p, size=%ld\n", &i, sizeof i);
- // CHECK-NOT: Libomptarget
- #pragma omp target enter data map(alloc: i)
- #pragma omp target exit data map(present, delete: i)
+// CHECK-NOT: Libomptarget
+#pragma omp target enter data map(alloc : i)
+#pragma omp target exit data map(present, delete : i)
// CHECK: i was present
fprintf(stderr, "i was present\n");
- // CHECK: Libomptarget message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] ([[#SIZE]] bytes)
- // CHECK: Libomptarget fatal error 1: failure of target construct while offloading is mandatory
- #pragma omp target exit data map(present, delete: i)
+// CHECK: Libomptarget message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] ([[#SIZE]] bytes)
+// CHECK: Libomptarget fatal error 1: failure of target construct while offloading is mandatory
+#pragma omp target exit data map(present, delete : i)
// CHECK-NOT: i was present
fprintf(stderr, "i was present\n");
// RUN: %libomptarget-run-fail-generic 2>&1 \
// RUN: | %fcheck-generic
-
#include <stdio.h>
int main() {
// CHECK: addr=0x[[#%x,HOST_ADDR:]], size=[[#%u,SIZE:]]
fprintf(stderr, "addr=%p, size=%ld\n", &i, sizeof i);
- // CHECK-NOT: Libomptarget
- #pragma omp target enter data map(alloc: i)
- #pragma omp target exit data map(present, release: i)
+// CHECK-NOT: Libomptarget
+#pragma omp target enter data map(alloc : i)
+#pragma omp target exit data map(present, release : i)
// CHECK: i was present
fprintf(stderr, "i was present\n");
- // CHECK: Libomptarget message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] ([[#SIZE]] bytes)
- // CHECK: Libomptarget fatal error 1: failure of target construct while offloading is mandatory
- #pragma omp target exit data map(present, release: i)
+// CHECK: Libomptarget message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] ([[#SIZE]] bytes)
+// CHECK: Libomptarget fatal error 1: failure of target construct while offloading is mandatory
+#pragma omp target exit data map(present, release : i)
// CHECK-NOT: i was present
fprintf(stderr, "i was present\n");
// RUN: %libomptarget-run-fail-generic 2>&1 \
// RUN: | %fcheck-generic
-
#include <stdio.h>
int main() {
fprintf(stderr, "addr=%p, size=%ld\n", &i, sizeof i);
// CHECK-NOT: Libomptarget
-#pragma omp target enter data map(alloc: i)
-#pragma omp target update CLAUSE(present: i)
-#pragma omp target exit data map(delete: i)
+#pragma omp target enter data map(alloc : i)
+#pragma omp target update CLAUSE(present : i)
+#pragma omp target exit data map(delete : i)
// CHECK: i is present
fprintf(stderr, "i is present\n");
// CHECK: Libomptarget message: device mapping required by 'present' motion modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] ([[#SIZE]] bytes)
// CHECK: Libomptarget fatal error 1: failure of target construct while offloading is mandatory
-#pragma omp target update CLAUSE(present: i)
+#pragma omp target update CLAUSE(present : i)
// CHECK-NOT: i is present
fprintf(stderr, "i is present\n");
// RUN: %libomptarget-run-fail-generic 2>&1 \
// RUN: | %fcheck-generic
-
// END.
#include <stdio.h>
#define BEFORE 0
-#define AFTER 1
+#define AFTER 1
#if EXTENDS == BEFORE
-# define SMALL 2:3
-# define LARGE 0:5
+#define SMALL 2 : 3
+#define LARGE 0 : 5
#elif EXTENDS == AFTER
-# define SMALL 0:3
-# define LARGE 0:5
+#define SMALL 0 : 3
+#define LARGE 0 : 5
#else
-# error EXTENDS undefined
+#error EXTENDS undefined
#endif
int main() {
fprintf(stderr, "addr=%p, size=%ld\n", arr, sizeof arr);
// CHECK-NOT: Libomptarget
-#pragma omp target data map(alloc: arr[LARGE])
+#pragma omp target data map(alloc : arr[LARGE])
{
-#pragma omp target update CLAUSE(present: arr[SMALL])
+#pragma omp target update CLAUSE(present : arr[SMALL])
}
// CHECK: arr is present
// CHECK: Libomptarget message: device mapping required by 'present' motion modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] ([[#SIZE]] bytes)
// CHECK: Libomptarget fatal error 1: failure of target construct while offloading is mandatory
-#pragma omp target data map(alloc: arr[SMALL])
+#pragma omp target data map(alloc : arr[SMALL])
{
-#pragma omp target update CLAUSE(present: arr[LARGE])
+#pragma omp target update CLAUSE(present : arr[LARGE])
}
// CHECK-NOT: arr is present
// REQUIRES: unified_shared_memory
-
#include <stdio.h>
// The runtime considers unified shared memory to be always present.
int i;
// CHECK-NOT: Libomptarget
-#pragma omp target data map(alloc: i)
-#pragma omp target map(present, alloc: i)
+#pragma omp target data map(alloc : i)
+#pragma omp target map(present, alloc : i)
;
// CHECK: i is present
fprintf(stderr, "i is present\n");
// CHECK-NOT: Libomptarget
-#pragma omp target map(present, alloc: i)
+#pragma omp target map(present, alloc : i)
;
// CHECK: is present
// RUN: %libomptarget-run-fail-generic 2>&1 \
// RUN: | %fcheck-generic
-
#include <stdio.h>
int main() {
fprintf(stderr, "addr=%p\n", arr);
// CHECK-NOT: Libomptarget
-#pragma omp target data map(alloc: arr[0:5])
-#pragma omp target map(present, alloc: arr[0:0])
+#pragma omp target data map(alloc : arr[0 : 5])
+#pragma omp target map(present, alloc : arr[0 : 0])
;
// CHECK: arr is present
// CHECK: Libomptarget error: Call to targetDataBegin failed, abort target.
// CHECK: Libomptarget error: Failed to process data before launching the kernel.
// CHECK: Libomptarget fatal error 1: failure of target construct while offloading is mandatory
-#pragma omp target data map(alloc: arr[0:0])
-#pragma omp target map(present, alloc: arr[0:0])
+#pragma omp target data map(alloc : arr[0 : 0])
+#pragma omp target map(present, alloc : arr[0 : 0])
;
// CHECK-NOT: arr is present
// RUN: %libomptarget-run-fail-generic 2>&1 \
// RUN: | %fcheck-generic
-
#include <stdio.h>
int main() {
fprintf(stderr, "addr=%p\n", arr);
// CHECK-NOT: Libomptarget
-#pragma omp target enter data map(alloc: arr[0:5])
-#pragma omp target exit data map(present, release: arr[0:0])
+#pragma omp target enter data map(alloc : arr[0 : 5])
+#pragma omp target exit data map(present, release : arr[0 : 0])
// CHECK: arr is present
fprintf(stderr, "arr is present\n");
//
// CHECK: Libomptarget message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] (0 bytes)
// CHECK: Libomptarget fatal error 1: failure of target construct while offloading is mandatory
-#pragma omp target enter data map(alloc: arr[0:0])
-#pragma omp target exit data map(present, release: arr[0:0])
+#pragma omp target enter data map(alloc : arr[0 : 0])
+#pragma omp target exit data map(present, release : arr[0 : 0])
// CHECK-NOT: arr is present
fprintf(stderr, "arr is present\n");
int main() {
int data1[3] = {1}, data2[3] = {2}, data3[3] = {3};
int sum[16] = {0};
-#pragma omp target teams distribute parallel for map(tofrom \
- : sum) \
+#pragma omp target teams distribute parallel for map(tofrom : sum) \
firstprivate(data1, data2, data3)
for (int i = 0; i < 16; ++i) {
for (int j = 0; j < 3; ++j) {
} DV;
void init(double vertexx[]) {
- #pragma omp target map(vertexx[0:100])
+#pragma omp target map(vertexx[0 : 100])
{
printf("In init: %lf, expected 100.0\n", vertexx[77]);
vertexx[77] = 77.0;
}
void change(DV *dvptr) {
- #pragma omp target map(dvptr->dataptr[0:100])
+#pragma omp target map(dvptr->dataptr[0 : 100])
{
printf("In change: %lf, expected 77.0\n", dvptr->dataptr[77]);
dvptr->dataptr[77] += 1.0;
DV dv;
dv.dataptr = &vertexx[0];
- #pragma omp target enter data map(to:vertexx[0:100])
+#pragma omp target enter data map(to : vertexx[0 : 100])
init(vertexx);
change(&dv);
- #pragma omp target exit data map(from:vertexx[0:100])
+#pragma omp target exit data map(from : vertexx[0 : 100])
// CHECK: Final: 78.0
printf("Final: %lf\n", vertexx[77]);
}
-
#include <stdio.h>
-void sum(int* input, int size, int* output)
-{
-#pragma omp target teams distribute parallel for reduction(+:output[0]) \
- map(to:input[0:size])
+void sum(int *input, int size, int *output) {
+#pragma omp target teams distribute parallel for reduction(+ : output[0]) \
+ map(to : input[0 : size])
for (int i = 0; i < size; i++)
output[0] += input[i];
}
-int main()
-{
+int main() {
const int size = 100;
int *array = new int[size];
int result = 0;
delete[] array;
return 0;
}
-
// RUN: %libomptarget-run-generic 2>&1 \
// RUN: | %fcheck-generic
-
// END.
#include <stdio.h>
#define BEFORE 0
-#define AFTER 1
+#define AFTER 1
#define SIZE 100
#if EXTENDS == BEFORE
-# define SMALL_BEG (SIZE-2)
-# define SMALL_END SIZE
-# define LARGE_BEG 0
-# define LARGE_END SIZE
+#define SMALL_BEG (SIZE - 2)
+#define SMALL_END SIZE
+#define LARGE_BEG 0
+#define LARGE_END SIZE
#elif EXTENDS == AFTER
-# define SMALL_BEG 0
-# define SMALL_END 2
-# define LARGE_BEG 0
-# define LARGE_END SIZE
+#define SMALL_BEG 0
+#define SMALL_END 2
+#define LARGE_BEG 0
+#define LARGE_END SIZE
#else
-# error EXTENDS undefined
+#error EXTENDS undefined
#endif
-#define SMALL SMALL_BEG:(SMALL_END-SMALL_BEG)
-#define LARGE LARGE_BEG:(LARGE_END-LARGE_BEG)
+#define SMALL \
+ SMALL_BEG: \
+ (SMALL_END - SMALL_BEG)
+#define LARGE \
+ LARGE_BEG: \
+ (LARGE_END - LARGE_BEG)
void check_not_present() {
int arr[SIZE];
// arr[LARGE] isn't (fully) present at the end of the target data region, so
// the device-to-host transfer should not be performed, or it might fail.
-#pragma omp target data map(tofrom: arr[LARGE])
+#pragma omp target data map(tofrom : arr[LARGE])
{
-#pragma omp target exit data map(delete: arr[LARGE])
-#pragma omp target enter data map(alloc: arr[SMALL])
-#pragma omp target map(alloc: arr[SMALL])
+#pragma omp target exit data map(delete : arr[LARGE])
+#pragma omp target enter data map(alloc : arr[SMALL])
+#pragma omp target map(alloc : arr[SMALL])
for (int i = SMALL_BEG; i < SMALL_END; ++i)
arr[i] = 88;
}
// arr[SMALL] is (fully) present at the end of the target data region, and the
// device-to-host transfer should be performed only for it even though more
// of the array is then present.
-#pragma omp target data map(tofrom: arr[SMALL])
+#pragma omp target data map(tofrom : arr[SMALL])
{
-#pragma omp target exit data map(delete: arr[SMALL])
-#pragma omp target enter data map(alloc: arr[LARGE])
-#pragma omp target map(alloc: arr[LARGE])
+#pragma omp target exit data map(delete : arr[SMALL])
+#pragma omp target enter data map(alloc : arr[LARGE])
+#pragma omp target map(alloc : arr[LARGE])
for (int i = LARGE_BEG; i < LARGE_END; ++i)
arr[i] = 88;
}
A = (float *)omp_target_alloc((FROM + LENGTH) * sizeof(float), device_id);
float *A_dev = NULL;
-#pragma omp target has_device_addr(A [FROM:LENGTH]) map(A_dev)
+#pragma omp target has_device_addr(A[FROM : LENGTH]) map(A_dev)
{ A_dev = A; }
// CHECK: Success
if (A_dev == NULL || A_dev != A)
short *xp = &x[0];
x[1] = 111;
-#pragma omp target data map(tofrom : xp [0:2]) use_device_addr(xp [0:2])
-#pragma omp target has_device_addr(xp [0:2])
+#pragma omp target data map(tofrom : xp[0 : 2]) use_device_addr(xp[0 : 2])
+#pragma omp target has_device_addr(xp[0 : 2])
{
xp[1] = 222;
// CHECK: 222
short a[10], b[10];
a[1] = 111;
b[1] = 111;
-#pragma omp target data map(to : a [0:2], b [0:2]) use_device_addr(a, b)
+#pragma omp target data map(to : a[0 : 2], b[0 : 2]) use_device_addr(a, b)
#pragma omp target has_device_addr(a) has_device_addr(b[0])
{
a[1] = 222;
// RUN: %libomptarget-run-generic 2>&1 \
// RUN: | %fcheck-generic
-
// END.
#include <omp.h>
int main() {
int arr[100];
-#pragma omp target data map(alloc: arr[50:2]) // partially mapped
+#pragma omp target data map(alloc : arr[50 : 2]) // partially mapped
{
// CHECK: arr[50] must present: 1
fprintf(stderr, "arr[50] must present: %d\n",
arr[50] = 5;
arr[51] = 6;
} // must treat as present (dec ref count) even though full size not present
- } // wouldn't delete if previous ref count dec didn't happen
+ } // wouldn't delete if previous ref count dec didn't happen
// CHECK: arr[50] still present: 0
fprintf(stderr, "arr[50] still present: %d\n",
// CHECK: 10 111
printf("%d %ld %p %p %p %p\n", dvb1.c[0].b.a[0], dvb1.c[0].b.d1, &dvb1,
&dvb1.c[0], &dvb1.c[0].b, &dvb1.c[0].b.a[0]);
-#pragma omp target map(to \
- : dvb1, dvb1.c [0:2]) \
- map(tofrom \
- : dvb1.c[0].b.a [0:10], dvb1.c[1].b.a [0:10])
+#pragma omp target map(to : dvb1, dvb1.c[0 : 2]) \
+ map(tofrom : dvb1.c[0].b.a[0 : 10], dvb1.c[1].b.a[0 : 10])
{
// CHECK: 10 111
printf("%d %ld %p %p %p %p\n", dvb1.c[0].b.a[0], dvb1.c[0].b.d1, &dvb1,
// RUN: %libomptarget-run-generic 2>&1 \
// RUN: | %fcheck-generic
-
// END.
#include <stdio.h>
#define BEFORE 0
-#define AFTER 1
+#define AFTER 1
#if EXTENDS == BEFORE
-# define SMALL 2:3
-# define LARGE 0:5
+#define SMALL 2 : 3
+#define LARGE 0 : 5
#elif EXTENDS == AFTER
-# define SMALL 0:3
-# define LARGE 0:5
+#define SMALL 0 : 3
+#define LARGE 0 : 5
#else
-# error EXTENDS undefined
+#error EXTENDS undefined
#endif
int main() {
int arr[5];
// CHECK-NOT: Libomptarget
-#pragma omp target data map(alloc: arr[LARGE])
+#pragma omp target data map(alloc : arr[LARGE])
{
#pragma omp target update CLAUSE(arr[SMALL])
}
fprintf(stderr, "success\n");
// CHECK-NOT: Libomptarget
-#pragma omp target data map(alloc: arr[SMALL])
+#pragma omp target data map(alloc : arr[SMALL])
{
#pragma omp target update CLAUSE(arr[LARGE])
}
x[1] = 111;
printf("%d, %p\n", xp[1], &xp[1]);
-#pragma omp target data use_device_addr(xp [1:3]) map(tofrom : x)
+#pragma omp target data use_device_addr(xp[1 : 3]) map(tofrom : x)
#pragma omp target is_device_ptr(xp)
{ xp[1] = 222; }
// CHECK: 222
// CHECK-NEXT: signed: xs=[[#NUM_THREADS-1]]{{$}}
int xs = -1;
int numThreads;
- #pragma omp target parallel for num_threads(NUM_THREADS_TRY) \
- map(tofrom:xs, numThreads)
+#pragma omp target parallel for num_threads(NUM_THREADS_TRY) \
+ map(tofrom : xs, numThreads)
for (int i = 0; i < omp_get_num_threads(); ++i) {
- #pragma omp atomic compare
- if (xs < i) { xs = i; }
+#pragma omp atomic compare
+ if (xs < i) {
+ xs = i;
+ }
if (i == 0)
numThreads = omp_get_num_threads();
}
// CHECK-NEXT: unsigned: xu=0x0{{$}}
unsigned xu = UINT_MAX;
- #pragma omp target parallel for num_threads(NUM_THREADS_TRY) \
- map(tofrom:xu)
+#pragma omp target parallel for num_threads(NUM_THREADS_TRY) map(tofrom : xu)
for (int i = 0; i < omp_get_num_threads(); ++i) {
- #pragma omp atomic compare
- if (xu > i) { xu = i; }
+#pragma omp atomic compare
+ if (xu > i) {
+ xu = i;
+ }
}
printf("unsigned: xu=0x%x\n", xu);
return 0;
sum_host += array[i];
}
-#pragma omp target teams distribute parallel for map(to: array[:size]) \
- reduction(+ : sum)
+#pragma omp target teams distribute parallel for map(to : array[ : size]) \
+ reduction(+ : sum)
for (int i = 0; i < size; i++)
sum += array[i];
std::cout << "hierarchical parallelism" << std::endl;
const int nblock(10), block_size(10);
T block_sum[nblock];
-#pragma omp target teams distribute map(to \
- : array[:size]) \
- map(from \
- : block_sum[:nblock])
+#pragma omp target teams distribute map(to : array[ : size]) \
+ map(from : block_sum[ : nblock])
for (int ib = 0; ib < nblock; ib++) {
T partial_sum = 0;
const int istart = ib * block_size;
// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
#if ADD_REDUCTION
-# define REDUCTION(...) reduction(__VA_ARGS__)
+#define REDUCTION(...) reduction(__VA_ARGS__)
#else
-# define REDUCTION(...)
+#define REDUCTION(...)
#endif
#include <stdio.h>
int main() {
int x = 0, y = 1;
- #pragma omp target teams num_teams(1) map(tofrom:x, y) REDUCTION(+:x)
+#pragma omp target teams num_teams(1) map(tofrom : x, y) REDUCTION(+ : x)
{
x += 5;
- #pragma omp parallel
+#pragma omp parallel
y = 6;
}
// CHECK: 5, 6
s.t.p[i] = i;
}
-#pragma omp target enter data map(to : s, s.t.p[:N])
+#pragma omp target enter data map(to : s, s.t.p[ : N])
#pragma omp target
{
}
}
-#pragma omp target update from(s.t.p[:N])
+#pragma omp target update from(s.t.p[ : N])
for (int i = 0; i < N; ++i) {
assert(s.t.p[i] == 2 * i);
s.t.p[i] += i;
}
-#pragma omp target update to(s.t.p[:N])
+#pragma omp target update to(s.t.p[ : N])
#pragma omp target
{
}
}
-#pragma omp target exit data map(from : s, s.t.p[:N])
+#pragma omp target exit data map(from : s, s.t.p[ : N])
for (int i = 0; i < N; ++i) {
assert(s.t.p[i] == 4 * i);
// CHECK-NOT: {{.}}
int main() {
int x = 0;
- #pragma omp target teams num_teams(2) reduction(+:x)
+#pragma omp target teams num_teams(2) reduction(+ : x)
x += 2;
printf("Hello World: %d\n", x);
return 0;
assert(src_ptr && "src_ptr is NULL");
assert(dst_ptr && "dst_ptr is NULL");
-#pragma omp target teams distribute parallel for device(src_device) \
- is_device_ptr(src_ptr)
+#pragma omp target teams distribute parallel for device(src_device) \
+ is_device_ptr(src_ptr)
for (int i = 0; i < N; ++i) {
src_ptr[i] = magic_num;
}
assert(buffer && "failed to allocate host buffer");
-#pragma omp target teams distribute parallel for device(dst_device) \
- map(from: buffer[0:N]) is_device_ptr(dst_ptr)
+#pragma omp target teams distribute parallel for device(dst_device) \
+ map(from : buffer[0 : N]) is_device_ptr(dst_ptr)
for (int i = 0; i < N; ++i) {
buffer[i] = dst_ptr[i] + magic_num;
}
printf("dlopen() failed: %s\n", dlerror());
return 1;
}
- Foo = (int (*)(void)) dlsym(Handle, "foo");
+ Foo = (int (*)(void))dlsym(Handle, "foo");
if (Handle == NULL) {
printf("dlsym() failed: %s\n", dlerror());
return 1;
// XFAIL: amdgcn-amd-amdhsa
// XFAIL: amdgcn-amd-amdhsa-LTO
-#include <stdio.h>
#include <omp.h>
+#include <stdio.h>
static void check(char *X, int Dev) {
printf(" host X = %c\n", *X);
- #pragma omp target device(Dev)
+#pragma omp target device(Dev)
printf("device X = %c\n", *X);
}
// CHECK: host X = h
// CHECK-NEXT: device X = d
char X = 'd';
- #pragma omp target enter data map(to:X)
+#pragma omp target enter data map(to : X)
X = 'h';
CHECK_DATA();
- //--------------------------------------------------
- // Check behavior when specifying host directly.
- //--------------------------------------------------
+//--------------------------------------------------
+// Check behavior when specifying host directly.
+//--------------------------------------------------
- // CHECK-NEXT: omp_is_initial_device() = 1
- // CHECK-NEXT: host X = h
- // CHECK-NEXT: device X = d
- #pragma omp target device(DevInit) map(always,tofrom:X)
+// CHECK-NEXT: omp_is_initial_device() = 1
+// CHECK-NEXT: host X = h
+// CHECK-NEXT: device X = d
+#pragma omp target device(DevInit) map(always, tofrom : X)
printf("omp_is_initial_device() = %d\n", omp_is_initial_device());
CHECK_DATA();
- // CHECK-NEXT: omp_is_initial_device() = 1
- // CHECK-NEXT: host X = h
- // CHECK-NEXT: device X = d
- #pragma omp target teams device(DevInit) num_teams(1) map(always,tofrom:X)
+// CHECK-NEXT: omp_is_initial_device() = 1
+// CHECK-NEXT: host X = h
+// CHECK-NEXT: device X = d
+#pragma omp target teams device(DevInit) num_teams(1) map(always, tofrom : X)
printf("omp_is_initial_device() = %d\n", omp_is_initial_device());
CHECK_DATA();
- // Check that __kmpc_push_target_tripcount_mapper doesn't fail. I'm not sure
- // how to check that it actually pushes to the initial device.
- #pragma omp target teams device(DevInit) num_teams(1)
- #pragma omp distribute
+// Check that __kmpc_push_target_tripcount_mapper doesn't fail. I'm not sure
+// how to check that it actually pushes to the initial device.
+#pragma omp target teams device(DevInit) num_teams(1)
+#pragma omp distribute
for (int i = 0; i < 2; ++i)
- ;
+ ;
- // CHECK-NEXT: host X = h
- // CHECK-NEXT: device X = d
- #pragma omp target data device(DevInit) map(always,tofrom:X)
+// CHECK-NEXT: host X = h
+// CHECK-NEXT: device X = d
+#pragma omp target data device(DevInit) map(always, tofrom : X)
;
CHECK_DATA();
- // CHECK-NEXT: host X = h
- // CHECK-NEXT: device X = d
- #pragma omp target enter data device(DevInit) map(always,to:X)
+// CHECK-NEXT: host X = h
+// CHECK-NEXT: device X = d
+#pragma omp target enter data device(DevInit) map(always, to : X)
;
CHECK_DATA();
- // CHECK-NEXT: host X = h
- // CHECK-NEXT: device X = d
- #pragma omp target exit data device(DevInit) map(always,from:X)
+// CHECK-NEXT: host X = h
+// CHECK-NEXT: device X = d
+#pragma omp target exit data device(DevInit) map(always, from : X)
;
CHECK_DATA();
- // CHECK-NEXT: host X = h
- // CHECK-NEXT: device X = d
- #pragma omp target update device(DevInit) to(X)
+// CHECK-NEXT: host X = h
+// CHECK-NEXT: device X = d
+#pragma omp target update device(DevInit) to(X)
;
CHECK_DATA();
- // CHECK-NEXT: host X = h
- // CHECK-NEXT: device X = d
- #pragma omp target update device(DevInit) from(X)
+// CHECK-NEXT: host X = h
+// CHECK-NEXT: device X = d
+#pragma omp target update device(DevInit) from(X)
;
CHECK_DATA();
omp_set_default_device(DevInit);
- // CHECK-NEXT: omp_is_initial_device() = 1
- // CHECK-NEXT: host X = h
- // CHECK-NEXT: device X = d
- #pragma omp target map(always,tofrom:X)
+// CHECK-NEXT: omp_is_initial_device() = 1
+// CHECK-NEXT: host X = h
+// CHECK-NEXT: device X = d
+#pragma omp target map(always, tofrom : X)
printf("omp_is_initial_device() = %d\n", omp_is_initial_device());
CHECK_DATA();
- // CHECK-NEXT: omp_is_initial_device() = 1
- // CHECK-NEXT: host X = h
- // CHECK-NEXT: device X = d
- #pragma omp target teams num_teams(1) map(always,tofrom:X)
+// CHECK-NEXT: omp_is_initial_device() = 1
+// CHECK-NEXT: host X = h
+// CHECK-NEXT: device X = d
+#pragma omp target teams num_teams(1) map(always, tofrom : X)
printf("omp_is_initial_device() = %d\n", omp_is_initial_device());
CHECK_DATA();
- // Check that __kmpc_push_target_tripcount_mapper doesn't fail. I'm not sure
- // how to check that it actually pushes to the initial device.
- #pragma omp target teams num_teams(1)
- #pragma omp distribute
+// Check that __kmpc_push_target_tripcount_mapper doesn't fail. I'm not sure
+// how to check that it actually pushes to the initial device.
+#pragma omp target teams num_teams(1)
+#pragma omp distribute
for (int i = 0; i < 2; ++i)
- ;
+ ;
- // CHECK-NEXT: host X = h
- // CHECK-NEXT: device X = d
- #pragma omp target data map(always,tofrom:X)
+// CHECK-NEXT: host X = h
+// CHECK-NEXT: device X = d
+#pragma omp target data map(always, tofrom : X)
;
CHECK_DATA();
- // CHECK-NEXT: host X = h
- // CHECK-NEXT: device X = d
- #pragma omp target enter data map(always,to:X)
+// CHECK-NEXT: host X = h
+// CHECK-NEXT: device X = d
+#pragma omp target enter data map(always, to : X)
;
CHECK_DATA();
- // CHECK-NEXT: host X = h
- // CHECK-NEXT: device X = d
- #pragma omp target exit data map(always,from:X)
+// CHECK-NEXT: host X = h
+// CHECK-NEXT: device X = d
+#pragma omp target exit data map(always, from : X)
;
CHECK_DATA();
- // CHECK-NEXT: host X = h
- // CHECK-NEXT: device X = d
- #pragma omp target update to(X)
+// CHECK-NEXT: host X = h
+// CHECK-NEXT: device X = d
+#pragma omp target update to(X)
;
CHECK_DATA();
- // CHECK-NEXT: host X = h
- // CHECK-NEXT: device X = d
- #pragma omp target update from(X)
+// CHECK-NEXT: host X = h
+// CHECK-NEXT: device X = d
+#pragma omp target update from(X)
;
CHECK_DATA();
// RUN: %fcheck-nvptx64-nvidia-cuda -allow-empty -check-prefix=INFO
// REQUIRES: nvptx64-nvidia-cuda
-#include <stdio.h>
#include <omp.h>
+#include <stdio.h>
#define N 64
// INFO: Libomptarget device 0 info: OpenMP Host-Device pointer mappings after block at info.c:[[#%u,]]:[[#%u,]]:
// INFO: Libomptarget device 0 info: Host Ptr Target Ptr Size (B) DynRefCount HoldRefCount Declaration
// INFO: Libomptarget device 0 info: [[#%#x,]] [[#%#x,]] 4 INF 0 global at unknown:0:0
-#pragma omp target data map(alloc:A[0:N]) map(ompx_hold,tofrom:B[0:N]) map(to:C[0:N])
+#pragma omp target data map(alloc : A[0 : N]) \
+ map(ompx_hold, tofrom : B[0 : N]) map(to : C[0 : N])
#pragma omp target firstprivate(val)
{ val = 1; }
__tgt_set_info_flag(0x0);
// INFO-NOT: Libomptarget device 0 info: {{.*}}
#pragma omp target
- { }
+ {}
return 0;
}
int main() {
// CHECK: x = 98
int x = 98;
- #pragma omp target exit data map(from:x)
+#pragma omp target exit data map(from : x)
printf("x = %d\n", x);
return 0;
}
/*
Test for looptripcount being popped from runtime stack.
*/
-#include <stdio.h>
#include <omp.h>
-int main()
-{
+#include <stdio.h>
+int main() {
int N = 128;
int NN = 1024;
int num_teams[NN];
printf("#pragma omp target teams distribute parallel for thread_limit(4)\n");
#pragma omp target teams distribute parallel for thread_limit(4)
- for (int j = 0; j< N; j++) {
+ for (int j = 0; j < N; j++) {
num_threads[j] = omp_get_num_threads();
num_teams[j] = omp_get_num_teams();
}
printf("num_threads %d num_teams %d\n", num_threads[0], num_teams[0]);
-// DEBUG: loop trip count is 128
+ // DEBUG: loop trip count is 128
printf("#pragma omp target teams distribute parallel for\n");
#pragma omp target teams distribute parallel for
- for (int j = 0; j< N; j++) {
+ for (int j = 0; j < N; j++) {
num_threads[j] = omp_get_num_threads();
num_teams[j] = omp_get_num_teams();
}
printf("num_threads %d num_teams %d\n", num_threads[0], num_teams[0]);
-// DEBUG: loop trip count is 128
+ // DEBUG: loop trip count is 128
return 0;
}
// CHECK: Libomptarget fatal error 1: failure of target construct while offloading is mandatory
int main(void) {
int X;
- #pragma omp DIR device(omp_get_initial_device())
+#pragma omp DIR device(omp_get_initial_device())
;
return 0;
}
}
int buffer[n];
#pragma omp target teams distribute parallel for is_device_ptr(p) \
- map(from \
- : buffer)
+ map(from : buffer)
for (int j = 0; j < n; ++j) {
buffer[j] = p[j];
}
// RUN: %libomptarget-compile-generic && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-generic 2>&1 | %fcheck-generic -allow-empty -check-prefix=DEBUG
// REQUIRES: libomptarget-debug
+#include <cassert>
#include <cstdio>
#include <cstdlib>
-#include <cassert>
// Data structure definitions copied from OpenMP RTL.
struct __tgt_target_non_contig {
int64_t stride;
};
-enum tgt_map_type {
- OMP_TGT_MAPTYPE_NON_CONTIG = 0x100000000000
-};
+enum tgt_map_type { OMP_TGT_MAPTYPE_NON_CONTIG = 0x100000000000 };
// OpenMP RTL interfaces
#ifdef __cplusplus
// DEBUG: offset 456
return 0;
}
-
// RUN: %libomptarget-compile-run-and-check-generic
-#include <stdio.h>
#include <omp.h>
+#include <stdio.h>
int main(void) {
int isHost = -1;
-#pragma omp target map(from: isHost)
+#pragma omp target map(from : isHost)
{ isHost = omp_is_initial_device(); }
if (isHost < 0) {
// RUN: %libomptarget-compilexx-run-and-check-generic
-#include <stdio.h>
#include <omp.h>
+#include <stdio.h>
int main(void) {
int isHost = 0;
-#pragma omp target map(from: isHost)
+#pragma omp target map(from : isHost)
{ isHost = omp_is_initial_device(); }
if (isHost < 0) {
(for whatever reason) the set must be consistent with previously
set values.
*/
-#include <stdio.h>
#include <omp.h>
+#include <stdio.h>
// ---------------------------------------------------------------------------
// Various definitions copied from OpenMP RTL
// CHECK-NEXT: update: x=[[#NUM_TEAMS]]{{$}}
int x = 0;
int numTeams;
- #pragma omp target teams num_teams(NUM_TEAMS_TRY) map(tofrom:x, numTeams)
+#pragma omp target teams num_teams(NUM_TEAMS_TRY) map(tofrom : x, numTeams)
{
- #pragma omp atomic update
+#pragma omp atomic update
++x;
if (omp_get_team_num() == 0)
numTeams = omp_get_num_teams();
bool xCaptured[numTeams];
memset(xCaptured, 0, sizeof xCaptured);
x = 0;
- #pragma omp target teams num_teams(NUM_TEAMS_TRY) map(tofrom:x, numTeams)
+#pragma omp target teams num_teams(NUM_TEAMS_TRY) map(tofrom : x, numTeams)
{
int v;
- #pragma omp atomic capture
+#pragma omp atomic capture
v = x++;
xCaptured[v] = true;
}
#pragma omp target update from(A) depend(out : A[0]) nowait
// B updated via exit, A just released
-#pragma omp target exit data map(release \
- : A) map(from \
- : B) depend(out \
- : A[0]) nowait
+#pragma omp target exit data map(release : A) map(from : B) depend(out : A[0]) \
+ nowait
} // if
} // parallel
return Sum != 2 * N * (2 + N - 1 + 2) / 2;
}
-
for (int i = 0; i < 1024; ++i)
data[i] = i;
-#pragma omp target map(tofrom: sum) map(to: data) depend(inout : data[0]) nowait
+#pragma omp target map(tofrom : sum) map(to : data) depend(inout : data[0]) \
+ nowait
{
for (int i = 0; i < 1024; ++i) {
sum += data[i];
}
}
-#pragma omp target map(tofrom: sum) map(to: data) depend(inout : data[0])
+#pragma omp target map(tofrom : sum) map(to : data) depend(inout : data[0])
{
for (int i = 0; i < 1024; ++i) {
sum += data[i];
// std::equal is lowered to libc function memcmp.
void test_memcpy() {
int r = 0;
-#pragma omp target map(from: r)
+#pragma omp target map(from : r)
{
int x[2] = {0, 0};
int y[2] = {0, 0};
int main(int argc, char *argv[]) {
int *data = (int *)malloc(N * sizeof(int));
-#pragma omp target map(from: data[0:N])
+#pragma omp target map(from : data[0 : N])
{
double start = omp_get_wtime();
for (int i = 0; i < N; ++i)
}
// CHECK: duration: {{.+[1-9]+}}
-
// UNSUPPORTED: amdgcn-amd-amdhsa
// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
-#include <stdio.h>
#include <omp.h>
+#include <stdio.h>
// ---------------------------------------------------------------------------
// Various definitions copied from OpenMP RTL
int rc = omp_target_associate_ptr(&x, x_dev, sizeof x, 0, dev);
assert(!rc && "expected omp_target_associate_ptr to succeed");
- // To determine whether x needs to be transfered, the runtime cannot simply
- // check whether unified shared memory is enabled and the 'close' modifier is
- // specified. It must check whether x was previously placed in device memory
- // by, for example, omp_target_associate_ptr.
- #pragma omp target map(always, tofrom: x)
+// To determine whether x needs to be transfered, the runtime cannot simply
+// check whether unified shared memory is enabled and the 'close' modifier is
+// specified. It must check whether x was previously placed in device memory
+// by, for example, omp_target_associate_ptr.
+#pragma omp target map(always, tofrom : x)
x += 1;
// CHECK: x=11
int dev = omp_get_default_device();
struct S s = {10, 20};
- #pragma omp target enter data map(close, to: s)
- #pragma omp target map(alloc: s)
+#pragma omp target enter data map(close, to : s)
+#pragma omp target map(alloc : s)
{
s.x = 11;
s.y = 21;
}
- // To determine whether x needs to be transfered or deleted, the runtime
- // cannot simply check whether unified shared memory is enabled and the
- // 'close' modifier is specified. It must check whether x was previously
- // placed in device memory by, for example, a 'close' modifier that isn't
- // specified here. The following struct member case checks a special code
- // path in the runtime implementation where members are transferred before
- // deletion of the struct.
- #pragma omp target exit data map(from: s.x, s.y)
+// To determine whether x needs to be transfered or deleted, the runtime
+// cannot simply check whether unified shared memory is enabled and the
+// 'close' modifier is specified. It must check whether x was previously
+// placed in device memory by, for example, a 'close' modifier that isn't
+// specified here. The following struct member case checks a special code
+// path in the runtime implementation where members are transferred before
+// deletion of the struct.
+#pragma omp target exit data map(from : s.x, s.y)
// CHECK: s.x=11, s.y=21
printf("s.x=%d, s.y=%d\n", s.x, s.y);
// amdgpu runtime crash
// UNSUPPORTED: amdgcn-amd-amdhsa
-
#include <omp.h>
#include <stdio.h>
// Test that updates on the device are not visible to host
// when only a TO mapping is used.
//
-#pragma omp target map(tofrom \
- : device_data, device_alloc) map(close, to \
- : alloc[:N], data \
- [:N])
+#pragma omp target map(tofrom : device_data, device_alloc) \
+ map(close, to : alloc[ : N], data[ : N])
{
device_data = &data[0];
device_alloc = &alloc[0];
data[i] += 1;
}
-#pragma omp target map(close, tofrom : alloc[:N], data[:N])
+#pragma omp target map(close, tofrom : alloc[ : N], data[ : N])
{
// CHECK: Alloc device values are correct: Succeeded
fails = 0;
// amdgpu runtime crash
// UNSUPPORTED: amdgcn-amd-amdhsa
-#include <stdio.h>
#include <omp.h>
+#include <stdio.h>
// ---------------------------------------------------------------------------
// Various definitions copied from OpenMP RTL