diff options
Diffstat (limited to 'libgomp/testsuite')
| -rw-r--r-- | libgomp/testsuite/lib/libgomp.exp | 7 | ||||
| -rw-r--r-- | libgomp/testsuite/libgomp.c++/pr101544-1-O0.C | 3 | ||||
| -rw-r--r-- | libgomp/testsuite/libgomp.c++/pr101544-1.C | 82 | ||||
| -rw-r--r-- | libgomp/testsuite/libgomp.c++/pr96835-1-O0.C | 3 | ||||
| -rw-r--r-- | libgomp/testsuite/libgomp.c++/pr96835-1.C | 45 | ||||
| -rw-r--r-- | libgomp/testsuite/libgomp.c-c++-common/get-mapped-ptr-1.c | 2 | ||||
| -rw-r--r-- | libgomp/testsuite/libgomp.c-c++-common/interop-1.c | 43 | ||||
| -rw-r--r-- | libgomp/testsuite/libgomp.c/interop-fr-1.c | 577 | ||||
| -rw-r--r-- | libgomp/testsuite/libgomp.fortran/get-mapped-ptr-1.f90 | 2 | ||||
| -rw-r--r-- | libgomp/testsuite/libgomp.oacc-c++/pr101544-1.C | 6 | ||||
| -rw-r--r-- | libgomp/testsuite/libgomp.oacc-c++/pr96835-1.C | 6 |
11 files changed, 769 insertions, 7 deletions
diff --git a/libgomp/testsuite/lib/libgomp.exp b/libgomp/testsuite/lib/libgomp.exp index fd21371..bc38e3c 100644 --- a/libgomp/testsuite/lib/libgomp.exp +++ b/libgomp/testsuite/lib/libgomp.exp @@ -233,11 +233,8 @@ proc libgomp_init { args } { # error-message parsing machinery. lappend ALWAYS_CFLAGS "additional_flags=-fmessage-length=0" - # Disable caret - lappend ALWAYS_CFLAGS "additional_flags=-fno-diagnostics-show-caret" - - # Disable color diagnostics - lappend ALWAYS_CFLAGS "additional_flags=-fdiagnostics-color=never" + # Disable caret, color, URL diagnostics + lappend ALWAYS_CFLAGS "additional_flags=-fdiagnostics-plain-output" # Help GCC to find offload compilers' 'mkoffload'. global offload_additional_options diff --git a/libgomp/testsuite/libgomp.c++/pr101544-1-O0.C b/libgomp/testsuite/libgomp.c++/pr101544-1-O0.C new file mode 100644 index 0000000..c8a73dc --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/pr101544-1-O0.C @@ -0,0 +1,3 @@ +// { dg-additional-options -O0 } + +#include "pr101544-1.C" diff --git a/libgomp/testsuite/libgomp.c++/pr101544-1.C b/libgomp/testsuite/libgomp.c++/pr101544-1.C new file mode 100644 index 0000000..fcd3e97 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/pr101544-1.C @@ -0,0 +1,82 @@ +// See also '../libgomp.oacc-c++/pr101544-1.C'. +#ifndef ALWAYS_INLINE +# define ALWAYS_INLINE +#endif + +//===--- declare_target_base_class.cpp --------------------------------------===// +// +// OpenMP API Version 4.5 Nov 2015 +// +// This test was suggested by members of NERSC. This test defines a declare +// target region which includes only a base class and a 'concrete' device +// pointer. +// +// Test suggestion comes from Chris Daily and Rahulkumar Gayatri from NERSC +////===----------------------------------------------------------------------===// + +#include <new> +#include <vector> +#include <iostream> + +#pragma omp declare target +//#pragma acc routine //TODO error: '#pragma acc routine' not immediately followed by function declaration or definition +class S { +public: + //#pragma acc routine //TODO error: '#pragma acc routine' must be at file scope + ALWAYS_INLINE + S() : _devPtr(nullptr) {} + //#pragma acc routine //TODO error: '#pragma acc routine' must be at file scope + ALWAYS_INLINE + double sag(double x, double y) { + return x + y; + } + S* cloneToDevice() { + S* ptr; +#pragma omp target map(ptr) +#pragma acc serial copy(ptr) + { + ptr = new S(); + } + _devPtr = ptr; + return ptr; + } +private: + S* _devPtr; +}; +//#pragma acc routine (S) //TODO error: 'class S' does not refer to a function +//#pragma acc routine (S::S) //TODO error: '#pragma acc routine' names a set of overloads +//#pragma acc routine (S::sag) //TODO error: '#pragma acc routine' names a set of overloads +#pragma omp end declare target + +int main() { + int errors = 0; + + S s; + S* devPtr = s.cloneToDevice(); + + std::vector<double> in(10, 0.0); + for(int i = 0; i < 10; i++) { + in[i] = i; + } + + std::vector<double> out(10, 0.0); + + double* inptr = in.data(); + double* outptr = out.data(); + +#pragma omp target teams distribute parallel for map(inptr[:10], outptr[:10]) is_device_ptr(devPtr) +#pragma acc parallel loop copy(inptr[:10], outptr[:10]) deviceptr(devPtr) + for(int i = 0; i < 10; i++) { + outptr[i] = devPtr->sag(inptr[i], inptr[i]); + } + + for(int i = 0; i < 10; i++) { + if (out[i] != i * 2) + { + ++errors; + std::cerr << "ERROR: " << "i = " << i << ": " << out[i] << " != " << (i * 2) << "\n"; + } + } + + return errors ? 1 : 0; +} diff --git a/libgomp/testsuite/libgomp.c++/pr96835-1-O0.C b/libgomp/testsuite/libgomp.c++/pr96835-1-O0.C new file mode 100644 index 0000000..85e4290 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/pr96835-1-O0.C @@ -0,0 +1,3 @@ +// { dg-additional-options -O0 } + +#include "pr96835-1.C" diff --git a/libgomp/testsuite/libgomp.c++/pr96835-1.C b/libgomp/testsuite/libgomp.c++/pr96835-1.C new file mode 100644 index 0000000..c9f6475 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/pr96835-1.C @@ -0,0 +1,45 @@ +// See also '../libgomp.oacc-c++/pr96835-1.C'. +#ifndef ALWAYS_INLINE +# define ALWAYS_INLINE +#endif + +#pragma omp declare target + +template<int sz> +struct vector { + int values_[sz]; + vector(); + ALWAYS_INLINE + vector(int const& init_val); + ALWAYS_INLINE + int dot(vector o) { + int res = 0; + for (int i = 0; i < sz; ++ i) + res += values_[i] * o.values_[i]; + return res; + } +}; + +template<int sz> +vector<sz>::vector(int const& init_val) { + for (int i = 0; i < sz; ++ i) values_[i] = init_val; +} +template<int sz> +vector<sz>::vector() : vector(0) { +} + +#pragma omp end declare target + +int main() { + int res = 0; + #pragma omp target map(from:res) + #pragma acc serial copyout(res) + { + vector<4> v1(1); + vector<4> v2(2); + res = v1.dot(v2); + } + if (res != 8) + __builtin_abort(); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/get-mapped-ptr-1.c b/libgomp/testsuite/libgomp.c-c++-common/get-mapped-ptr-1.c index 4708ae8..90d1a72 100644 --- a/libgomp/testsuite/libgomp.c-c++-common/get-mapped-ptr-1.c +++ b/libgomp/testsuite/libgomp.c-c++-common/get-mapped-ptr-1.c @@ -21,7 +21,7 @@ main () if (omp_target_associate_ptr (q, p, sizeof (int), 0, d) != 0) return 0; - if (omp_get_mapped_ptr (q, -5) != NULL) + if (omp_get_mapped_ptr (q, -6) != NULL) abort (); if (omp_get_mapped_ptr (q, omp_get_num_devices () + 1) != NULL) diff --git a/libgomp/testsuite/libgomp.c-c++-common/interop-1.c b/libgomp/testsuite/libgomp.c-c++-common/interop-1.c new file mode 100644 index 0000000..149f387 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/interop-1.c @@ -0,0 +1,43 @@ +/* { dg-do run } */ + +#include <omp.h> +#include <stdlib.h> + +int +main () +{ + int dev = omp_get_num_devices (); + int x[6]; + omp_interop_t obj1 = omp_interop_none; +#pragma omp interop init(targetsync : obj1) depend(in : x) device(dev) + if (obj1 != omp_interop_none) + abort (); + +#pragma omp interop use(obj1) +#pragma omp interop destroy(obj1) depend(out : x) + if (obj1 != omp_interop_none) + abort (); + + omp_set_default_device (dev); + omp_interop_t obj2; + +#pragma omp interop init( \ + target, targetsync, \ + prefer_type({fr("hip"), attr("ompx_gnu_prio:1", "ompx_gnu_debug")}, \ + {attr("ompx_gnu_nicest"), attr("ompx_something")}) : obj1, \ + obj2) nowait + if (obj1 != omp_interop_none || obj2 != omp_interop_none) + abort (); +#pragma omp interop use(obj1, obj2) nowait + + omp_interop_t obj3 = __omp_interop_t_max__; + +#pragma omp interop init(target : obj3) use(obj2) destroy(obj1) nowait + if (obj1 != omp_interop_none || obj3 != omp_interop_none) + abort (); +#pragma omp interop destroy(obj3, obj2) nowait + if (obj2 != omp_interop_none || obj3 != omp_interop_none) + abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/interop-fr-1.c b/libgomp/testsuite/libgomp.c/interop-fr-1.c new file mode 100644 index 0000000..9310c95 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/interop-fr-1.c @@ -0,0 +1,577 @@ +/* { dg-do run } */ + +#include <assert.h> +#include <stdio.h> +#include <stdlib.h> +#include <string.h> +#include <omp.h> +#include "../libgomp.c-c++-common/on_device_arch.h" + +#define DEFAULT_DEVICE -99 + +/* The following assumes that when a nvptx device is available, + cuda/cuda_driver/hip are supported. + And that likewise when a gcn device is available that the + plugin also can not only the HSA but also the HIP library + such that hsa/hip are supported. + For the host, omp_interop_none is expected. + + Otherwise, it only does some basic tests without checking + that the returned result really makes sense. */ + +void check_host (int); +void check_nvptx (int); +void check_gcn (int); + +void check_type (omp_interop_t obj) +{ + const char *type; + + type = omp_get_interop_type_desc (obj, omp_ipr_fr_id); + if (obj != omp_interop_none) + assert (strcmp (type, "omp_interop_t") == 0); + else + assert (type == NULL); + + type = omp_get_interop_type_desc (obj, omp_ipr_fr_name); + if (obj != omp_interop_none) + assert (strcmp (type, "const char *") == 0); + else + assert (type == NULL); + + type = omp_get_interop_type_desc (obj, omp_ipr_vendor); + if (obj != omp_interop_none) + assert (strcmp (type, "int") == 0); + else + assert (type == NULL); + + type = omp_get_interop_type_desc (obj, omp_ipr_vendor_name); + if (obj != omp_interop_none) + assert (strcmp (type, "const char *") == 0); + else + assert (type == NULL); + + type = omp_get_interop_type_desc (obj, omp_ipr_device_num); + if (obj != omp_interop_none) + assert (strcmp (type, "int") == 0); + else + assert (type == NULL); + + if (obj != omp_interop_none) + return; + assert (omp_get_interop_type_desc (obj, omp_ipr_platform) == NULL); + assert (omp_get_interop_type_desc (obj, omp_ipr_device) == NULL); + assert (omp_get_interop_type_desc (obj, omp_ipr_device_context) == NULL); + assert (omp_get_interop_type_desc (obj, omp_ipr_targetsync) == NULL); +} + +void +do_check (int dev) +{ + int num_dev = omp_get_num_devices (); + const char *dev_type; + if (dev != DEFAULT_DEVICE) + omp_set_default_device (dev); + int is_nvptx = on_device_arch_nvptx (); + int is_gcn = on_device_arch_gcn (); + int is_host; + + if (dev != DEFAULT_DEVICE) + is_host = dev == -1 || dev == num_dev; + else + { + int def_dev = omp_get_default_device (); + is_host = def_dev == -1 || def_dev == num_dev; + } + + assert (is_nvptx + is_gcn + is_host == 1); + + if (num_dev > 0 && dev != DEFAULT_DEVICE) + { + if (is_host) + omp_set_default_device (0); + else + omp_set_default_device (-1); + } + + if (is_host) + dev_type = "host"; + else if (is_nvptx) + dev_type = "nvptx"; + else if (is_gcn) + dev_type = "gcn"; + + printf ("Running on the %s device (%d)\n", dev_type, dev); + if (is_host) + check_host (dev); + else if (is_nvptx) + check_nvptx (dev); + else if (is_gcn) + check_gcn (dev); +} + + +void +check_host (int dev) +{ + omp_interop_t obj = (omp_interop_t) -1L; + if (dev == DEFAULT_DEVICE) { + #pragma omp interop init(target : obj) + } else { + #pragma omp interop init(target : obj) device(dev) + } + assert (obj == omp_interop_none); + check_type (obj); + + obj = (omp_interop_t) -1L; + if (dev == DEFAULT_DEVICE) { + #pragma omp interop init(target, prefer_type({attr("ompx_foo")}, {attr("ompx_bar"), fr("cuda"), attr("ompx_foobar")},{fr("cuda_driver")}, {fr("hip")}, {fr("hsa")}) : obj) + } else { + #pragma omp interop init(target, prefer_type({attr("ompx_foo")}, {attr("ompx_bar"), fr("cuda"), attr("ompx_foobar")},{fr("cuda_driver")}, {fr("hip")}, {fr("hsa")}) : obj) device(dev) + } + assert (obj == omp_interop_none); + check_type (obj); + + obj = (omp_interop_t) -1L; + if (dev == DEFAULT_DEVICE) { + #pragma omp interop init(targetsync : obj) + } else { + #pragma omp interop init(targetsync : obj) device(dev) + } + assert (obj == omp_interop_none); + check_type (obj); + + obj = (omp_interop_t) -1L; + if (dev == DEFAULT_DEVICE) { + #pragma omp interop init(targetsync, prefer_type("cuda","cuda_driver", "hip", "hsa") : obj) + } else { + #pragma omp interop init(targetsync, prefer_type("cuda","cuda_driver", "hip", "hsa") : obj) device(dev) + } + assert (obj == omp_interop_none); + check_type (obj); +} + + +void +check_nvptx (int dev) +{ + for (int variant = 0; variant <= 7; variant++) + { + omp_interop_t obj = (omp_interop_t) -1L; + switch (variant) + { + /* Expect 'cuda'. */ + case 0: + { + if (dev == DEFAULT_DEVICE) { + #pragma omp interop init(target : obj) + } else { + #pragma omp interop init(target : obj) device(dev) + } + break; + } + case 1: + { + if (dev == DEFAULT_DEVICE) { + #pragma omp interop init(targetsync : obj) + } else { + #pragma omp interop init(targetsync : obj) device(dev) + } + break; + } + case 2: + { + if (dev == DEFAULT_DEVICE) { + #pragma omp interop init(target, prefer_type({attr("ompx_foo")}, {fr("hsa")}, {attr("ompx_bar"), fr("cuda"), attr("ompx_foobar")},{fr("cuda_driver")}, {fr("hip")}) : obj) + } else { + #pragma omp interop init(target, prefer_type({attr("ompx_foo")}, {fr("hsa")}, {attr("ompx_bar"), fr("cuda"), attr("ompx_foobar")},{fr("cuda_driver")}, {fr("hip")}) : obj) device(dev) + } + break; + } + case 3: + { + if (dev == DEFAULT_DEVICE) { + #pragma omp interop init(targetsync, prefer_type("hsa", "cuda", "cuda_driver", "hip") : obj) + } else { + #pragma omp interop init(targetsync, prefer_type("hsa", "cuda", "cuda_driver", "hip") : obj) device(dev) + } + break; + } + + /* Expect 'cuda_driver'. */ + case 4: + { + if (dev == DEFAULT_DEVICE) { + #pragma omp interop init(target, prefer_type("hsa", "cuda_driver", "hip", "cuda") : obj) + } else { + #pragma omp interop init(target, prefer_type("hsa", "cuda_driver", "hip", "cuda") : obj) device(dev) + } + break; + } + case 5: + { + if (dev == DEFAULT_DEVICE) { + #pragma omp interop init(targetsync, prefer_type("hsa", "cuda_driver", "hip", "cuda") : obj) + } else { + #pragma omp interop init(targetsync, prefer_type("hsa", "cuda_driver", "hip", "cuda") : obj) device(dev) + } + break; + } + + /* Expect 'hip'. */ + case 6: + { + if (dev == DEFAULT_DEVICE) { + #pragma omp interop init(target, prefer_type("hsa", "hip", "cuda", "cuda_driver") : obj) + } else { + #pragma omp interop init(target, prefer_type("hsa", "hip", "cuda", "cuda_driver") : obj) device(dev) + } + break; + } + case 7: + { + if (dev == DEFAULT_DEVICE) { + #pragma omp interop init(targetsync, prefer_type("hsa", "hip", "cuda", "cuda_driver") : obj) + } else { + #pragma omp interop init(targetsync, prefer_type("hsa", "hip", "cuda", "cuda_driver") : obj) device(dev) + } + break; + } + default: + abort (); + } + assert (obj != omp_interop_none && obj != (omp_interop_t) -1L); + + omp_interop_rc_t ret_code = omp_irc_no_value; + omp_interop_fr_t fr = (omp_interop_fr_t) omp_get_interop_int (obj, omp_ipr_fr_id, &ret_code); + + assert (ret_code == omp_irc_success); + if (variant >= 0 && variant <= 3) + assert (fr == omp_ifr_cuda); + else if (variant <= 5) + assert (fr == omp_ifr_cuda_driver); + else if (variant <= 7) + assert (fr == omp_ifr_hip); + else + assert (0); + + ret_code = omp_irc_no_value; + const char *fr_name = omp_get_interop_str (obj, omp_ipr_fr_name, &ret_code); + + assert (ret_code == omp_irc_success); + if (fr == omp_ifr_cuda) + assert (strcmp (fr_name, "cuda") == 0); + else if (fr == omp_ifr_cuda_driver) + assert (strcmp (fr_name, "cuda_driver") == 0); + else if (fr == omp_ifr_hip) + assert (strcmp (fr_name, "hip") == 0); + else + assert (0); + + ret_code = omp_irc_no_value; + int vendor = (int) omp_get_interop_int (obj, omp_ipr_vendor, &ret_code); + assert (ret_code == omp_irc_success); + assert (vendor == 11); /* Nvidia */ + + ret_code = omp_irc_no_value; + const char *vendor_name = omp_get_interop_str (obj, omp_ipr_vendor_name, &ret_code); + assert (ret_code == omp_irc_success); + assert (strcmp (vendor_name, "nvidia") == 0); + + ret_code = omp_irc_no_value; + int dev_num = (int) omp_get_interop_int (obj, omp_ipr_device_num, &ret_code); + assert (ret_code == omp_irc_success); + if (dev == DEFAULT_DEVICE) + assert (dev_num == omp_get_default_device ()); + else + assert (dev_num == dev); + + /* Platform: N/A. */ + ret_code = omp_irc_success; + (void) omp_get_interop_int (obj, omp_ipr_platform, &ret_code); + assert (ret_code == omp_irc_no_value); + ret_code = omp_irc_success; + (void) omp_get_interop_ptr (obj, omp_ipr_platform, &ret_code); + assert (ret_code == omp_irc_no_value); + ret_code = omp_irc_success; + (void) omp_get_interop_str (obj, omp_ipr_platform, &ret_code); + assert (ret_code == omp_irc_no_value); + + /* Device: int / CUdevice / hipDevice_t -- all internally an 'int'. */ + ret_code = omp_irc_no_value; + int fr_device = (int) omp_get_interop_int (obj, omp_ipr_device, &ret_code); + + /* CUDA also starts from 0 and goes to < n with cudaGetDeviceCount(&cn). */ + assert (ret_code == omp_irc_success); + assert (fr_device >= 0 && fr_device < omp_get_num_devices ()); + + /* Device context: N/A / CUcontext / hipCtx_t -- a pointer. */ + ret_code = omp_irc_out_of_range; + void *ctx = omp_get_interop_ptr (obj, omp_ipr_device_context, &ret_code); + + if (fr == omp_ifr_cuda) + { + assert (ret_code == omp_irc_no_value); + assert (ctx == NULL); + } + else + { + assert (ret_code == omp_irc_success); + assert (ctx != NULL); + } + + /* Stream/targetsync: cudaStream_t / CUstream / hipStream_t -- a pointer. */ + ret_code = omp_irc_out_of_range; + void *stream = omp_get_interop_ptr (obj, omp_ipr_targetsync, &ret_code); + + if (variant % 2 == 0) /* no targetsync */ + { + assert (ret_code == omp_irc_no_value); + assert (stream == NULL); + } + else + { + assert (ret_code == omp_irc_success); + assert (stream != NULL); + } + + check_type (obj); + if (fr == omp_ifr_cuda) + { + assert (strcmp (omp_get_interop_type_desc (obj, omp_ipr_platform), "N/A") == 0); + assert (strcmp (omp_get_interop_type_desc (obj, omp_ipr_device), "int") == 0); + assert (strcmp (omp_get_interop_type_desc (obj, omp_ipr_device_context), "N/A") == 0); + assert (strcmp (omp_get_interop_type_desc (obj, omp_ipr_targetsync), "cudaStream_t") == 0); + } + else if (fr == omp_ifr_cuda_driver) + { + assert (strcmp (omp_get_interop_type_desc (obj, omp_ipr_platform), "N/A") == 0); + assert (strcmp (omp_get_interop_type_desc (obj, omp_ipr_device), "CUdevice") == 0); + assert (strcmp (omp_get_interop_type_desc (obj, omp_ipr_device_context), "CUcontext") == 0); + assert (strcmp (omp_get_interop_type_desc (obj, omp_ipr_targetsync), "CUstream") == 0); + } + else + { + assert (strcmp (omp_get_interop_type_desc (obj, omp_ipr_platform), "N/A") == 0); + assert (strcmp (omp_get_interop_type_desc (obj, omp_ipr_device), "hipDevice_t") == 0); + assert (strcmp (omp_get_interop_type_desc (obj, omp_ipr_device_context), "hipCtx_t") == 0); + assert (strcmp (omp_get_interop_type_desc (obj, omp_ipr_targetsync), "hipStream_t") == 0); + } + + if (dev == DEFAULT_DEVICE) { + #pragma omp interop use(obj) + #pragma omp interop destroy(obj) + } else { + #pragma omp interop use(obj) device(dev) + #pragma omp interop destroy(obj) device(dev) + } + } +} + + +void +check_gcn (int dev) +{ + for (int variant = 0; variant <= 5; variant++) + { + omp_interop_t obj = (omp_interop_t) -1L; + switch (variant) + { + /* Expect 'hip'. */ + case 0: + { + if (dev == DEFAULT_DEVICE) { + #pragma omp interop init(target : obj) + } else { + #pragma omp interop init(target : obj) device(dev) + } + break; + } + case 1: + { + if (dev == DEFAULT_DEVICE) { + #pragma omp interop init(targetsync : obj) + } else { + #pragma omp interop init(targetsync : obj) device(dev) + } + break; + } + case 2: + { + if (dev == DEFAULT_DEVICE) { + #pragma omp interop init(target, prefer_type({attr("ompx_foo")}, {fr("cuda")}, {fr("cuda_driver")}, {attr("ompx_bar"), fr("hip"), attr("ompx_foobar")},{fr("hsa")}) : obj) + } else { + #pragma omp interop init(target, prefer_type({attr("ompx_foo")}, {fr("cuda")}, {fr("cuda_driver")}, {attr("ompx_bar"), fr("hip"), attr("ompx_foobar")},{fr("hsa")}) : obj) device(dev) + } + break; + } + case 3: + { + if (dev == DEFAULT_DEVICE) { + #pragma omp interop init(targetsync, prefer_type("cuda", "cuda_driver", "hip", "hsa") : obj) + } else { + #pragma omp interop init(targetsync, prefer_type("cuda", "cuda_driver", "hip", "hsa") : obj) device(dev) + } + break; + } + + /* Expect 'hsa'. */ + case 4: + { + if (dev == DEFAULT_DEVICE) { + #pragma omp interop init(target, prefer_type("cuda", "cuda_driver", "hsa", "hip") : obj) + } else { + #pragma omp interop init(target, prefer_type("cuda", "cuda_driver", "hsa", "hip") : obj) device(dev) + } + break; + } + case 5: + { + if (dev == DEFAULT_DEVICE) { + #pragma omp interop init(targetsync, prefer_type("cuda", "cuda_driver", "hsa", "hip") : obj) + } else { + #pragma omp interop init(targetsync, prefer_type("cuda", "cuda_driver", "hsa", "hip") : obj) device(dev) + } + break; + } + default: + abort (); + } + assert (obj != omp_interop_none && obj != (omp_interop_t) -1L); + + omp_interop_rc_t ret_code = omp_irc_no_value; + omp_interop_fr_t fr = (omp_interop_fr_t) omp_get_interop_int (obj, omp_ipr_fr_id, &ret_code); + + assert (ret_code == omp_irc_success); + if (variant >= 0 && variant <= 3) + assert (fr == omp_ifr_hip); + else if (variant <= 5) + assert (fr == omp_ifr_hsa); + else + assert (0); + + ret_code = omp_irc_no_value; + const char *fr_name = omp_get_interop_str (obj, omp_ipr_fr_name, &ret_code); + + assert (ret_code == omp_irc_success); + if (fr == omp_ifr_hip) + assert (strcmp (fr_name, "hip") == 0); + else if (fr == omp_ifr_hsa) + assert (strcmp (fr_name, "hsa") == 0); + else + assert (0); + + ret_code = omp_irc_no_value; + int vendor = (int) omp_get_interop_int (obj, omp_ipr_vendor, &ret_code); + assert (ret_code == omp_irc_success); + assert (vendor == 1); /* Amd */ + + ret_code = omp_irc_no_value; + const char *vendor_name = omp_get_interop_str (obj, omp_ipr_vendor_name, &ret_code); + assert (ret_code == omp_irc_success); + assert (strcmp (vendor_name, "amd") == 0); + + ret_code = omp_irc_no_value; + int dev_num = (int) omp_get_interop_int (obj, omp_ipr_device_num, &ret_code); + assert (ret_code == omp_irc_success); + if (dev == DEFAULT_DEVICE) + assert (dev_num == omp_get_default_device ()); + else + assert (dev_num == dev); + + /* Platform: N/A. */ + ret_code = omp_irc_success; + (void) omp_get_interop_int (obj, omp_ipr_platform, &ret_code); + assert (ret_code == omp_irc_no_value); + ret_code = omp_irc_success; + (void) omp_get_interop_ptr (obj, omp_ipr_platform, &ret_code); + assert (ret_code == omp_irc_no_value); + ret_code = omp_irc_success; + (void) omp_get_interop_str (obj, omp_ipr_platform, &ret_code); + assert (ret_code == omp_irc_no_value); + + /* Device: hipDevice_t / hsa_agent_t* -- hip is internally an 'int'. */ + ret_code = omp_irc_no_value; + if (fr == omp_ifr_hip) + { + /* HIP also starts from 0 and goes to < n as with cudaGetDeviceCount(&cn). */ + int fr_device = (int) omp_get_interop_int (obj, omp_ipr_device, &ret_code); + assert (ret_code == omp_irc_success); + assert (fr_device >= 0 && fr_device < omp_get_num_devices ()); + } + else + { + void *agent = omp_get_interop_ptr (obj, omp_ipr_device, &ret_code); + assert (ret_code == omp_irc_success); + assert (agent != NULL); + } + + /* Device context: hipCtx_t / N/A -- a pointer. */ + ret_code = omp_irc_out_of_range; + void *ctx = omp_get_interop_ptr (obj, omp_ipr_device_context, &ret_code); + if (fr == omp_ifr_hip) + { + assert (ret_code == omp_irc_success); + assert (ctx != NULL); + } + else + { + assert (ret_code == omp_irc_no_value); + assert (ctx == NULL); + } + + /* Stream/targetsync: cudaStream_t / CUstream / hipStream_t -- a pointer. */ + ret_code = omp_irc_out_of_range; + void *stream = omp_get_interop_ptr (obj, omp_ipr_targetsync, &ret_code); + + if (variant % 2 == 0) /* no targetsync */ + { + assert (ret_code == omp_irc_no_value); + assert (stream == NULL); + } + else + { + assert (ret_code == omp_irc_success); + assert (stream != NULL); + } + + check_type (obj); + if (fr == omp_ifr_hip) + { + assert (strcmp (omp_get_interop_type_desc (obj, omp_ipr_platform), "N/A") == 0); + assert (strcmp (omp_get_interop_type_desc (obj, omp_ipr_device), "hipDevice_t") == 0); + assert (strcmp (omp_get_interop_type_desc (obj, omp_ipr_device_context), "hipCtx_t") == 0); + assert (strcmp (omp_get_interop_type_desc (obj, omp_ipr_targetsync), "hipStream_t") == 0); + } + else + { + assert (strcmp (omp_get_interop_type_desc (obj, omp_ipr_platform), "N/A") == 0); + assert (strcmp (omp_get_interop_type_desc (obj, omp_ipr_device), "hsa_agent_t *") == 0); + assert (strcmp (omp_get_interop_type_desc (obj, omp_ipr_device_context), "N/A") == 0); + assert (strcmp (omp_get_interop_type_desc (obj, omp_ipr_targetsync), "hsa_queue_t *") == 0); + } + + if (dev == DEFAULT_DEVICE) { + #pragma omp interop use(obj) + #pragma omp interop destroy(obj) + } else { + #pragma omp interop use(obj) device(dev) + #pragma omp interop destroy(obj) device(dev) + } + } +} + + +int +main () +{ + do_check (DEFAULT_DEVICE); + int ndev = omp_get_num_devices (); + for (int dev = -1; dev < ndev; dev++) + do_check (dev); + for (int dev = -1; dev < ndev; dev++) + { + omp_set_default_device (dev); + do_check (DEFAULT_DEVICE); + } +} diff --git a/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-1.f90 b/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-1.f90 index 9b6334f..be60acc 100644 --- a/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-1.f90 +++ b/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-1.f90 @@ -19,7 +19,7 @@ program main if (omp_target_associate_ptr (c_loc (q), p, c_sizeof (q), & 0_c_size_t, d) == 0) then - if(c_associated (omp_get_mapped_ptr (c_loc (q), -5))) & + if(c_associated (omp_get_mapped_ptr (c_loc (q), -6))) & stop 1 if(c_associated (omp_get_mapped_ptr (c_loc (q), & diff --git a/libgomp/testsuite/libgomp.oacc-c++/pr101544-1.C b/libgomp/testsuite/libgomp.oacc-c++/pr101544-1.C new file mode 100644 index 0000000..d4d28a6 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c++/pr101544-1.C @@ -0,0 +1,6 @@ +// { dg-additional-options -fno-inline } for stable results regarding OpenACC 'routine'. +// But actually, as none of the '#pragma acc routine' syntax is accepted, force inlining: +#define ALWAYS_INLINE __attribute__((always_inline)) + +#include "../libgomp.c++/pr101544-1.C" +//TODO { dg-prune-output {using 'vector_length \(32\)', ignoring 1} } diff --git a/libgomp/testsuite/libgomp.oacc-c++/pr96835-1.C b/libgomp/testsuite/libgomp.oacc-c++/pr96835-1.C new file mode 100644 index 0000000..0a6ee22 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c++/pr96835-1.C @@ -0,0 +1,6 @@ +// { dg-additional-options -fno-inline } for stable results regarding OpenACC 'routine'. +// But actually, as none of the '#pragma acc routine' syntax is accepted (see '../libgomp.c++/pr101544-1.C'), force inlining: +#define ALWAYS_INLINE inline __attribute__((always_inline)) + +#include "../libgomp.c++/pr96835-1.C" +//TODO { dg-prune-output {using 'vector_length \(32\)', ignoring 1} } |