diff options
Diffstat (limited to 'libgomp/testsuite')
8 files changed, 350 insertions, 0 deletions
diff --git a/libgomp/testsuite/libgomp.c++/static-aggr-constructor-destructor-1.C b/libgomp/testsuite/libgomp.c++/static-aggr-constructor-destructor-1.C new file mode 100644 index 0000000..403a071 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/static-aggr-constructor-destructor-1.C @@ -0,0 +1,72 @@ +// { dg-do run } +// { dg-additional-options "-fdump-tree-gimple -fdump-tree-optimized" } +// { dg-additional-options -foffload-options=-fdump-tree-optimized { target { offload_device_nvptx || offload_target_amdgcn } } } + +// { dg-final { scan-tree-dump-times "omp_is_initial_device" 1 "gimple" } } +// { dg-final { scan-tree-dump-times "_GLOBAL__off_I_v1" 1 "gimple" } } +// { dg-final { scan-tree-dump-times "__omp_target_static_init_and_destruction" 2 "gimple" } } +// { dg-final { scan-tree-dump-times "__attribute__\\(\\(\[^\n\r]*omp declare target nohost" 2 "gimple" } } + +// { dg-final { scan-tree-dump-not "omp_is_initial_device" "optimized" } } +// { dg-final { scan-tree-dump-not "__omp_target_static_init_and_destruction" "optimized" } } +// FIXME: should be '-not' not '-times' 1: +// { dg-final { scan-tree-dump-times "void _GLOBAL__off_I_v1" 1 "optimized" } } +// { dg-final { scan-tree-dump-times "__attribute__\\(\\(\[^\n\r]*omp declare target nohost" 1 "optimized" } } + +// { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump-not "omp_initial_device;" "optimized" { target offload_target_amdgcn } } } +// { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "v1\\._x = 5;" "optimized" { target offload_target_amdgcn } } } +// { dg-final { only_for_offload_target nvptx-none scan-offload-tree-dump-not "omp_initial_device;" "optimized" { target offload_target_nvptx } } } +// { dg-final { only_for_offload_target nvptx-none scan-offload-tree-dump "v1\\._x = 5;" "optimized" { target offload_target_nvptx } } } + +#include <cassert> +#include <omp.h> + +#pragma omp declare target + +struct str { + str(int x) : _x(x) { } + int add(str o) { return _x + o._x; } + int _x; +} v1(5); + +#pragma omp end declare target + +void check_host() +{ + assert (v1._x == 5); +} + +void check_devs() +{ + for (int dev = omp_initial_device; dev <= omp_get_num_devices(); dev++) + { + int res = 99, dev_num = 98; + #pragma omp target map(from: res, dev_num) device(dev) + { + res = v1._x; + dev_num = omp_get_device_num(); + } + assert (res == 5); + if (dev == omp_initial_device) + assert (dev_num == omp_get_num_devices()); + else + assert (dev_num == dev); + } +} + +int main() +{ + int res = -1; + str v2(2); + +#pragma omp target map(from:res) + { + res = v1.add(v2); + } + + assert (res == 7); + check_host(); + check_devs(); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/static-aggr-constructor-destructor-2.C b/libgomp/testsuite/libgomp.c++/static-aggr-constructor-destructor-2.C new file mode 100644 index 0000000..6dd4260 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/static-aggr-constructor-destructor-2.C @@ -0,0 +1,50 @@ +// { dg-do run } +// { dg-additional-options "-fdump-tree-gimple -fdump-tree-optimized" } +// { dg-additional-options -foffload-options=-fdump-tree-optimized { target { offload_device_nvptx || offload_target_amdgcn } } } + +// { dg-final { scan-tree-dump-times "omp_is_initial_device" 1 "gimple" } } +// { dg-final { scan-tree-dump-times "_GLOBAL__off_I_v1" 1 "gimple" } } +// { dg-final { scan-tree-dump-times "__omp_target_static_init_and_destruction" 2 "gimple" } } +// { dg-final { scan-tree-dump-times "__attribute__\\(\\(\[^\n\r]*omp declare target nohost" 2 "gimple" } } + +// { dg-final { scan-tree-dump-not "omp_is_initial_device" "optimized" } } +// { dg-final { scan-tree-dump-not "__omp_target_static_init_and_destruction" "optimized" } } +// FIXME: should be '-not' not '-times' 1: +// { dg-final { scan-tree-dump-times "void _GLOBAL__off_I_" 1 "optimized" } } +// { dg-final { scan-tree-dump-times "__attribute__\\(\\(\[^\n\r]*omp declare target nohost" 1 "optimized" } } + +// { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump-not "omp_initial_device;" "optimized" { target offload_target_amdgcn } } } +// { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump "v1\\._x = 5;" "optimized" { target offload_target_amdgcn } } } +// { dg-final { only_for_offload_target nvptx-none scan-offload-tree-dump-not "omp_initial_device;" "optimized" { target offload_target_nvptx } } } +// { dg-final { only_for_offload_target nvptx-none scan-offload-tree-dump "v1\\._x = 5;" "optimized" { target offload_target_nvptx } } } + + +#include <cassert> + +#pragma omp declare target + +template<typename T> +struct str { + str(T x) : _x(x) { } + T add(str o) { return _x + o._x; } + T _x; +}; + +str<long> v1(5); + +#pragma omp end declare target + +int main() +{ + long res = -1; + str<long> v2(2); + +#pragma omp target map(from:res) + { + res = v1.add(v2); + } + + assert (res == 7); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/static-aggr-constructor-destructor-3.C b/libgomp/testsuite/libgomp.c++/static-aggr-constructor-destructor-3.C new file mode 100644 index 0000000..8d4aff2 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/static-aggr-constructor-destructor-3.C @@ -0,0 +1,36 @@ +// { dg-do run } + +#include <cassert> + +#pragma omp declare target + +struct item { + item(item *p, int v) : prev(p), val(v) { } + int get() { return prev ? prev->get() * val : val; } + item *prev; + int val; +}; + +/* This case demonstrates why constructing on the host and then copying to + the target would be less desirable. With on-target construction, "prev" + for each 'item' will be a device pointer, not a host pointer. */ +item hubert1(nullptr, 3); +item hubert2(&hubert1, 5); +item hubert3(&hubert2, 7); +item hubert4(&hubert3, 11); + +#pragma omp end declare target + +int main() +{ + int res = -1; + +#pragma omp target map(from:res) + { + res = hubert4.get (); + } + + assert (res == 1155); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-is-initial-host-2.c b/libgomp/testsuite/libgomp.c-c++-common/target-is-initial-host-2.c new file mode 100644 index 0000000..313d188 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-is-initial-host-2.c @@ -0,0 +1,43 @@ +/* { dg-do run } */ +/* { dg-additional-options "-fno-builtin-omp_is_initial_device" } */ + +/* Check whether 'omp_is_initial_device()' is NOT compile-time optimized. */ + +/* { dg-additional-options "-fdump-tree-gimple -fdump-tree-optimized" } */ +/* { dg-additional-options -foffload-options=-fdump-tree-optimized { target { offload_device_nvptx || offload_target_amdgcn } } } */ + +/* { dg-final { scan-tree-dump-times "omp_is_initial_device" 1 "gimple" } } */ + +/* { dg-final { scan-tree-dump-times "omp_is_initial_device" 1 "optimized" } } */ + +/* { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump-times "omp_is_initial_device" 1 "optimized" { target offload_target_amdgcn } } } */ +/* { dg-final { only_for_offload_target nvptx-none scan-offload-tree-dump-times "omp_is_initial_device" 1 "optimized" { target offload_target_nvptx } } } */ + + +#include <omp.h> + +int +main () +{ + int is_initial, dev_num, initial; + initial = omp_get_initial_device(); + for (int dev = omp_initial_device; dev <= omp_get_num_devices(); dev++) + { + is_initial = dev_num = 99; + #pragma omp target map(from: is_initial, dev_num) device(dev) + { + is_initial = omp_is_initial_device (); + dev_num = omp_get_device_num (); + } + if (dev == omp_initial_device || dev == initial) + { + if (dev_num != initial || is_initial != 1) + __builtin_abort (); + } + else + { + if (dev_num != dev || is_initial != 0) + __builtin_abort (); + } + } +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-is-initial-host.c b/libgomp/testsuite/libgomp.c-c++-common/target-is-initial-host.c new file mode 100644 index 0000000..423727c --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-is-initial-host.c @@ -0,0 +1,42 @@ +/* { dg-do run } */ + +/* Check whether 'omp_is_initial_device()' is properly compile-time optimized. */ + +/* { dg-additional-options "-fdump-tree-gimple -fdump-tree-optimized" } */ +/* { dg-additional-options -foffload-options=-fdump-tree-optimized { target { offload_device_nvptx || offload_target_amdgcn } } } */ + +/* { dg-final { scan-tree-dump-times "omp_is_initial_device" 1 "gimple" } } */ + +/* { dg-final { scan-tree-dump-not "omp_is_initial_device" "optimized" } } */ + +/* { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump-not "omp_is_initial_device" "optimized" { target offload_target_amdgcn } } } */ +/* { dg-final { only_for_offload_target nvptx-none scan-offload-tree-dump-not "omp_is_initial_device" "optimized" { target offload_target_nvptx } } } */ + + +#include <omp.h> + +int +main () +{ + int is_initial, dev_num, initial; + initial = omp_get_initial_device(); + for (int dev = omp_initial_device; dev <= omp_get_num_devices(); dev++) + { + is_initial = dev_num = 99; + #pragma omp target map(from: is_initial, dev_num) device(dev) + { + is_initial = omp_is_initial_device (); + dev_num = omp_get_device_num (); + } + if (dev == omp_initial_device || dev == initial) + { + if (dev_num != initial || is_initial != 1) + __builtin_abort (); + } + else + { + if (dev_num != dev || is_initial != 0) + __builtin_abort (); + } + } +} diff --git a/libgomp/testsuite/libgomp.fortran/target-is-initial-host-2.f90 b/libgomp/testsuite/libgomp.fortran/target-is-initial-host-2.f90 new file mode 100644 index 0000000..e06ced2 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/target-is-initial-host-2.f90 @@ -0,0 +1,37 @@ +! { dg-additional-options "-fno-builtin-omp_is_initial_device" } + +! Check whether 'omp_is_initial_device()' is NOT compile-time optimized. */ + +! { dg-additional-options "-fdump-tree-gimple -fdump-tree-optimized" } +! { dg-additional-options -foffload-options=-fdump-tree-optimized { target { offload_device_nvptx || offload_target_amdgcn } } } + +! { dg-final { scan-tree-dump-times "omp_is_initial_device" 1 "gimple" } } + +! { dg-final { scan-tree-dump-times "omp_is_initial_device" 1 "optimized" } } + +! { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump-times "omp_is_initial_device" 1 "optimized" { target offload_target_amdgcn } } } +! { dg-final { only_for_offload_target nvptx-none scan-offload-tree-dump-times "omp_is_initial_device" 1 "optimized" { target offload_target_nvptx } } } + + +program main + use omp_lib + implicit none (type, external) + integer :: dev_num, initial, dev + logical :: is_initial + + initial = omp_get_initial_device() + do dev = omp_initial_device, omp_get_num_devices() + dev_num = 99 + !$omp target map(from: is_initial, dev_num) device(dev) + is_initial = omp_is_initial_device () + dev_num = omp_get_device_num () + !$omp end target + if (dev == omp_initial_device .or. dev == initial) then + if (dev_num /= initial .or. .not. is_initial) & + stop 1 + else + if (dev_num /= dev .or. is_initial) & + stop 2 + end if + end do +end diff --git a/libgomp/testsuite/libgomp.fortran/target-is-initial-host.f b/libgomp/testsuite/libgomp.fortran/target-is-initial-host.f new file mode 100644 index 0000000..fec4a3f --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/target-is-initial-host.f @@ -0,0 +1,35 @@ +! Check whether 'omp_is_initial_device()' is properly compile-time optimized. */ + +! { dg-additional-options "-fdump-tree-gimple -fdump-tree-optimized" } +! { dg-additional-options -foffload-options=-fdump-tree-optimized { target { offload_device_nvptx || offload_target_amdgcn } } } + +! { dg-final { scan-tree-dump-times "omp_is_initial_device" 1 "gimple" } } + +! { dg-final { scan-tree-dump-not "omp_is_initial_device" "optimized" } } + +! { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump-not "omp_is_initial_device" "optimized" { target offload_target_amdgcn } } } +! { dg-final { only_for_offload_target nvptx-none scan-offload-tree-dump-not "omp_is_initial_device" "optimized" { target offload_target_nvptx } } } + + + program main + implicit none (type, external) + include "omp_lib.h" + integer :: dev_num, initial, dev + logical :: is_initial + + initial = omp_get_initial_device() + do dev = omp_initial_device, omp_get_num_devices() + dev_num = 99 +!$omp target map(from: is_initial, dev_num) device(dev) + is_initial = omp_is_initial_device () + dev_num = omp_get_device_num () +!$omp end target + if (dev == omp_initial_device .or. dev == initial) then + if (dev_num /= initial .or. .not. is_initial) & + & stop 1 + else + if (dev_num /= dev .or. is_initial) & + & stop 2 + end if + end do + end diff --git a/libgomp/testsuite/libgomp.fortran/target-is-initial-host.f90 b/libgomp/testsuite/libgomp.fortran/target-is-initial-host.f90 new file mode 100644 index 0000000..f8a645f --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/target-is-initial-host.f90 @@ -0,0 +1,35 @@ +! Check whether 'omp_is_initial_device()' is properly compile-time optimized. */ + +! { dg-additional-options "-fdump-tree-gimple -fdump-tree-optimized" } +! { dg-additional-options -foffload-options=-fdump-tree-optimized { target { offload_device_nvptx || offload_target_amdgcn } } } + +! { dg-final { scan-tree-dump-times "omp_is_initial_device" 1 "gimple" } } + +! { dg-final { scan-tree-dump-not "omp_is_initial_device" "optimized" } } + +! { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump-not "omp_is_initial_device" "optimized" { target offload_target_amdgcn } } } +! { dg-final { only_for_offload_target nvptx-none scan-offload-tree-dump-not "omp_is_initial_device" "optimized" { target offload_target_nvptx } } } + + +program main + use omp_lib + implicit none (type, external) + integer :: dev_num, initial, dev + logical :: is_initial + + initial = omp_get_initial_device() + do dev = omp_initial_device, omp_get_num_devices() + dev_num = 99 + !$omp target map(from: is_initial, dev_num) device(dev) + is_initial = omp_is_initial_device () + dev_num = omp_get_device_num () + !$omp end target + if (dev == omp_initial_device .or. dev == initial) then + if (dev_num /= initial .or. .not. is_initial) & + stop 1 + else + if (dev_num /= dev .or. is_initial) & + stop 2 + end if + end do +end |