diff options
Diffstat (limited to 'libgomp/testsuite/libgomp.c')
7 files changed, 263 insertions, 0 deletions
diff --git a/libgomp/testsuite/libgomp.c/declare-variant-3-sm61.c b/libgomp/testsuite/libgomp.c/declare-variant-3-sm61.c new file mode 100644 index 0000000..e6941d3 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm61.c @@ -0,0 +1,8 @@ +/* { dg-do link { target { offload_target_nvptx } } } */ +/* { dg-additional-options -foffload=nvptx-none } */ +/* { dg-additional-options "-foffload=-misa=sm_61 -foffload=-mptx=_" } */ +/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */ + +#include "declare-variant-3.h" + +/* { dg-final { only_for_offload_target nvptx-none scan-offload-tree-dump "= f61 \\(\\);" "optimized" } } */ diff --git a/libgomp/testsuite/libgomp.c/declare-variant-3.h b/libgomp/testsuite/libgomp.c/declare-variant-3.h index c9c8f4a..f5695a2 100644 --- a/libgomp/testsuite/libgomp.c/declare-variant-3.h +++ b/libgomp/testsuite/libgomp.c/declare-variant-3.h @@ -37,6 +37,13 @@ f53 (void) __attribute__ ((noipa)) int +f61 (void) +{ + return 61; +} + +__attribute__ ((noipa)) +int f70 (void) { return 70; @@ -68,6 +75,7 @@ f89 (void) #pragma omp declare variant (f37) match (device={isa("sm_37")}) #pragma omp declare variant (f52) match (device={isa("sm_52")}) #pragma omp declare variant (f53) match (device={isa("sm_53")}) +#pragma omp declare variant (f61) match (device={isa("sm_61")}) #pragma omp declare variant (f70) match (device={isa("sm_70")}) #pragma omp declare variant (f75) match (device={isa("sm_75")}) #pragma omp declare variant (f80) match (device={isa("sm_80")}) diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx942.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx942.c new file mode 100644 index 0000000..d1df550 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx942.c @@ -0,0 +1,8 @@ +/* { dg-do link { target { offload_target_amdgcn } } } */ +/* { dg-additional-options -foffload=amdgcn-amdhsa } */ +/* { dg-additional-options -foffload=-march=gfx942 } */ +/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */ + +#include "declare-variant-4.h" + +/* { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "= gfx942 \\(\\);" "optimized" } } */ diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4.h b/libgomp/testsuite/libgomp.c/declare-variant-4.h index 53788d2..2257f4c 100644 --- a/libgomp/testsuite/libgomp.c/declare-variant-4.h +++ b/libgomp/testsuite/libgomp.c/declare-variant-4.h @@ -37,6 +37,13 @@ gfx90c (void) __attribute__ ((noipa)) int +gfx942 (void) +{ + return 0x942; +} + +__attribute__ ((noipa)) +int gfx1030 (void) { return 0x1030; @@ -68,6 +75,7 @@ gfx1103 (void) #pragma omp declare variant(gfx908) match(device = {isa("gfx908")}) #pragma omp declare variant(gfx90a) match(device = {isa("gfx90a")}) #pragma omp declare variant(gfx90c) match(device = {isa("gfx90c")}) +#pragma omp declare variant(gfx942) match(device = {isa("gfx942")}) #pragma omp declare variant(gfx1030) match(device = {isa("gfx1030")}) #pragma omp declare variant(gfx1036) match(device = {isa("gfx1036")}) #pragma omp declare variant(gfx1100) match(device = {isa("gfx1100")}) diff --git a/libgomp/testsuite/libgomp.c/target-map-zero-sized-2.c b/libgomp/testsuite/libgomp.c/target-map-zero-sized-2.c new file mode 100644 index 0000000..3220828 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-map-zero-sized-2.c @@ -0,0 +1,74 @@ +int +main () +{ + int i, n; + int data[] = {1,2}; + struct S { int **ptrset; }; + +// ----------------------------------- + +/* The produced mapping for sptr1->ptrset[i][:n] + + GOMP_MAP_STRUCT (size = 1) + GOMP_MAP_ZERO_LEN_ARRAY_SECTION + GOMP_MAP_ZERO_LEN_ARRAY_SECTION + GOMP_MAP_ATTACH + GOMP_MAP_ATTACH -> attaching to 2nd GOMP_MAP_ZERO_LEN_ARRAY_SECTION + +which get split into 3 separate map_vars call; in particular, +the latter is separate and points to an unmpapped variable. + +Thus, it failed with: + libgomp: pointer target not mapped for attach */ + + struct S s1, *sptr1; + s1.ptrset = (int **) __builtin_malloc (sizeof(void*) * 3); + s1.ptrset[0] = data; + s1.ptrset[1] = data; + s1.ptrset[2] = data; + sptr1 = &s1; + + i = 1; + n = 0; + #pragma omp target enter data map(sptr1[:1], sptr1->ptrset[:3]) + #pragma omp target enter data map(sptr1->ptrset[i][:n]) + + #pragma omp target exit data map(sptr1->ptrset[i][:n]) + #pragma omp target exit data map(sptr1[:1], sptr1->ptrset[:3]) + + __builtin_free (s1.ptrset); + +// ----------------------------------- + +/* The produced mapping for sptr2->ptrset[i][:n] is similar: + + GOMP_MAP_STRUCT (size = 1) + GOMP_MAP_ZERO_LEN_ARRAY_SECTION + GOMP_MAP_TO ! this one has now a finite size + GOMP_MAP_ATTACH + GOMP_MAP_ATTACH -> attach to the GOMP_MAP_TO + +As the latter GOMP_MAP_ATTACH has now a pointer target, +the attachment worked. */ + + struct S s2, *sptr2; + s2.ptrset = (int **) __builtin_malloc (sizeof(void*) * 3); + s2.ptrset[0] = data; + s2.ptrset[1] = data; + s2.ptrset[2] = data; + sptr2 = &s2; + + i = 1; + n = 2; + #pragma omp target enter data map(sptr2[:1], sptr2->ptrset[:3]) + #pragma omp target enter data map(sptr2->ptrset[i][:n]) + + #pragma omp target + if (sptr2->ptrset[1][0] != 1 || sptr2->ptrset[1][1] != 2) + __builtin_abort (); + + #pragma omp target exit data map(sptr2->ptrset[i][:n]) + #pragma omp target exit data map(sptr2[:1], sptr2->ptrset[:3]) + + __builtin_free (s2.ptrset); +} diff --git a/libgomp/testsuite/libgomp.c/target-map-zero-sized-3.c b/libgomp/testsuite/libgomp.c/target-map-zero-sized-3.c new file mode 100644 index 0000000..580c6ad --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-map-zero-sized-3.c @@ -0,0 +1,50 @@ +int +main () +{ + int i, n; + int data[] = {1,2}; + struct S { + int **ptrset; + int **ptrset2; + }; + + /* This is the same as target-map-zero-sized-3.c, but by mixing + mapped and non-mapped items, the mapping before the ATTACH + might (or here: is) not actually associated with the the + pointer used for attaching. Thus, if one does a simple + + if (openmp_p + && (pragma_kind & GOMP_MAP_VARS_ENTER_DATA) + && mapnum == 1) + check in target.c's gomp_map_vars_internal will fail + as mapnum > 1 but still the map associated with this + ATTACH is in a different set. */ + + struct S s1, *sptr1; + s1.ptrset = (int **) __builtin_malloc (sizeof(void*) * 3); + s1.ptrset2 = (int **) __builtin_malloc (sizeof(void*) * 3); + s1.ptrset[0] = data; + s1.ptrset[1] = data; + s1.ptrset[2] = data; + s1.ptrset2[0] = data; + s1.ptrset2[1] = data; + s1.ptrset2[2] = data; + sptr1 = &s1; + + i = 1; + n = 0; + #pragma omp target enter data map(data) + #pragma omp target enter data map(sptr1[:1], sptr1->ptrset[:3], sptr1->ptrset2[:3]) + #pragma omp target enter data map(sptr1->ptrset[i][:n], sptr1->ptrset2[i][:n]) + + #pragma omp target map(sptr1->ptrset[i][:n], sptr1->ptrset2[i][:n]) + if (sptr1->ptrset2[1][0] != 1 || sptr1->ptrset2[1][1] != 2) + __builtin_abort (); + + #pragma omp target exit data map(sptr1->ptrset[i][:n], sptr1->ptrset2[i][:n]) + #pragma omp target exit data map(sptr1[:1], sptr1->ptrset[:3], sptr1->ptrset2[:3]) + #pragma omp target exit data map(data) + + __builtin_free (s1.ptrset); + __builtin_free (s1.ptrset2); +} diff --git a/libgomp/testsuite/libgomp.c/target-map-zero-sized.c b/libgomp/testsuite/libgomp.c/target-map-zero-sized.c new file mode 100644 index 0000000..7c4ab80 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-map-zero-sized.c @@ -0,0 +1,107 @@ +/* { dg-do run } */ +/* { dg-additional-options "-O0" } */ + +/* Issue showed up in the real world when large data was distributed + over multiple MPI progresses - such that for one process n == 0 + happend at run time. + + Before map(var[:0]) and map(var[:n]) with n > 0 was handled, + this patch now also handles map(var[:n]) with n == 0. + + Failed before with "libgomp: pointer target not mapped for attach". */ + +/* Here, the base address is shifted - which should have no effect, + but must work as well. */ +void +with_offset () +{ + struct S { + int *ptr1, *ptr2; + }; + struct S s1, s2; + int *a, *b, *c, *d; + s1.ptr1 = (int *) 0L; + s1.ptr2 = (int *) 0xdeedbeef; + s2.ptr1 = (int *) 0L; + s2.ptr2 = (int *) 0xdeedbeef; + a = (int *) 0L; + b = (int *) 0xdeedbeef; + c = (int *) 0L; + d = (int *) 0xdeedbeef; + + int n1, n2, n3, n4; + n1 = n2 = n3 = n4 = 0; + + #pragma omp target enter data map(s1.ptr1[4:n1], s1.ptr2[6:n2], a[3:n3], b[2:n4]) + + #pragma omp target map(s2.ptr1[4:n1], s2.ptr2[2:n2], c[6:n3], d[9:n4]) + { + if (s2.ptr1 != (void *) 0L || s2.ptr2 != (void *) 0xdeedbeef + || c != (void *) 0L || d != (void *) 0xdeedbeef) + __builtin_abort (); + } + + #pragma omp target map(s1.ptr1[4:n1], s1.ptr2[6:n2], a[3:n3], b[2:n4]) + { + if (s1.ptr1 != (void *) 0L || s1.ptr2 != (void *) 0xdeedbeef + || a != (void *) 0L || b != (void *) 0xdeedbeef) + __builtin_abort (); + } + + #pragma omp target + { + if (s1.ptr1 != (void *) 0L || s1.ptr2 != (void *) 0xdeedbeef + || a != (void *) 0L || b != (void *) 0xdeedbeef) + __builtin_abort (); + } + + #pragma omp target exit data map(s1.ptr1[4:n1], s1.ptr2[6:n2], a[3:n3], b[2:n4]) +} + +int +main () +{ + struct S { + int *ptr1, *ptr2; + }; + struct S s1, s2; + int *a, *b, *c, *d; + s1.ptr1 = (int *) 0L; + s1.ptr2 = (int *) 0xdeedbeef; + s2.ptr1 = (int *) 0L; + s2.ptr2 = (int *) 0xdeedbeef; + a = (int *) 0L; + b = (int *) 0xdeedbeef; + c = (int *) 0L; + d = (int *) 0xdeedbeef; + + int n1, n2, n3, n4; + n1 = n2 = n3 = n4 = 0; + + #pragma omp target enter data map(s1.ptr1[:n1], s1.ptr2[:n2], a[:n3], b[:n4]) + + #pragma omp target map(s2.ptr1[:n1], s2.ptr2[:n2], c[:n3], d[:n4]) + { + if (s2.ptr1 != (void *) 0L || s2.ptr2 != (void *) 0xdeedbeef + || c != (void *) 0L || d != (void *) 0xdeedbeef) + __builtin_abort (); + } + + #pragma omp target map(s1.ptr1[:n1], s1.ptr2[:n2], a[:n3], b[:n4]) + { + if (s1.ptr1 != (void *) 0L || s1.ptr2 != (void *) 0xdeedbeef + || a != (void *) 0L || b != (void *) 0xdeedbeef) + __builtin_abort (); + } + + #pragma omp target + { + if (s1.ptr1 != (void *) 0L || s1.ptr2 != (void *) 0xdeedbeef + || a != (void *) 0L || b != (void *) 0xdeedbeef) + __builtin_abort (); + } + + #pragma omp target exit data map(s1.ptr1[:n1], s1.ptr2[:n2], a[:n3], b[:n4]) + + with_offset (); +} |