diff options
Diffstat (limited to 'libgomp/testsuite/libgomp.c')
58 files changed, 1834 insertions, 16 deletions
diff --git a/libgomp/testsuite/libgomp.c/alloc-managed-1.c b/libgomp/testsuite/libgomp.c/alloc-managed-1.c new file mode 100644 index 0000000..88ddcf3 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/alloc-managed-1.c @@ -0,0 +1,29 @@ +/* { dg-do run } */ +/* { dg-require-effective-target omp_managedmem } */ +/* { dg-additional-options "-foffload-options=amdgcn-amdhsa=-mxnack=on" { target offload_target_amdgcn_with_xnack } } */ + +/* Check that omp_alloc can allocate Managed Memory, and that host and target + can see the data, at the same address, without a mapping. */ + +#include <omp.h> +#include <stdint.h> + +int +main () +{ + int *a = (int *) omp_alloc(sizeof(int), ompx_gnu_managed_mem_alloc); + if (!a) + __builtin_abort (); + + *a = 42; + uintptr_t a_p = (uintptr_t)a; + + #pragma omp target is_device_ptr(a) + { + if (*a != 42 || a_p != (uintptr_t)a) + __builtin_abort (); + } + + omp_free(a, ompx_gnu_managed_mem_alloc); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/alloc-managed-2.c b/libgomp/testsuite/libgomp.c/alloc-managed-2.c new file mode 100644 index 0000000..660f6e6 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/alloc-managed-2.c @@ -0,0 +1,39 @@ +/* { dg-do run } */ +/* { dg-require-effective-target omp_managedmem } */ +/* { dg-additional-options -foffload-options=amdgcn-amdhsa=-mxnack=on { target offload_target_amdgcn_with_xnack } } */ + +/* Check that omp_calloc can allocate Managed Memory, and that host and target + can see the data, at the same address, without a mapping. */ + +#include <omp.h> +#include <stdint.h> + +int +main () +{ + int *a = (int *) omp_calloc(5, sizeof(int), ompx_gnu_managed_mem_alloc); + if (!a) + __builtin_abort (); + + /* Check that memory is zero-initialized */ + for (int i = 0; i < 5; i++) + if (a[i] != 0) + __builtin_abort (); + + a[0] = 42; + a[4] = 99; + uintptr_t a_p = (uintptr_t)a; + + #pragma omp target is_device_ptr(a) + { + if (a[0] != 42 || a[4] != 99 || a_p != (uintptr_t)a) + __builtin_abort (); + /* Check zero-initialization on device side */ + for (int i = 1; i < 4; i++) + if (a[i] != 0) + __builtin_abort (); + } + + omp_free(a, ompx_gnu_managed_mem_alloc); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/alloc-managed-3.c b/libgomp/testsuite/libgomp.c/alloc-managed-3.c new file mode 100644 index 0000000..fefdeb3 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/alloc-managed-3.c @@ -0,0 +1,45 @@ +/* { dg-do run } */ +/* { dg-require-effective-target omp_managedmem } */ +/* { dg-additional-options -foffload-options=amdgcn-amdhsa=-mxnack=on { target offload_target_amdgcn_with_xnack } } */ + +/* Check that omp_realloc can allocate Managed Memory, and that host and target + can see the data, at the same address, without a mapping. */ + +#include <omp.h> +#include <stdint.h> + +int +main () +{ + int *a = (int *) omp_alloc(2 * sizeof(int), ompx_gnu_managed_mem_alloc); + if (!a) + __builtin_abort (); + + a[0] = 42; + a[1] = 43; + + /* Reallocate to larger size */ + int *b = (int *) omp_realloc(a, 5 * sizeof(int), ompx_gnu_managed_mem_alloc, + ompx_gnu_managed_mem_alloc); + if (!b) + __builtin_abort (); + + /* Check that original data is preserved */ + if (b[0] != 42 || b[1] != 43) + __builtin_abort (); + + b[2] = 44; + b[3] = 45; + b[4] = 46; + uintptr_t b_p = (uintptr_t)b; + + #pragma omp target is_device_ptr(b) + { + if (b[0] != 42 || b[1] != 43 || b[2] != 44 || b[3] != 45 || b[4] != 46 + || b_p != (uintptr_t)b) + __builtin_abort (); + } + + omp_free(b, ompx_gnu_managed_mem_alloc); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/alloc-managed-4.c b/libgomp/testsuite/libgomp.c/alloc-managed-4.c new file mode 100644 index 0000000..577e3e2 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/alloc-managed-4.c @@ -0,0 +1,23 @@ +/* { dg-do run } */ +/* { dg-require-effective-target omp_managedmem } */ +/* { dg-additional-options -foffload-options=amdgcn-amdhsa=-mxnack=on { target offload_target_amdgcn_with_xnack } } */ +/* { dg-shouldfail "" } */ +/* { dg-output "libgomp: attempted to free managed memory at 0x\[0-9a-f\]+, but the default device is set to the host device" } */ + +/* Check that omp_free emits an error if the default device has been changed + to the host device. */ + +#include <omp.h> +#include <stdint.h> + +int +main () +{ + int *a = (int *) omp_alloc(2 * sizeof(int), ompx_gnu_managed_mem_alloc); + if (!a) + __builtin_abort (); + + omp_set_default_device (omp_initial_device); + omp_free(a, ompx_gnu_managed_mem_alloc); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/alloc-pinned-1.c b/libgomp/testsuite/libgomp.c/alloc-pinned-1.c index 672f245..693f903 100644 --- a/libgomp/testsuite/libgomp.c/alloc-pinned-1.c +++ b/libgomp/testsuite/libgomp.c/alloc-pinned-1.c @@ -2,6 +2,8 @@ /* { dg-skip-if "Pinning not implemented on this host" { ! *-*-linux-gnu* } } */ +/* { dg-additional-options -DOFFLOAD_DEVICE_NVPTX { target offload_device_nvptx } } */ + /* Test that pinned memory works. */ #include <stdio.h> @@ -63,10 +65,16 @@ verify0 (char *p, size_t s) int main () { +#ifdef OFFLOAD_DEVICE_NVPTX + /* Go big or go home. + The OS ulimit does not affect memory locked via CUDA for NVPTX devices. */ + const int SIZE = 40 * 1024 * 1024; +#else /* Allocate at least a page each time, allowing space for overhead, but stay within the ulimit. */ const int SIZE = PAGE_SIZE - 128; CHECK_SIZE (SIZE * 5); // This is intended to help diagnose failures +#endif const omp_alloctrait_t traits[] = { { omp_atk_pinned, 1 } @@ -88,21 +96,39 @@ main () abort (); int amount = get_pinned_mem (); +#ifdef OFFLOAD_DEVICE_NVPTX + /* This doesn't show up as process 'VmLck'ed memory. */ + if (amount != 0) + abort (); +#else if (amount == 0) abort (); +#endif p = omp_realloc (p, SIZE * 2, allocator, allocator); int amount2 = get_pinned_mem (); +#ifdef OFFLOAD_DEVICE_NVPTX + /* This doesn't show up as process 'VmLck'ed memory. */ + if (amount2 != 0) + abort (); +#else if (amount2 <= amount) abort (); +#endif /* SIZE*2 ensures that it doesn't slot into the space possibly vacated by realloc. */ p = omp_calloc (1, SIZE * 2, allocator); +#ifdef OFFLOAD_DEVICE_NVPTX + /* This doesn't show up as process 'VmLck'ed memory. */ + if (get_pinned_mem () != 0) + abort (); +#else if (get_pinned_mem () <= amount2) abort (); +#endif verify0 (p, SIZE * 2); diff --git a/libgomp/testsuite/libgomp.c/alloc-pinned-2.c b/libgomp/testsuite/libgomp.c/alloc-pinned-2.c index b6d1d83..e7ac64e 100644 --- a/libgomp/testsuite/libgomp.c/alloc-pinned-2.c +++ b/libgomp/testsuite/libgomp.c/alloc-pinned-2.c @@ -2,6 +2,8 @@ /* { dg-skip-if "Pinning not implemented on this host" { ! *-*-linux-gnu* } } */ +/* { dg-additional-options -DOFFLOAD_DEVICE_NVPTX { target offload_device_nvptx } } */ + /* Test that pinned memory works (pool_size code path). */ #include <stdio.h> @@ -63,10 +65,16 @@ verify0 (char *p, size_t s) int main () { +#ifdef OFFLOAD_DEVICE_NVPTX + /* Go big or go home. + The OS ulimit does not affect memory locked via CUDA for NVPTX devices. */ + const int SIZE = 40 * 1024 * 1024; +#else /* Allocate at least a page each time, allowing space for overhead, but stay within the ulimit. */ const int SIZE = PAGE_SIZE - 128; CHECK_SIZE (SIZE * 5); // This is intended to help diagnose failures +#endif const omp_alloctrait_t traits[] = { { omp_atk_pinned, 1 }, @@ -89,16 +97,28 @@ main () abort (); int amount = get_pinned_mem (); +#ifdef OFFLOAD_DEVICE_NVPTX + /* This doesn't show up as process 'VmLck'ed memory. */ + if (amount != 0) + abort (); +#else if (amount == 0) abort (); +#endif p = omp_realloc (p, SIZE * 2, allocator, allocator); if (!p) abort (); int amount2 = get_pinned_mem (); +#ifdef OFFLOAD_DEVICE_NVPTX + /* This doesn't show up as process 'VmLck'ed memory. */ + if (amount2 != 0) + abort (); +#else if (amount2 <= amount) abort (); +#endif /* SIZE*2 ensures that it doesn't slot into the space possibly vacated by realloc. */ @@ -106,8 +126,14 @@ main () if (!p) abort (); +#ifdef OFFLOAD_DEVICE_NVPTX + /* This doesn't show up as process 'VmLck'ed memory. */ + if (get_pinned_mem () != 0) + abort (); +#else if (get_pinned_mem () <= amount2) abort (); +#endif verify0 (p, SIZE * 2); diff --git a/libgomp/testsuite/libgomp.c/alloc-pinned-3.c b/libgomp/testsuite/libgomp.c/alloc-pinned-3.c index 11dc818..250cb55 100644 --- a/libgomp/testsuite/libgomp.c/alloc-pinned-3.c +++ b/libgomp/testsuite/libgomp.c/alloc-pinned-3.c @@ -1,5 +1,7 @@ /* { dg-do run } */ +/* { dg-additional-options -DOFFLOAD_DEVICE_NVPTX { target offload_device_nvptx } } */ + /* Test that pinned memory fails correctly. */ #include <stdio.h> @@ -75,8 +77,15 @@ verify0 (char *p, size_t s) int main () { +#ifdef OFFLOAD_DEVICE_NVPTX + /* Go big or go home. + The OS ulimit does not affect memory locked via CUDA for NVPTX devices. */ + const int SIZE = 40 * 1024 * 1024; +#else /* This needs to be large enough to cover multiple pages. */ const int SIZE = PAGE_SIZE * 4; +#endif + const int PIN_LIMIT = PAGE_SIZE * 2; /* Pinned memory, no fallback. */ const omp_alloctrait_t traits1[] = { @@ -101,23 +110,34 @@ main () #endif /* Ensure that the limit is smaller than the allocation. */ - set_pin_limit (SIZE / 2); + set_pin_limit (PIN_LIMIT); // Sanity check if (get_pinned_mem () != 0) abort (); - // Should fail void *p1 = omp_alloc (SIZE, allocator1); +#ifdef OFFLOAD_DEVICE_NVPTX + // Doesn't care about 'set_pin_limit'. + if (!p1) + abort (); +#else + // Should fail if (p1) abort (); +#endif - // Should fail void *p2 = omp_calloc (1, SIZE, allocator1); +#ifdef OFFLOAD_DEVICE_NVPTX + // Doesn't care about 'set_pin_limit'. + if (!p2) + abort (); +#else + // Should fail if (p2) abort (); +#endif - // Should fall back void *p3 = omp_alloc (SIZE, allocator2); if (!p3) abort (); @@ -128,16 +148,29 @@ main () abort (); verify0 (p4, SIZE); - // Should fail to realloc void *notpinned = omp_alloc (SIZE, omp_default_mem_alloc); void *p5 = omp_realloc (notpinned, SIZE, allocator1, omp_default_mem_alloc); +#ifdef OFFLOAD_DEVICE_NVPTX + // Doesn't care about 'set_pin_limit'; does reallocate. + if (!notpinned || !p5 || p5 == notpinned) + abort (); +#else + // Should fail to realloc if (!notpinned || p5) abort (); +#endif - // Should fall back to no realloc needed +#ifdef OFFLOAD_DEVICE_NVPTX + void *p6 = omp_realloc (p5, SIZE, allocator2, allocator1); + // Does reallocate. + if (p5 == p6) + abort (); +#else void *p6 = omp_realloc (notpinned, SIZE, allocator2, omp_default_mem_alloc); + // Should fall back to no realloc needed if (p6 != notpinned) abort (); +#endif // No memory should have been pinned int amount = get_pinned_mem (); diff --git a/libgomp/testsuite/libgomp.c/alloc-pinned-4.c b/libgomp/testsuite/libgomp.c/alloc-pinned-4.c index 2ecd01f..b7a9966 100644 --- a/libgomp/testsuite/libgomp.c/alloc-pinned-4.c +++ b/libgomp/testsuite/libgomp.c/alloc-pinned-4.c @@ -1,5 +1,7 @@ /* { dg-do run } */ +/* { dg-additional-options -DOFFLOAD_DEVICE_NVPTX { target offload_device_nvptx } } */ + /* Test that pinned memory fails correctly, pool_size code path. */ #include <stdio.h> @@ -75,8 +77,15 @@ verify0 (char *p, size_t s) int main () { +#ifdef OFFLOAD_DEVICE_NVPTX + /* Go big or go home. + The OS ulimit does not affect memory locked via CUDA for NVPTX devices. */ + const int SIZE = 40 * 1024 * 1024; +#else /* This needs to be large enough to cover multiple pages. */ const int SIZE = PAGE_SIZE * 4; +#endif + const int PIN_LIMIT = PAGE_SIZE * 2; /* Pinned memory, no fallback. */ const omp_alloctrait_t traits1[] = { @@ -103,21 +112,33 @@ main () #endif /* Ensure that the limit is smaller than the allocation. */ - set_pin_limit (SIZE / 2); + set_pin_limit (PIN_LIMIT); // Sanity check if (get_pinned_mem () != 0) abort (); - // Should fail void *p = omp_alloc (SIZE, allocator1); +#ifdef OFFLOAD_DEVICE_NVPTX + // Doesn't care about 'set_pin_limit'. + if (!p) + abort (); +#else + // Should fail if (p) abort (); +#endif - // Should fail p = omp_calloc (1, SIZE, allocator1); +#ifdef OFFLOAD_DEVICE_NVPTX + // Doesn't care about 'set_pin_limit'. + if (!p) + abort (); +#else + // Should fail if (p) abort (); +#endif // Should fall back p = omp_alloc (SIZE, allocator2); @@ -130,16 +151,29 @@ main () abort (); verify0 (p, SIZE); - // Should fail to realloc void *notpinned = omp_alloc (SIZE, omp_default_mem_alloc); p = omp_realloc (notpinned, SIZE, allocator1, omp_default_mem_alloc); +#ifdef OFFLOAD_DEVICE_NVPTX + // Doesn't care about 'set_pin_limit'; does reallocate. + if (!notpinned || !p || p == notpinned) + abort (); +#else + // Should fail to realloc if (!notpinned || p) abort (); +#endif - // Should fall back to no realloc needed +#ifdef OFFLOAD_DEVICE_NVPTX + void *p_ = omp_realloc (p, SIZE, allocator2, allocator1); + // Does reallocate. + if (p_ == p) + abort (); +#else p = omp_realloc (notpinned, SIZE, allocator2, omp_default_mem_alloc); + // Should fall back to no realloc needed if (p != notpinned) abort (); +#endif // No memory should have been pinned int amount = get_pinned_mem (); diff --git a/libgomp/testsuite/libgomp.c/alloc-pinned-5.c b/libgomp/testsuite/libgomp.c/alloc-pinned-5.c index 0ba2feb..cc77764 100644 --- a/libgomp/testsuite/libgomp.c/alloc-pinned-5.c +++ b/libgomp/testsuite/libgomp.c/alloc-pinned-5.c @@ -2,6 +2,8 @@ /* { dg-skip-if "Pinning not implemented on this host" { ! *-*-linux-gnu* } } */ +/* { dg-additional-options -DOFFLOAD_DEVICE_NVPTX { target offload_device_nvptx } } */ + /* Test that ompx_gnu_pinned_mem_alloc works. */ #include <stdio.h> @@ -63,10 +65,16 @@ verify0 (char *p, size_t s) int main () { +#ifdef OFFLOAD_DEVICE_NVPTX + /* Go big or go home. + The OS ulimit does not affect memory locked via CUDA for NVPTX devices. */ + const int SIZE = 40 * 1024 * 1024; +#else /* Allocate at least a page each time, allowing space for overhead, but stay within the ulimit. */ const int SIZE = PAGE_SIZE - 128; CHECK_SIZE (SIZE * 5); +#endif // Sanity check if (get_pinned_mem () != 0) @@ -77,22 +85,40 @@ main () abort (); int amount = get_pinned_mem (); +#ifdef OFFLOAD_DEVICE_NVPTX + /* This doesn't show up as process 'VmLck'ed memory. */ + if (amount != 0) + abort (); +#else if (amount == 0) abort (); +#endif p = omp_realloc (p, SIZE * 2, ompx_gnu_pinned_mem_alloc, ompx_gnu_pinned_mem_alloc); int amount2 = get_pinned_mem (); +#ifdef OFFLOAD_DEVICE_NVPTX + /* This doesn't show up as process 'VmLck'ed memory. */ + if (amount2 != 0) + abort (); +#else if (amount2 <= amount) abort (); +#endif /* SIZE*2 ensures that it doesn't slot into the space possibly vacated by realloc. */ p = omp_calloc (1, SIZE * 2, ompx_gnu_pinned_mem_alloc); +#ifdef OFFLOAD_DEVICE_NVPTX + /* This doesn't show up as process 'VmLck'ed memory. */ + if (get_pinned_mem () != 0) + abort (); +#else if (get_pinned_mem () <= amount2) abort (); +#endif verify0 (p, SIZE * 2); diff --git a/libgomp/testsuite/libgomp.c/alloc-pinned-6.c b/libgomp/testsuite/libgomp.c/alloc-pinned-6.c index 99f1269..6dd5544 100644 --- a/libgomp/testsuite/libgomp.c/alloc-pinned-6.c +++ b/libgomp/testsuite/libgomp.c/alloc-pinned-6.c @@ -1,4 +1,5 @@ /* { dg-do run } */ +/* { dg-additional-options -DOFFLOAD_DEVICE_NVPTX { target offload_device_nvptx } } */ /* Test that ompx_gnu_pinned_mem_alloc fails correctly. */ @@ -66,32 +67,57 @@ set_pin_limit (int size) int main () { +#ifdef OFFLOAD_DEVICE_NVPTX + /* Go big or go home. + The OS ulimit does not affect memory locked via CUDA for NVPTX devices. */ + const int SIZE = 40 * 1024 * 1024; +#else /* Allocate at least a page each time, but stay within the ulimit. */ const int SIZE = PAGE_SIZE * 4; +#endif + const int PIN_LIMIT = PAGE_SIZE*2; /* Ensure that the limit is smaller than the allocation. */ - set_pin_limit (SIZE / 2); + set_pin_limit (PIN_LIMIT); // Sanity check if (get_pinned_mem () != 0) abort (); - // Should fail void *p = omp_alloc (SIZE, ompx_gnu_pinned_mem_alloc); +#ifdef OFFLOAD_DEVICE_NVPTX + // Doesn't care about 'set_pin_limit'. + if (!p) + abort (); +#else + // Should fail if (p) abort (); +#endif - // Should fail p = omp_calloc (1, SIZE, ompx_gnu_pinned_mem_alloc); +#ifdef OFFLOAD_DEVICE_NVPTX + // Doesn't care about 'set_pin_limit'. + if (!p) + abort (); +#else + // Should fail if (p) abort (); +#endif - // Should fail to realloc void *notpinned = omp_alloc (SIZE, omp_default_mem_alloc); p = omp_realloc (notpinned, SIZE, ompx_gnu_pinned_mem_alloc, omp_default_mem_alloc); +#ifdef OFFLOAD_DEVICE_NVPTX + // Doesn't care about 'set_pin_limit'; does reallocate. + if (!notpinned || !p || p == notpinned) + abort (); +#else + // Should fail to realloc if (!notpinned || p) abort (); +#endif // No memory should have been pinned int amount = get_pinned_mem (); diff --git a/libgomp/testsuite/libgomp.c/alloc-pinned-8.c b/libgomp/testsuite/libgomp.c/alloc-pinned-8.c new file mode 100644 index 0000000..0fc737b --- /dev/null +++ b/libgomp/testsuite/libgomp.c/alloc-pinned-8.c @@ -0,0 +1,122 @@ +/* { dg-do run } */ + +/* { dg-skip-if "Pinning not implemented on this host" { ! *-*-linux-gnu* } } */ + +/* { dg-additional-options -DOFFLOAD_DEVICE_NVPTX { target offload_device_nvptx } } */ + +/* Test that pinned memory works for small allocations. */ + +#include <stdio.h> +#include <stdlib.h> + +#ifdef __linux__ +#include <sys/types.h> +#include <unistd.h> + +#include <sys/mman.h> +#include <sys/resource.h> + +#define PAGE_SIZE sysconf(_SC_PAGESIZE) +#define CHECK_SIZE(SIZE) { \ + struct rlimit limit; \ + if (getrlimit (RLIMIT_MEMLOCK, &limit) \ + || limit.rlim_cur <= SIZE) \ + fprintf (stderr, "insufficient lockable memory; please increase ulimit\n"); \ + } + +int +get_pinned_mem () +{ + int pid = getpid (); + char buf[100]; + sprintf (buf, "/proc/%d/status", pid); + + FILE *proc = fopen (buf, "r"); + if (!proc) + abort (); + while (fgets (buf, 100, proc)) + { + int val; + if (sscanf (buf, "VmLck: %d", &val)) + { + fclose (proc); + return val; + } + } + abort (); +} +#else +#error "OS unsupported" +#endif + +static void +verify0 (char *p, size_t s) +{ + for (size_t i = 0; i < s; ++i) + if (p[i] != 0) + abort (); +} + +#include <omp.h> + +int +main () +{ + /* Choose a small size where all our allocations fit on one page. */ + const int SIZE = 10; +#ifndef OFFLOAD_DEVICE_NVPTX + CHECK_SIZE (SIZE*4); +#endif + + const omp_alloctrait_t traits[] = { + { omp_atk_pinned, 1 } + }; + omp_allocator_handle_t allocator = omp_init_allocator (omp_default_mem_space, 1, traits); + + // Sanity check + if (get_pinned_mem () != 0) + abort (); + + void *p = omp_alloc (SIZE, allocator); + if (!p) + abort (); + + int amount = get_pinned_mem (); +#ifdef OFFLOAD_DEVICE_NVPTX + /* This doesn't show up as process 'VmLck'ed memory. */ + if (amount != 0) + abort (); +#else + if (amount == 0) + abort (); +#endif + + p = omp_realloc (p, SIZE * 2, allocator, allocator); + + int amount2 = get_pinned_mem (); +#ifdef OFFLOAD_DEVICE_NVPTX + /* This doesn't show up as process 'VmLck'ed memory. */ + if (amount2 != 0) + abort (); +#else + /* A small allocation should not allocate another page. */ + if (amount2 != amount) + abort (); +#endif + + p = omp_calloc (1, SIZE, allocator); + +#ifdef OFFLOAD_DEVICE_NVPTX + /* This doesn't show up as process 'VmLck'ed memory. */ + if (get_pinned_mem () != 0) + abort (); +#else + /* A small allocation should not allocate another page. */ + if (get_pinned_mem () != amount2) + abort (); +#endif + + verify0 (p, SIZE); + + return 0; +} 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-gfx10-3-generic.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx10-3-generic.c new file mode 100644 index 0000000..b7b95e6 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx10-3-generic.c @@ -0,0 +1,25 @@ +/* { dg-do link { target { offload_target_amdgcn } } } */ +/* { dg-additional-options -foffload=amdgcn-amdhsa } */ +/* { dg-additional-options -foffload=-march=gfx10-3-generic } */ +/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */ + +#include "declare-variant-4.h" + +/* { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "= gfx10_3_generic \\(\\);" "optimized" } } */ + + +/* This code will link nicely if the multilib for that GPU architecture + has been build for GCC. In that case, scan-offload-tree-dump will + PASS and the linking will yield an XPASS message due to following line: */ + +/* { dg-excess-errors "ld: error: unable to find library -lgomp|gcn mkoffload: fatal error" } */ + +/* If the multi-lib config is not available (as this is a generic config), + scan-offload-tree-dump will PASS - but linking fails with the + following error (XFAIL): + ld: error: unable to find library -lgomp + collect2: error: ld returned 1 exit status + gcn mkoffload: fatal error: ...-gnu-accel-amdgcn-amdhsa-gcc returned 1 exit status + compilation terminated. + lto-wrapper: fatal error: .../amdgcn-amdhsa/mkoffload returned 1 exit status + compilation terminated. */ diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1030.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1030.c index d98d5ef..3703e96 100644 --- a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1030.c +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1030.c @@ -6,3 +6,28 @@ #include "declare-variant-4.h" /* { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "= gfx1030 \\(\\);" "optimized" } } */ + + +/* This code will link nicely if the multilib for that GPU architecture + has been build for GCC. In that case, scan-offload-tree-dump will + PASS and the linking will yield an XPASS message due to following line: */ + +/* { dg-excess-errors "ld: error: unable to find library -lgomp|gcn mkoffload: fatal error" } */ + +/* If the multi-lib config is not available, there are two options: + + * If the generic multi-lib is available, mkoffload fails early, + yielding UNRESOLVED for scan-offload-tree-dump and an XFAIL + for the message: + gcn mkoffload: fatal error: GCC was built without library support + for '-march=gfx...'; consider compiling for the associated + generic architecture '-march=gfx...-generic' instead + + * Or compling succeeds - then scan-offload-tree-dump will PASS - + but linking fails with the following error (XFAIL): + ld: error: unable to find library -lgomp + collect2: error: ld returned 1 exit status + gcn mkoffload: fatal error: ...-gnu-accel-amdgcn-amdhsa-gcc returned 1 exit status + compilation terminated. + lto-wrapper: fatal error: .../amdgcn-amdhsa/mkoffload returned 1 exit status + compilation terminated. */ diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1031.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1031.c new file mode 100644 index 0000000..e0d6289 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1031.c @@ -0,0 +1,33 @@ +/* { dg-do link { target { offload_target_amdgcn } } } */ +/* { dg-additional-options -foffload=amdgcn-amdhsa } */ +/* { dg-additional-options -foffload=-march=gfx1031 } */ +/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */ + +#include "declare-variant-4.h" + +/* { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "= gfx1031 \\(\\);" "optimized" } } */ + + +/* This code will link nicely if the multilib for that GPU architecture + has been build for GCC. In that case, scan-offload-tree-dump will + PASS and the linking will yield an XPASS message due to following line: */ + +/* { dg-excess-errors "ld: error: unable to find library -lgomp|gcn mkoffload: fatal error" } */ + +/* If the multi-lib config is not available, there are two options: + + * If the generic multi-lib is available, mkoffload fails early, + yielding UNRESOLVED for scan-offload-tree-dump and an XFAIL + for the message: + gcn mkoffload: fatal error: GCC was built without library support + for '-march=gfx...'; consider compiling for the associated + generic architecture '-march=gfx...-generic' instead + + * Or compling succeeds - then scan-offload-tree-dump will PASS - + but linking fails with the following error (XFAIL): + ld: error: unable to find library -lgomp + collect2: error: ld returned 1 exit status + gcn mkoffload: fatal error: ...-gnu-accel-amdgcn-amdhsa-gcc returned 1 exit status + compilation terminated. + lto-wrapper: fatal error: .../amdgcn-amdhsa/mkoffload returned 1 exit status + compilation terminated. */ diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1032.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1032.c new file mode 100644 index 0000000..46174cc2 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1032.c @@ -0,0 +1,33 @@ +/* { dg-do link { target { offload_target_amdgcn } } } */ +/* { dg-additional-options -foffload=amdgcn-amdhsa } */ +/* { dg-additional-options -foffload=-march=gfx1032 } */ +/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */ + +#include "declare-variant-4.h" + +/* { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "= gfx1032 \\(\\);" "optimized" } } */ + + +/* This code will link nicely if the multilib for that GPU architecture + has been build for GCC. In that case, scan-offload-tree-dump will + PASS and the linking will yield an XPASS message due to following line: */ + +/* { dg-excess-errors "ld: error: unable to find library -lgomp|gcn mkoffload: fatal error" } */ + +/* If the multi-lib config is not available, there are two options: + + * If the generic multi-lib is available, mkoffload fails early, + yielding UNRESOLVED for scan-offload-tree-dump and an XFAIL + for the message: + gcn mkoffload: fatal error: GCC was built without library support + for '-march=gfx...'; consider compiling for the associated + generic architecture '-march=gfx...-generic' instead + + * Or compling succeeds - then scan-offload-tree-dump will PASS - + but linking fails with the following error (XFAIL): + ld: error: unable to find library -lgomp + collect2: error: ld returned 1 exit status + gcn mkoffload: fatal error: ...-gnu-accel-amdgcn-amdhsa-gcc returned 1 exit status + compilation terminated. + lto-wrapper: fatal error: .../amdgcn-amdhsa/mkoffload returned 1 exit status + compilation terminated. */ diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1033.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1033.c new file mode 100644 index 0000000..1bd6e66 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1033.c @@ -0,0 +1,33 @@ +/* { dg-do link { target { offload_target_amdgcn } } } */ +/* { dg-additional-options -foffload=amdgcn-amdhsa } */ +/* { dg-additional-options -foffload=-march=gfx1033 } */ +/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */ + +#include "declare-variant-4.h" + +/* { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "= gfx1033 \\(\\);" "optimized" } } */ + + +/* This code will link nicely if the multilib for that GPU architecture + has been build for GCC. In that case, scan-offload-tree-dump will + PASS and the linking will yield an XPASS message due to following line: */ + +/* { dg-excess-errors "ld: error: unable to find library -lgomp|gcn mkoffload: fatal error" } */ + +/* If the multi-lib config is not available, there are two options: + + * If the generic multi-lib is available, mkoffload fails early, + yielding UNRESOLVED for scan-offload-tree-dump and an XFAIL + for the message: + gcn mkoffload: fatal error: GCC was built without library support + for '-march=gfx...'; consider compiling for the associated + generic architecture '-march=gfx...-generic' instead + + * Or compling succeeds - then scan-offload-tree-dump will PASS - + but linking fails with the following error (XFAIL): + ld: error: unable to find library -lgomp + collect2: error: ld returned 1 exit status + gcn mkoffload: fatal error: ...-gnu-accel-amdgcn-amdhsa-gcc returned 1 exit status + compilation terminated. + lto-wrapper: fatal error: .../amdgcn-amdhsa/mkoffload returned 1 exit status + compilation terminated. */ diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1034.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1034.c new file mode 100644 index 0000000..4f67a73 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1034.c @@ -0,0 +1,33 @@ +/* { dg-do link { target { offload_target_amdgcn } } } */ +/* { dg-additional-options -foffload=amdgcn-amdhsa } */ +/* { dg-additional-options -foffload=-march=gfx1034 } */ +/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */ + +#include "declare-variant-4.h" + +/* { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "= gfx1034 \\(\\);" "optimized" } } */ + + +/* This code will link nicely if the multilib for that GPU architecture + has been build for GCC. In that case, scan-offload-tree-dump will + PASS and the linking will yield an XPASS message due to following line: */ + +/* { dg-excess-errors "ld: error: unable to find library -lgomp|gcn mkoffload: fatal error" } */ + +/* If the multi-lib config is not available, there are two options: + + * If the generic multi-lib is available, mkoffload fails early, + yielding UNRESOLVED for scan-offload-tree-dump and an XFAIL + for the message: + gcn mkoffload: fatal error: GCC was built without library support + for '-march=gfx...'; consider compiling for the associated + generic architecture '-march=gfx...-generic' instead + + * Or compling succeeds - then scan-offload-tree-dump will PASS - + but linking fails with the following error (XFAIL): + ld: error: unable to find library -lgomp + collect2: error: ld returned 1 exit status + gcn mkoffload: fatal error: ...-gnu-accel-amdgcn-amdhsa-gcc returned 1 exit status + compilation terminated. + lto-wrapper: fatal error: .../amdgcn-amdhsa/mkoffload returned 1 exit status + compilation terminated. */ diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1035.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1035.c new file mode 100644 index 0000000..a69d5e7 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1035.c @@ -0,0 +1,33 @@ +/* { dg-do link { target { offload_target_amdgcn } } } */ +/* { dg-additional-options -foffload=amdgcn-amdhsa } */ +/* { dg-additional-options -foffload=-march=gfx1035 } */ +/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */ + +#include "declare-variant-4.h" + +/* { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "= gfx1035 \\(\\);" "optimized" } } */ + + +/* This code will link nicely if the multilib for that GPU architecture + has been build for GCC. In that case, scan-offload-tree-dump will + PASS and the linking will yield an XPASS message due to following line: */ + +/* { dg-excess-errors "ld: error: unable to find library -lgomp|gcn mkoffload: fatal error" } */ + +/* If the multi-lib config is not available, there are two options: + + * If the generic multi-lib is available, mkoffload fails early, + yielding UNRESOLVED for scan-offload-tree-dump and an XFAIL + for the message: + gcn mkoffload: fatal error: GCC was built without library support + for '-march=gfx...'; consider compiling for the associated + generic architecture '-march=gfx...-generic' instead + + * Or compling succeeds - then scan-offload-tree-dump will PASS - + but linking fails with the following error (XFAIL): + ld: error: unable to find library -lgomp + collect2: error: ld returned 1 exit status + gcn mkoffload: fatal error: ...-gnu-accel-amdgcn-amdhsa-gcc returned 1 exit status + compilation terminated. + lto-wrapper: fatal error: .../amdgcn-amdhsa/mkoffload returned 1 exit status + compilation terminated. */ diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1036.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1036.c index 93b8641..8c258c4 100644 --- a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1036.c +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1036.c @@ -6,3 +6,28 @@ #include "declare-variant-4.h" /* { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "= gfx1036 \\(\\);" "optimized" } } */ + + +/* This code will link nicely if the multilib for that GPU architecture + has been build for GCC. In that case, scan-offload-tree-dump will + PASS and the linking will yield an XPASS message due to following line: */ + +/* { dg-excess-errors "ld: error: unable to find library -lgomp|gcn mkoffload: fatal error" } */ + +/* If the multi-lib config is not available, there are two options: + + * If the generic multi-lib is available, mkoffload fails early, + yielding UNRESOLVED for scan-offload-tree-dump and an XFAIL + for the message: + gcn mkoffload: fatal error: GCC was built without library support + for '-march=gfx...'; consider compiling for the associated + generic architecture '-march=gfx...-generic' instead + + * Or compling succeeds - then scan-offload-tree-dump will PASS - + but linking fails with the following error (XFAIL): + ld: error: unable to find library -lgomp + collect2: error: ld returned 1 exit status + gcn mkoffload: fatal error: ...-gnu-accel-amdgcn-amdhsa-gcc returned 1 exit status + compilation terminated. + lto-wrapper: fatal error: .../amdgcn-amdhsa/mkoffload returned 1 exit status + compilation terminated. */ diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx11-generic.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx11-generic.c new file mode 100644 index 0000000..fa9efb4 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx11-generic.c @@ -0,0 +1,25 @@ +/* { dg-do link { target { offload_target_amdgcn } } } */ +/* { dg-additional-options -foffload=amdgcn-amdhsa } */ +/* { dg-additional-options -foffload=-march=gfx11-generic } */ +/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */ + +#include "declare-variant-4.h" + +/* { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "= gfx11_generic \\(\\);" "optimized" } } */ + + +/* This code will link nicely if the multilib for that GPU architecture + has been build for GCC. In that case, scan-offload-tree-dump will + PASS and the linking will yield an XPASS message due to following line: */ + +/* { dg-excess-errors "ld: error: unable to find library -lgomp|gcn mkoffload: fatal error" } */ + +/* If the multi-lib config is not available (as this is a generic config), + scan-offload-tree-dump will PASS - but linking fails with the + following error (XFAIL): + ld: error: unable to find library -lgomp + collect2: error: ld returned 1 exit status + gcn mkoffload: fatal error: ...-gnu-accel-amdgcn-amdhsa-gcc returned 1 exit status + compilation terminated. + lto-wrapper: fatal error: .../amdgcn-amdhsa/mkoffload returned 1 exit status + compilation terminated. */ diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1100.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1100.c index 6ade352..f0b7c6d 100644 --- a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1100.c +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1100.c @@ -6,3 +6,28 @@ #include "declare-variant-4.h" /* { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "= gfx1100 \\(\\);" "optimized" } } */ + + +/* This code will link nicely if the multilib for that GPU architecture + has been build for GCC. In that case, scan-offload-tree-dump will + PASS and the linking will yield an XPASS message due to following line: */ + +/* { dg-excess-errors "ld: error: unable to find library -lgomp|gcn mkoffload: fatal error" } */ + +/* If the multi-lib config is not available, there are two options: + + * If the generic multi-lib is available, mkoffload fails early, + yielding UNRESOLVED for scan-offload-tree-dump and an XFAIL + for the message: + gcn mkoffload: fatal error: GCC was built without library support + for '-march=gfx...'; consider compiling for the associated + generic architecture '-march=gfx...-generic' instead + + * Or compling succeeds - then scan-offload-tree-dump will PASS - + but linking fails with the following error (XFAIL): + ld: error: unable to find library -lgomp + collect2: error: ld returned 1 exit status + gcn mkoffload: fatal error: ...-gnu-accel-amdgcn-amdhsa-gcc returned 1 exit status + compilation terminated. + lto-wrapper: fatal error: .../amdgcn-amdhsa/mkoffload returned 1 exit status + compilation terminated. */ diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1101.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1101.c new file mode 100644 index 0000000..213e904 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1101.c @@ -0,0 +1,33 @@ +/* { dg-do link { target { offload_target_amdgcn } } } */ +/* { dg-additional-options -foffload=amdgcn-amdhsa } */ +/* { dg-additional-options -foffload=-march=gfx1101 } */ +/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */ + +#include "declare-variant-4.h" + +/* { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "= gfx1101 \\(\\);" "optimized" } } */ + + +/* This code will link nicely if the multilib for that GPU architecture + has been build for GCC. In that case, scan-offload-tree-dump will + PASS and the linking will yield an XPASS message due to following line: */ + +/* { dg-excess-errors "ld: error: unable to find library -lgomp|gcn mkoffload: fatal error" } */ + +/* If the multi-lib config is not available, there are two options: + + * If the generic multi-lib is available, mkoffload fails early, + yielding UNRESOLVED for scan-offload-tree-dump and an XFAIL + for the message: + gcn mkoffload: fatal error: GCC was built without library support + for '-march=gfx...'; consider compiling for the associated + generic architecture '-march=gfx...-generic' instead + + * Or compling succeeds - then scan-offload-tree-dump will PASS - + but linking fails with the following error (XFAIL): + ld: error: unable to find library -lgomp + collect2: error: ld returned 1 exit status + gcn mkoffload: fatal error: ...-gnu-accel-amdgcn-amdhsa-gcc returned 1 exit status + compilation terminated. + lto-wrapper: fatal error: .../amdgcn-amdhsa/mkoffload returned 1 exit status + compilation terminated. */ diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1102.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1102.c new file mode 100644 index 0000000..3f68dc8 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1102.c @@ -0,0 +1,33 @@ +/* { dg-do link { target { offload_target_amdgcn } } } */ +/* { dg-additional-options -foffload=amdgcn-amdhsa } */ +/* { dg-additional-options -foffload=-march=gfx1102 } */ +/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */ + +#include "declare-variant-4.h" + +/* { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "= gfx1102 \\(\\);" "optimized" } } */ + + +/* This code will link nicely if the multilib for that GPU architecture + has been build for GCC. In that case, scan-offload-tree-dump will + PASS and the linking will yield an XPASS message due to following line: */ + +/* { dg-excess-errors "ld: error: unable to find library -lgomp|gcn mkoffload: fatal error" } */ + +/* If the multi-lib config is not available, there are two options: + + * If the generic multi-lib is available, mkoffload fails early, + yielding UNRESOLVED for scan-offload-tree-dump and an XFAIL + for the message: + gcn mkoffload: fatal error: GCC was built without library support + for '-march=gfx...'; consider compiling for the associated + generic architecture '-march=gfx...-generic' instead + + * Or compling succeeds - then scan-offload-tree-dump will PASS - + but linking fails with the following error (XFAIL): + ld: error: unable to find library -lgomp + collect2: error: ld returned 1 exit status + gcn mkoffload: fatal error: ...-gnu-accel-amdgcn-amdhsa-gcc returned 1 exit status + compilation terminated. + lto-wrapper: fatal error: .../amdgcn-amdhsa/mkoffload returned 1 exit status + compilation terminated. */ diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1103.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1103.c index 6a6dc4f..c1eed44 100644 --- a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1103.c +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1103.c @@ -6,3 +6,28 @@ #include "declare-variant-4.h" /* { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "= gfx1103 \\(\\);" "optimized" } } */ + + +/* This code will link nicely if the multilib for that GPU architecture + has been build for GCC. In that case, scan-offload-tree-dump will + PASS and the linking will yield an XPASS message due to following line: */ + +/* { dg-excess-errors "ld: error: unable to find library -lgomp|gcn mkoffload: fatal error" } */ + +/* If the multi-lib config is not available, there are two options: + + * If the generic multi-lib is available, mkoffload fails early, + yielding UNRESOLVED for scan-offload-tree-dump and an XFAIL + for the message: + gcn mkoffload: fatal error: GCC was built without library support + for '-march=gfx...'; consider compiling for the associated + generic architecture '-march=gfx...-generic' instead + + * Or compling succeeds - then scan-offload-tree-dump will PASS - + but linking fails with the following error (XFAIL): + ld: error: unable to find library -lgomp + collect2: error: ld returned 1 exit status + gcn mkoffload: fatal error: ...-gnu-accel-amdgcn-amdhsa-gcc returned 1 exit status + compilation terminated. + lto-wrapper: fatal error: .../amdgcn-amdhsa/mkoffload returned 1 exit status + compilation terminated. */ diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1150.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1150.c new file mode 100644 index 0000000..39d64ca --- /dev/null +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1150.c @@ -0,0 +1,33 @@ +/* { dg-do link { target { offload_target_amdgcn } } } */ +/* { dg-additional-options -foffload=amdgcn-amdhsa } */ +/* { dg-additional-options -foffload=-march=gfx1150 } */ +/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */ + +#include "declare-variant-4.h" + +/* { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "= gfx1150 \\(\\);" "optimized" } } */ + + +/* This code will link nicely if the multilib for that GPU architecture + has been build for GCC. In that case, scan-offload-tree-dump will + PASS and the linking will yield an XPASS message due to following line: */ + +/* { dg-excess-errors "ld: error: unable to find library -lgomp|gcn mkoffload: fatal error" } */ + +/* If the multi-lib config is not available, there are two options: + + * If the generic multi-lib is available, mkoffload fails early, + yielding UNRESOLVED for scan-offload-tree-dump and an XFAIL + for the message: + gcn mkoffload: fatal error: GCC was built without library support + for '-march=gfx...'; consider compiling for the associated + generic architecture '-march=gfx...-generic' instead + + * Or compling succeeds - then scan-offload-tree-dump will PASS - + but linking fails with the following error (XFAIL): + ld: error: unable to find library -lgomp + collect2: error: ld returned 1 exit status + gcn mkoffload: fatal error: ...-gnu-accel-amdgcn-amdhsa-gcc returned 1 exit status + compilation terminated. + lto-wrapper: fatal error: .../amdgcn-amdhsa/mkoffload returned 1 exit status + compilation terminated. */ diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1151.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1151.c new file mode 100644 index 0000000..2a0c732 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1151.c @@ -0,0 +1,33 @@ +/* { dg-do link { target { offload_target_amdgcn } } } */ +/* { dg-additional-options -foffload=amdgcn-amdhsa } */ +/* { dg-additional-options -foffload=-march=gfx1151 } */ +/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */ + +#include "declare-variant-4.h" + +/* { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "= gfx1151 \\(\\);" "optimized" } } */ + + +/* This code will link nicely if the multilib for that GPU architecture + has been build for GCC. In that case, scan-offload-tree-dump will + PASS and the linking will yield an XPASS message due to following line: */ + +/* { dg-excess-errors "ld: error: unable to find library -lgomp|gcn mkoffload: fatal error" } */ + +/* If the multi-lib config is not available, there are two options: + + * If the generic multi-lib is available, mkoffload fails early, + yielding UNRESOLVED for scan-offload-tree-dump and an XFAIL + for the message: + gcn mkoffload: fatal error: GCC was built without library support + for '-march=gfx...'; consider compiling for the associated + generic architecture '-march=gfx...-generic' instead + + * Or compling succeeds - then scan-offload-tree-dump will PASS - + but linking fails with the following error (XFAIL): + ld: error: unable to find library -lgomp + collect2: error: ld returned 1 exit status + gcn mkoffload: fatal error: ...-gnu-accel-amdgcn-amdhsa-gcc returned 1 exit status + compilation terminated. + lto-wrapper: fatal error: .../amdgcn-amdhsa/mkoffload returned 1 exit status + compilation terminated. */ diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1152.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1152.c new file mode 100644 index 0000000..3c987dd --- /dev/null +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1152.c @@ -0,0 +1,33 @@ +/* { dg-do link { target { offload_target_amdgcn } } } */ +/* { dg-additional-options -foffload=amdgcn-amdhsa } */ +/* { dg-additional-options -foffload=-march=gfx1152 } */ +/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */ + +#include "declare-variant-4.h" + +/* { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "= gfx1152 \\(\\);" "optimized" } } */ + + +/* This code will link nicely if the multilib for that GPU architecture + has been build for GCC. In that case, scan-offload-tree-dump will + PASS and the linking will yield an XPASS message due to following line: */ + +/* { dg-excess-errors "ld: error: unable to find library -lgomp|gcn mkoffload: fatal error" } */ + +/* If the multi-lib config is not available, there are two options: + + * If the generic multi-lib is available, mkoffload fails early, + yielding UNRESOLVED for scan-offload-tree-dump and an XFAIL + for the message: + gcn mkoffload: fatal error: GCC was built without library support + for '-march=gfx...'; consider compiling for the associated + generic architecture '-march=gfx...-generic' instead + + * Or compling succeeds - then scan-offload-tree-dump will PASS - + but linking fails with the following error (XFAIL): + ld: error: unable to find library -lgomp + collect2: error: ld returned 1 exit status + gcn mkoffload: fatal error: ...-gnu-accel-amdgcn-amdhsa-gcc returned 1 exit status + compilation terminated. + lto-wrapper: fatal error: .../amdgcn-amdhsa/mkoffload returned 1 exit status + compilation terminated. */ diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1153.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1153.c new file mode 100644 index 0000000..7d38b82 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx1153.c @@ -0,0 +1,33 @@ +/* { dg-do link { target { offload_target_amdgcn } } } */ +/* { dg-additional-options -foffload=amdgcn-amdhsa } */ +/* { dg-additional-options -foffload=-march=gfx1153 } */ +/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */ + +#include "declare-variant-4.h" + +/* { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "= gfx1153 \\(\\);" "optimized" } } */ + + +/* This code will link nicely if the multilib for that GPU architecture + has been build for GCC. In that case, scan-offload-tree-dump will + PASS and the linking will yield an XPASS message due to following line: */ + +/* { dg-excess-errors "ld: error: unable to find library -lgomp|gcn mkoffload: fatal error" } */ + +/* If the multi-lib config is not available, there are two options: + + * If the generic multi-lib is available, mkoffload fails early, + yielding UNRESOLVED for scan-offload-tree-dump and an XFAIL + for the message: + gcn mkoffload: fatal error: GCC was built without library support + for '-march=gfx...'; consider compiling for the associated + generic architecture '-march=gfx...-generic' instead + + * Or compling succeeds - then scan-offload-tree-dump will PASS - + but linking fails with the following error (XFAIL): + ld: error: unable to find library -lgomp + collect2: error: ld returned 1 exit status + gcn mkoffload: fatal error: ...-gnu-accel-amdgcn-amdhsa-gcc returned 1 exit status + compilation terminated. + lto-wrapper: fatal error: .../amdgcn-amdhsa/mkoffload returned 1 exit status + compilation terminated. */ diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx9-4-generic.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx9-4-generic.c new file mode 100644 index 0000000..07d1254 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx9-4-generic.c @@ -0,0 +1,25 @@ +/* { dg-do link { target { offload_target_amdgcn } } } */ +/* { dg-additional-options -foffload=amdgcn-amdhsa } */ +/* { dg-additional-options -foffload=-march=gfx9-4-generic } */ +/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */ + +#include "declare-variant-4.h" + +/* { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "= gfx9_4_generic \\(\\);" "optimized" } } */ + + +/* This code will link nicely if the multilib for that GPU architecture + has been build for GCC. In that case, scan-offload-tree-dump will + PASS and the linking will yield an XPASS message due to following line: */ + +/* { dg-excess-errors "ld: error: unable to find library -lgomp|gcn mkoffload: fatal error" } */ + +/* If the multi-lib config is not available (as this is a generic config), + scan-offload-tree-dump will PASS - but linking fails with the + following error (XFAIL): + ld: error: unable to find library -lgomp + collect2: error: ld returned 1 exit status + gcn mkoffload: fatal error: ...-gnu-accel-amdgcn-amdhsa-gcc returned 1 exit status + compilation terminated. + lto-wrapper: fatal error: .../amdgcn-amdhsa/mkoffload returned 1 exit status + compilation terminated. */ diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx9-generic.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx9-generic.c new file mode 100644 index 0000000..d6ba097 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx9-generic.c @@ -0,0 +1,25 @@ +/* { dg-do link { target { offload_target_amdgcn } } } */ +/* { dg-additional-options -foffload=amdgcn-amdhsa } */ +/* { dg-additional-options -foffload=-march=gfx9-generic } */ +/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */ + +#include "declare-variant-4.h" + +/* { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "= gfx9_generic \\(\\);" "optimized" } } */ + + +/* This code will link nicely if the multilib for that GPU architecture + has been build for GCC. In that case, scan-offload-tree-dump will + PASS and the linking will yield an XPASS message due to following line: */ + +/* { dg-excess-errors "ld: error: unable to find library -lgomp|gcn mkoffload: fatal error" } */ + +/* If the multi-lib config is not available (as this is a generic config), + scan-offload-tree-dump will PASS - but linking fails with the + following error (XFAIL): + ld: error: unable to find library -lgomp + collect2: error: ld returned 1 exit status + gcn mkoffload: fatal error: ...-gnu-accel-amdgcn-amdhsa-gcc returned 1 exit status + compilation terminated. + lto-wrapper: fatal error: .../amdgcn-amdhsa/mkoffload returned 1 exit status + compilation terminated. */ diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c index f3f5244..37005fc 100644 --- a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx900.c @@ -6,3 +6,28 @@ #include "declare-variant-4.h" /* { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "= gfx900 \\(\\);" "optimized" } } */ + + +/* This code will link nicely if the multilib for that GPU architecture + has been build for GCC. In that case, scan-offload-tree-dump will + PASS and the linking will yield an XPASS message due to following line: */ + +/* { dg-excess-errors "ld: error: unable to find library -lgomp|gcn mkoffload: fatal error" } */ + +/* If the multi-lib config is not available, there are two options: + + * If the generic multi-lib is available, mkoffload fails early, + yielding UNRESOLVED for scan-offload-tree-dump and an XFAIL + for the message: + gcn mkoffload: fatal error: GCC was built without library support + for '-march=gfx...'; consider compiling for the associated + generic architecture '-march=gfx...-generic' instead + + * Or compling succeeds - then scan-offload-tree-dump will PASS - + but linking fails with the following error (XFAIL): + ld: error: unable to find library -lgomp + collect2: error: ld returned 1 exit status + gcn mkoffload: fatal error: ...-gnu-accel-amdgcn-amdhsa-gcc returned 1 exit status + compilation terminated. + lto-wrapper: fatal error: .../amdgcn-amdhsa/mkoffload returned 1 exit status + compilation terminated. */ diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx902.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx902.c new file mode 100644 index 0000000..82981c5 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx902.c @@ -0,0 +1,33 @@ +/* { dg-do link { target { offload_target_amdgcn } } } */ +/* { dg-additional-options -foffload=amdgcn-amdhsa } */ +/* { dg-additional-options -foffload=-march=gfx902 } */ +/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */ + +#include "declare-variant-4.h" + +/* { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "= gfx902 \\(\\);" "optimized" } } */ + + +/* This code will link nicely if the multilib for that GPU architecture + has been build for GCC. In that case, scan-offload-tree-dump will + PASS and the linking will yield an XPASS message due to following line: */ + +/* { dg-excess-errors "ld: error: unable to find library -lgomp|gcn mkoffload: fatal error" } */ + +/* If the multi-lib config is not available, there are two options: + + * If the generic multi-lib is available, mkoffload fails early, + yielding UNRESOLVED for scan-offload-tree-dump and an XFAIL + for the message: + gcn mkoffload: fatal error: GCC was built without library support + for '-march=gfx...'; consider compiling for the associated + generic architecture '-march=gfx...-generic' instead + + * Or compling succeeds - then scan-offload-tree-dump will PASS - + but linking fails with the following error (XFAIL): + ld: error: unable to find library -lgomp + collect2: error: ld returned 1 exit status + gcn mkoffload: fatal error: ...-gnu-accel-amdgcn-amdhsa-gcc returned 1 exit status + compilation terminated. + lto-wrapper: fatal error: .../amdgcn-amdhsa/mkoffload returned 1 exit status + compilation terminated. */ diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx904.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx904.c new file mode 100644 index 0000000..89815fe --- /dev/null +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx904.c @@ -0,0 +1,33 @@ +/* { dg-do link { target { offload_target_amdgcn } } } */ +/* { dg-additional-options -foffload=amdgcn-amdhsa } */ +/* { dg-additional-options -foffload=-march=gfx904 } */ +/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */ + +#include "declare-variant-4.h" + +/* { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "= gfx904 \\(\\);" "optimized" } } */ + + +/* This code will link nicely if the multilib for that GPU architecture + has been build for GCC. In that case, scan-offload-tree-dump will + PASS and the linking will yield an XPASS message due to following line: */ + +/* { dg-excess-errors "ld: error: unable to find library -lgomp|gcn mkoffload: fatal error" } */ + +/* If the multi-lib config is not available, there are two options: + + * If the generic multi-lib is available, mkoffload fails early, + yielding UNRESOLVED for scan-offload-tree-dump and an XFAIL + for the message: + gcn mkoffload: fatal error: GCC was built without library support + for '-march=gfx...'; consider compiling for the associated + generic architecture '-march=gfx...-generic' instead + + * Or compling succeeds - then scan-offload-tree-dump will PASS - + but linking fails with the following error (XFAIL): + ld: error: unable to find library -lgomp + collect2: error: ld returned 1 exit status + gcn mkoffload: fatal error: ...-gnu-accel-amdgcn-amdhsa-gcc returned 1 exit status + compilation terminated. + lto-wrapper: fatal error: .../amdgcn-amdhsa/mkoffload returned 1 exit status + compilation terminated. */ diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c index ac43388..aeef690 100644 --- a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx906.c @@ -6,3 +6,28 @@ #include "declare-variant-4.h" /* { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "= gfx906 \\(\\);" "optimized" } } */ + + +/* This code will link nicely if the multilib for that GPU architecture + has been build for GCC. In that case, scan-offload-tree-dump will + PASS and the linking will yield an XPASS message due to following line: */ + +/* { dg-excess-errors "ld: error: unable to find library -lgomp|gcn mkoffload: fatal error" } */ + +/* If the multi-lib config is not available, there are two options: + + * If the generic multi-lib is available, mkoffload fails early, + yielding UNRESOLVED for scan-offload-tree-dump and an XFAIL + for the message: + gcn mkoffload: fatal error: GCC was built without library support + for '-march=gfx...'; consider compiling for the associated + generic architecture '-march=gfx...-generic' instead + + * Or compling succeeds - then scan-offload-tree-dump will PASS - + but linking fails with the following error (XFAIL): + ld: error: unable to find library -lgomp + collect2: error: ld returned 1 exit status + gcn mkoffload: fatal error: ...-gnu-accel-amdgcn-amdhsa-gcc returned 1 exit status + compilation terminated. + lto-wrapper: fatal error: .../amdgcn-amdhsa/mkoffload returned 1 exit status + compilation terminated. */ diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c index f60741f..799b546 100644 --- a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx908.c @@ -6,3 +6,28 @@ #include "declare-variant-4.h" /* { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "= gfx908 \\(\\);" "optimized" } } */ + + +/* This code will link nicely if the multilib for that GPU architecture + has been build for GCC. In that case, scan-offload-tree-dump will + PASS and the linking will yield an XPASS message due to following line: */ + +/* { dg-excess-errors "ld: error: unable to find library -lgomp|gcn mkoffload: fatal error" } */ + +/* If the multi-lib config is not available, there are two options: + + * If the generic multi-lib is available, mkoffload fails early, + yielding UNRESOLVED for scan-offload-tree-dump and an XFAIL + for the message: + gcn mkoffload: fatal error: GCC was built without library support + for '-march=gfx...'; consider compiling for the associated + generic architecture '-march=gfx...-generic' instead + + * Or compling succeeds - then scan-offload-tree-dump will PASS - + but linking fails with the following error (XFAIL): + ld: error: unable to find library -lgomp + collect2: error: ld returned 1 exit status + gcn mkoffload: fatal error: ...-gnu-accel-amdgcn-amdhsa-gcc returned 1 exit status + compilation terminated. + lto-wrapper: fatal error: .../amdgcn-amdhsa/mkoffload returned 1 exit status + compilation terminated. */ diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx909.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx909.c new file mode 100644 index 0000000..e8a6f63 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx909.c @@ -0,0 +1,33 @@ +/* { dg-do link { target { offload_target_amdgcn } } } */ +/* { dg-additional-options -foffload=amdgcn-amdhsa } */ +/* { dg-additional-options -foffload=-march=gfx909 } */ +/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */ + +#include "declare-variant-4.h" + +/* { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "= gfx909 \\(\\);" "optimized" } } */ + + +/* This code will link nicely if the multilib for that GPU architecture + has been build for GCC. In that case, scan-offload-tree-dump will + PASS and the linking will yield an XPASS message due to following line: */ + +/* { dg-excess-errors "ld: error: unable to find library -lgomp|gcn mkoffload: fatal error" } */ + +/* If the multi-lib config is not available, there are two options: + + * If the generic multi-lib is available, mkoffload fails early, + yielding UNRESOLVED for scan-offload-tree-dump and an XFAIL + for the message: + gcn mkoffload: fatal error: GCC was built without library support + for '-march=gfx...'; consider compiling for the associated + generic architecture '-march=gfx...-generic' instead + + * Or compling succeeds - then scan-offload-tree-dump will PASS - + but linking fails with the following error (XFAIL): + ld: error: unable to find library -lgomp + collect2: error: ld returned 1 exit status + gcn mkoffload: fatal error: ...-gnu-accel-amdgcn-amdhsa-gcc returned 1 exit status + compilation terminated. + lto-wrapper: fatal error: .../amdgcn-amdhsa/mkoffload returned 1 exit status + compilation terminated. */ diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c index 832d174..de5626e 100644 --- a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx90a.c @@ -6,3 +6,28 @@ #include "declare-variant-4.h" /* { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "= gfx90a \\(\\);" "optimized" } } */ + + +/* This code will link nicely if the multilib for that GPU architecture + has been build for GCC. In that case, scan-offload-tree-dump will + PASS and the linking will yield an XPASS message due to following line: */ + +/* { dg-excess-errors "ld: error: unable to find library -lgomp|gcn mkoffload: fatal error" } */ + +/* If the multi-lib config is not available, there are two options: + + * If the generic multi-lib is available, mkoffload fails early, + yielding UNRESOLVED for scan-offload-tree-dump and an XFAIL + for the message: + gcn mkoffload: fatal error: GCC was built without library support + for '-march=gfx...'; consider compiling for the associated + generic architecture '-march=gfx...-generic' instead + + * Or compling succeeds - then scan-offload-tree-dump will PASS - + but linking fails with the following error (XFAIL): + ld: error: unable to find library -lgomp + collect2: error: ld returned 1 exit status + gcn mkoffload: fatal error: ...-gnu-accel-amdgcn-amdhsa-gcc returned 1 exit status + compilation terminated. + lto-wrapper: fatal error: .../amdgcn-amdhsa/mkoffload returned 1 exit status + compilation terminated. */ diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx90c.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx90c.c index 44629a8..dfad7ec 100644 --- a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx90c.c +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx90c.c @@ -6,3 +6,28 @@ #include "declare-variant-4.h" /* { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "= gfx90c \\(\\);" "optimized" } } */ + + +/* This code will link nicely if the multilib for that GPU architecture + has been build for GCC. In that case, scan-offload-tree-dump will + PASS and the linking will yield an XPASS message due to following line: */ + +/* { dg-excess-errors "ld: error: unable to find library -lgomp|gcn mkoffload: fatal error" } */ + +/* If the multi-lib config is not available, there are two options: + + * If the generic multi-lib is available, mkoffload fails early, + yielding UNRESOLVED for scan-offload-tree-dump and an XFAIL + for the message: + gcn mkoffload: fatal error: GCC was built without library support + for '-march=gfx...'; consider compiling for the associated + generic architecture '-march=gfx...-generic' instead + + * Or compling succeeds - then scan-offload-tree-dump will PASS - + but linking fails with the following error (XFAIL): + ld: error: unable to find library -lgomp + collect2: error: ld returned 1 exit status + gcn mkoffload: fatal error: ...-gnu-accel-amdgcn-amdhsa-gcc returned 1 exit status + compilation terminated. + lto-wrapper: fatal error: .../amdgcn-amdhsa/mkoffload returned 1 exit status + compilation terminated. */ 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..c8c7446 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx942.c @@ -0,0 +1,33 @@ +/* { 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" } } */ + + +/* This code will link nicely if the multilib for that GPU architecture + has been build for GCC. In that case, scan-offload-tree-dump will + PASS and the linking will yield an XPASS message due to following line: */ + +/* { dg-excess-errors "ld: error: unable to find library -lgomp|gcn mkoffload: fatal error" } */ + +/* If the multi-lib config is not available, there are two options: + + * If the generic multi-lib is available, mkoffload fails early, + yielding UNRESOLVED for scan-offload-tree-dump and an XFAIL + for the message: + gcn mkoffload: fatal error: GCC was built without library support + for '-march=gfx...'; consider compiling for the associated + generic architecture '-march=gfx...-generic' instead + + * Or compling succeeds - then scan-offload-tree-dump will PASS - + but linking fails with the following error (XFAIL): + ld: error: unable to find library -lgomp + collect2: error: ld returned 1 exit status + gcn mkoffload: fatal error: ...-gnu-accel-amdgcn-amdhsa-gcc returned 1 exit status + compilation terminated. + lto-wrapper: fatal error: .../amdgcn-amdhsa/mkoffload returned 1 exit status + compilation terminated. */ diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4-gfx950.c b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx950.c new file mode 100644 index 0000000..af81f11 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/declare-variant-4-gfx950.c @@ -0,0 +1,33 @@ +/* { dg-do link { target { offload_target_amdgcn } } } */ +/* { dg-additional-options -foffload=amdgcn-amdhsa } */ +/* { dg-additional-options -foffload=-march=gfx950 } */ +/* { dg-additional-options "-foffload=-fdump-tree-optimized" } */ + +#include "declare-variant-4.h" + +/* { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "= gfx950 \\(\\);" "optimized" } } */ + + +/* This code will link nicely if the multilib for that GPU architecture + has been build for GCC. In that case, scan-offload-tree-dump will + PASS and the linking will yield an XPASS message due to following line: */ + +/* { dg-excess-errors "ld: error: unable to find library -lgomp|gcn mkoffload: fatal error" } */ + +/* If the multi-lib config is not available, there are two options: + + * If the generic multi-lib is available, mkoffload fails early, + yielding UNRESOLVED for scan-offload-tree-dump and an XFAIL + for the message: + gcn mkoffload: fatal error: GCC was built without library support + for '-march=gfx...'; consider compiling for the associated + generic architecture '-march=gfx...-generic' instead + + * Or compling succeeds - then scan-offload-tree-dump will PASS - + but linking fails with the following error (XFAIL): + ld: error: unable to find library -lgomp + collect2: error: ld returned 1 exit status + gcn mkoffload: fatal error: ...-gnu-accel-amdgcn-amdhsa-gcc returned 1 exit status + compilation terminated. + lto-wrapper: fatal error: .../amdgcn-amdhsa/mkoffload returned 1 exit status + compilation terminated. */ diff --git a/libgomp/testsuite/libgomp.c/declare-variant-4.h b/libgomp/testsuite/libgomp.c/declare-variant-4.h index 53788d2..dd97edb 100644 --- a/libgomp/testsuite/libgomp.c/declare-variant-4.h +++ b/libgomp/testsuite/libgomp.c/declare-variant-4.h @@ -9,6 +9,20 @@ gfx900 (void) __attribute__ ((noipa)) int +gfx902 (void) +{ + return 0x902; +} + +__attribute__ ((noipa)) +int +gfx904 (void) +{ + return 0x904; +} + +__attribute__ ((noipa)) +int gfx906 (void) { return 0x906; @@ -23,6 +37,13 @@ gfx908 (void) __attribute__ ((noipa)) int +gfx909 (void) +{ + return 0x909; +} + +__attribute__ ((noipa)) +int gfx90a (void) { return 0x90a; @@ -37,6 +58,20 @@ gfx90c (void) __attribute__ ((noipa)) int +gfx942 (void) +{ + return 0x942; +} + +__attribute__ ((noipa)) +int +gfx950 (void) +{ + return 0x950; +} + +__attribute__ ((noipa)) +int gfx1030 (void) { return 0x1030; @@ -44,6 +79,41 @@ gfx1030 (void) __attribute__ ((noipa)) int +gfx1031 (void) +{ + return 0x1031; +} + +__attribute__ ((noipa)) +int +gfx1032 (void) +{ + return 0x1032; +} + +__attribute__ ((noipa)) +int +gfx1033 (void) +{ + return 0x1033; +} + +__attribute__ ((noipa)) +int +gfx1034 (void) +{ + return 0x1034; +} + +__attribute__ ((noipa)) +int +gfx1035 (void) +{ + return 0x1035; +} + +__attribute__ ((noipa)) +int gfx1036 (void) { return 0x1036; @@ -58,20 +128,111 @@ gfx1100 (void) __attribute__ ((noipa)) int +gfx1101 (void) +{ + return 0x1101; +} + +__attribute__ ((noipa)) +int +gfx1102 (void) +{ + return 0x1102; +} + +__attribute__ ((noipa)) +int gfx1103 (void) { return 0x1103; } +__attribute__ ((noipa)) +int +gfx1150 (void) +{ + return 0x1150; +} + +__attribute__ ((noipa)) +int +gfx1151 (void) +{ + return 0x1151; +} + +__attribute__ ((noipa)) +int +gfx1152 (void) +{ + return 0x1152; +} + +__attribute__ ((noipa)) +int +gfx1153 (void) +{ + return 0x1153; +} + +__attribute__ ((noipa)) +int +gfx9_generic (void) +{ + return 0x90ff; +} + +__attribute__ ((noipa)) +int +gfx9_4_generic (void) +{ + return 0x94ff; +} + +__attribute__ ((noipa)) +int +gfx10_3_generic (void) +{ + return 0x103ff; +} + +__attribute__ ((noipa)) +int +gfx11_generic (void) +{ + return 0x110ff; +} + + #pragma omp declare variant(gfx900) match(device = {isa("gfx900")}) +#pragma omp declare variant(gfx902) match(device = {isa("gfx902")}) +#pragma omp declare variant(gfx904) match(device = {isa("gfx904")}) #pragma omp declare variant(gfx906) match(device = {isa("gfx906")}) #pragma omp declare variant(gfx908) match(device = {isa("gfx908")}) +#pragma omp declare variant(gfx909) match(device = {isa("gfx909")}) #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(gfx950) match(device = {isa("gfx950")}) #pragma omp declare variant(gfx1030) match(device = {isa("gfx1030")}) +#pragma omp declare variant(gfx1031) match(device = {isa("gfx1031")}) +#pragma omp declare variant(gfx1032) match(device = {isa("gfx1032")}) +#pragma omp declare variant(gfx1033) match(device = {isa("gfx1033")}) +#pragma omp declare variant(gfx1034) match(device = {isa("gfx1034")}) +#pragma omp declare variant(gfx1035) match(device = {isa("gfx1035")}) #pragma omp declare variant(gfx1036) match(device = {isa("gfx1036")}) #pragma omp declare variant(gfx1100) match(device = {isa("gfx1100")}) +#pragma omp declare variant(gfx1101) match(device = {isa("gfx1101")}) +#pragma omp declare variant(gfx1102) match(device = {isa("gfx1102")}) #pragma omp declare variant(gfx1103) match(device = {isa("gfx1103")}) +#pragma omp declare variant(gfx1150) match(device = {isa("gfx1150")}) +#pragma omp declare variant(gfx1151) match(device = {isa("gfx1151")}) +#pragma omp declare variant(gfx1152) match(device = {isa("gfx1152")}) +#pragma omp declare variant(gfx1153) match(device = {isa("gfx1153")}) +#pragma omp declare variant(gfx9_generic) match(device = {isa("gfx9-generic")}) +#pragma omp declare variant(gfx9_4_generic) match(device = {isa("gfx9-4-generic")}) +#pragma omp declare variant(gfx10_3_generic) match(device = {isa("gfx10-3-generic")}) +#pragma omp declare variant(gfx11_generic) match(device = {isa("gfx11-generic")}) __attribute__ ((noipa)) int f (void) diff --git a/libgomp/testsuite/libgomp.c/device_uid.c b/libgomp/testsuite/libgomp.c/device_uid.c index 0412d06..83aba0f 100644 --- a/libgomp/testsuite/libgomp.c/device_uid.c +++ b/libgomp/testsuite/libgomp.c/device_uid.c @@ -5,10 +5,12 @@ int main() { const char **strs = (const char **) malloc (sizeof (char*) * (omp_get_num_devices () + 1)); - for (int i = omp_invalid_device - 1; i <= omp_get_num_devices () + 1; i++) + for (int i = omp_default_device - 1; i <= omp_get_num_devices () + 1; i++) { const char *str = omp_get_uid_from_device (i); int dev = omp_get_device_from_uid (str); + if (i == omp_default_device) + i = omp_get_default_device (); // __builtin_printf("%i -> %s -> %d\n", i, str, dev); if (i < omp_initial_device || i > omp_get_num_devices ()) { diff --git a/libgomp/testsuite/libgomp.c/interop-cuda-full.c b/libgomp/testsuite/libgomp.c/interop-cuda-full.c index 38aa6b1..c48a934 100644 --- a/libgomp/testsuite/libgomp.c/interop-cuda-full.c +++ b/libgomp/testsuite/libgomp.c/interop-cuda-full.c @@ -1,3 +1,6 @@ +/* { dg-do run { target { offload_device_nvptx } } } */ +/* { dg-do link { target { ! offload_device_nvptx } } } */ + /* { dg-require-effective-target openacc_cuda } */ /* { dg-require-effective-target openacc_cudart } */ /* { dg-additional-options "-lcuda -lcudart" } */ diff --git a/libgomp/testsuite/libgomp.c/interop-cuda-libonly.c b/libgomp/testsuite/libgomp.c/interop-cuda-libonly.c index 17cbb15..bc257a2 100644 --- a/libgomp/testsuite/libgomp.c/interop-cuda-libonly.c +++ b/libgomp/testsuite/libgomp.c/interop-cuda-libonly.c @@ -1,3 +1,6 @@ +/* { dg-do run { target { offload_device_nvptx } } } */ +/* { dg-do link { target { ! offload_device_nvptx } } } */ + /* { dg-require-effective-target openacc_libcudart } */ /* { dg-require-effective-target openacc_libcuda } */ /* { dg-additional-options "-lcuda -lcudart" } */ diff --git a/libgomp/testsuite/libgomp.c/interop-hip-amd-full.c b/libgomp/testsuite/libgomp.c/interop-hip-amd-full.c index d7725fc..bd44f44 100644 --- a/libgomp/testsuite/libgomp.c/interop-hip-amd-full.c +++ b/libgomp/testsuite/libgomp.c/interop-hip-amd-full.c @@ -1,3 +1,6 @@ +/* { dg-do run { target { offload_device_gcn } } } */ +/* { dg-do link { target { ! offload_device_gcn } } } */ + /* { dg-require-effective-target gomp_hip_header_amd } */ /* { dg-require-effective-target gomp_libamdhip64 } */ /* { dg-additional-options "-lamdhip64" } */ diff --git a/libgomp/testsuite/libgomp.c/interop-hip-amd-no-hip-header.c b/libgomp/testsuite/libgomp.c/interop-hip-amd-no-hip-header.c index 2584537..91ad987 100644 --- a/libgomp/testsuite/libgomp.c/interop-hip-amd-no-hip-header.c +++ b/libgomp/testsuite/libgomp.c/interop-hip-amd-no-hip-header.c @@ -1,3 +1,6 @@ +/* { dg-do run { target { offload_device_gcn } } } */ +/* { dg-do link { target { ! offload_device_gcn } } } */ + /* { dg-require-effective-target gomp_libamdhip64 } */ /* { dg-additional-options "-lamdhip64" } */ diff --git a/libgomp/testsuite/libgomp.c/interop-hip-nvidia-full.c b/libgomp/testsuite/libgomp.c/interop-hip-nvidia-full.c index 79af47d..d5dc236 100644 --- a/libgomp/testsuite/libgomp.c/interop-hip-nvidia-full.c +++ b/libgomp/testsuite/libgomp.c/interop-hip-nvidia-full.c @@ -1,3 +1,6 @@ +/* { dg-do run { target { offload_device_nvptx } } } */ +/* { dg-do link { target { ! offload_device_nvptx } } } */ + /* { dg-require-effective-target openacc_cudart } */ /* { dg-require-effective-target openacc_cuda } */ /* { dg-require-effective-target gomp_hip_header_nvidia } */ diff --git a/libgomp/testsuite/libgomp.c/interop-hip-nvidia-no-headers.c b/libgomp/testsuite/libgomp.c/interop-hip-nvidia-no-headers.c index 4586398..7cff2cb 100644 --- a/libgomp/testsuite/libgomp.c/interop-hip-nvidia-no-headers.c +++ b/libgomp/testsuite/libgomp.c/interop-hip-nvidia-no-headers.c @@ -1,3 +1,6 @@ +/* { dg-do run { target { offload_device_nvptx } } } */ +/* { dg-do link { target { ! offload_device_nvptx } } } */ + /* { dg-require-effective-target openacc_libcudart } */ /* { dg-require-effective-target openacc_libcuda } */ /* { dg-additional-options "-lcuda -lcudart" } */ diff --git a/libgomp/testsuite/libgomp.c/interop-hip-nvidia-no-hip-header.c b/libgomp/testsuite/libgomp.c/interop-hip-nvidia-no-hip-header.c index 4186984..7b7dc74 100644 --- a/libgomp/testsuite/libgomp.c/interop-hip-nvidia-no-hip-header.c +++ b/libgomp/testsuite/libgomp.c/interop-hip-nvidia-no-hip-header.c @@ -1,3 +1,6 @@ +/* { dg-do run { target { offload_device_nvptx } } } */ +/* { dg-do link { target { ! offload_device_nvptx } } } */ + /* { dg-require-effective-target openacc_cudart } */ /* { dg-require-effective-target openacc_cuda } */ /* { dg-additional-options "-lcuda -lcudart" } */ diff --git a/libgomp/testsuite/libgomp.c/ipcp-cb-spec1.c b/libgomp/testsuite/libgomp.c/ipcp-cb-spec1.c new file mode 100644 index 0000000..ff82f4c --- /dev/null +++ b/libgomp/testsuite/libgomp.c/ipcp-cb-spec1.c @@ -0,0 +1,18 @@ +/* Test that GOMP_task is special cased when cpyfn is NULL. */ + +/* { dg-do run } */ +/* { dg-options "-O3 -fopenmp -std=gnu99 -fdump-ipa-cp-details" } */ +/* { dg-require-effective-target fopenmp } */ + +void test(int c) { + for (int i = 0; i < c; i++) + if (!__builtin_constant_p(c)) + __builtin_abort(); +} +int main() { +#pragma omp task + test(7); + return 0; +} + +/* { dg-final { scan-ipa-dump "Creating a specialized node of main._omp_fn" "cp" } } */ diff --git a/libgomp/testsuite/libgomp.c/ipcp-cb-spec2.c b/libgomp/testsuite/libgomp.c/ipcp-cb-spec2.c new file mode 100644 index 0000000..30894d7 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/ipcp-cb-spec2.c @@ -0,0 +1,20 @@ +/* Check that GOMP_task doesn't produce callback edges when cpyfn is not + NULL. */ + +/* { dg-do run } */ +/* { dg-options "-O3 -fopenmp -std=gnu99 -fdump-ipa-cp-details" } */ +/* { dg-require-effective-target fopenmp } */ + +void test(int *a) { + for (int i = 0; i < 100; i++) { + a[i] = i; + } +} +int main() { + int a[100]; + __builtin_memset (a, 0, sizeof (a)); + #pragma omp task + test (a); +} + +/* { dg-final { scan-ipa-dump-not "Created callback edge" "cp" } } */ diff --git a/libgomp/testsuite/libgomp.c/ipcp-cb1.c b/libgomp/testsuite/libgomp.c/ipcp-cb1.c new file mode 100644 index 0000000..e390f04 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/ipcp-cb1.c @@ -0,0 +1,24 @@ +/* Test that we can propagate constants into outlined OpenMP kernels. + This tests the underlying callback attribute and its related edges. */ + +/* { dg-do run } */ +/* { dg-options "-O3 -fopenmp -std=gnu99 -fdump-ipa-cp-details" } */ +/* { dg-require-effective-target fopenmp } */ + +int a[100]; +void test(int c) { +#pragma omp parallel for + for (int i = 0; i < c; i++) { + if (!__builtin_constant_p(c)) { + __builtin_abort(); + } + a[i] = i; + } +} +int main() { + test(100); + return a[5] - 5; +} + +/* { dg-final { scan-ipa-dump "Creating a specialized node of test._omp_fn" "cp" } } */ +/* { dg-final { scan-ipa-dump "Aggregate replacements: 0\\\[0]=100\\(by_ref\\)" "cp" } } */ diff --git a/libgomp/testsuite/libgomp.c/pr122281.c b/libgomp/testsuite/libgomp.c/pr122281.c new file mode 100644 index 0000000..68fc3be --- /dev/null +++ b/libgomp/testsuite/libgomp.c/pr122281.c @@ -0,0 +1,43 @@ +/* { dg-do run } */ +/* { dg-additional-options "-O3" } */ + +/* PR libgomp/122281 */ +/* PR middle-end/105001 */ + +/* If SIMT is supported, the inner 'omp simd' is duplicated into + one SIMT and one SIMD variant. SIMT is currently only supported + with nvidia GPUs. (This only happens with -O1 or higher.) + + The duplication failed for the SIMD case as a tree was shared and + the initialization only happened in the SIMT branch, i.e. when + compiling for a SIMT-device, all non-SIMD (offload or host devices) + accesses failed (segfault) for the atomic update. */ + +#include <omp.h> + +int __attribute__((noinline, noclone)) +f(int *A, int n, int dev) { + int cnt = 0; + #pragma omp target map(cnt) map(to:A[0:n]) device(dev) + { + #pragma omp parallel for simd + for (int i = 0; i < n; i++) + if (A[i] != 0) + { + #pragma omp atomic + cnt++; + } + } + return cnt; +} + +int main() { + int n = 10; + int A[10] = {11,22,33,44,55,66,77,88,99,110}; + + /* Run over all devices, including the host; the host should be SIMD, + some non-host devices might be SIMT. */ + for (int dev = omp_initial_device; dev <= omp_get_num_devices(); dev++) + if (f (A, n, dev) != 10) + __builtin_abort(); +} 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 (); +} |