9f9d128f45
OpenMP permits that a 'target device(ancestor:1)' is called without being enclosed in a target region - using the current device (i.e. the host) in that case. This commit adds a testcase for this. In case of nvptx, the missing on-device 'GOMP_target_ext' call causes that it and also the associated on-device GOMP_REV_OFFLOAD_VAR variable are not linked in from nvptx's libgomp.a. Thus, handle the failing cuModuleGetGlobal gracefully by disabling reverse offload and assuming that the failure is fine. libgomp/ChangeLog: * plugin/plugin-nvptx.c (GOMP_OFFLOAD_load_image): Use unsigned int for 'i' to match 'fn_entries'; regard absent GOMP_REV_OFFLOAD_VAR as valid and the code having no reverse-offload code. * testsuite/libgomp.c-c++-common/reverse-offload-2.c: New test.
49 lines
1 KiB
C
49 lines
1 KiB
C
/* { dg-do run } */
|
|
/* { dg-additional-options "-foffload-options=nvptx-none=-misa=sm_35" { target { offload_target_nvptx } } } */
|
|
|
|
#pragma omp requires reverse_offload
|
|
|
|
int
|
|
main ()
|
|
{
|
|
int A[10];
|
|
int y;
|
|
|
|
for (int i = 0; i < 10; i++)
|
|
A[i] = 2*i;
|
|
|
|
y = 42;
|
|
|
|
/* Pointlessly copy to the default device. */
|
|
#pragma omp target data map(to: A)
|
|
{
|
|
/* Not enclosed in a target region (= i.e. running on the host); the
|
|
following is valid - it runs on the current device (= host). */
|
|
#pragma omp target device ( ancestor:1 ) firstprivate(y) map(to: A)
|
|
{
|
|
if (y != 42)
|
|
__builtin_abort ();
|
|
for (int i = 0; i < 10; i++)
|
|
if (A[i] != 2*i)
|
|
__builtin_abort ();
|
|
for (int i = 0; i < 10; i++)
|
|
if (A[i] != 2*i)
|
|
A[i] = 4*i;
|
|
y = 31;
|
|
}
|
|
|
|
if (y != 42)
|
|
__builtin_abort ();
|
|
for (int i = 0; i < 10; i++)
|
|
if (A[i] != 2*i)
|
|
__builtin_abort ();
|
|
}
|
|
|
|
if (y != 42)
|
|
__builtin_abort ();
|
|
for (int i = 0; i < 10; i++)
|
|
if (A[i] != 2*i)
|
|
__builtin_abort ();
|
|
|
|
return 0;
|
|
}
|