mirror of
https://gcc.gnu.org/git/gcc.git
synced 2024-11-23 10:54:07 +08:00
6b43f556f3
Reverse offload requests at least -misa=sm_35; with this patch, a warning instead of an error is shown, still permitting reverse offload for all other configured device types. This is achieved by not calling GOMP_offload_register_ver (and stopping generating pointless 'static const char' variables, once known.) The tool_name as progname changes adds "nvptx " and "gcn " to the "mkoffload: warning/error:" diagnostic. gcc/ChangeLog: * config/nvptx/mkoffload.cc (process): Replace a fatal_error by a warning + not enabling offloading if -misa=sm_30 prevents reverse offload. (main): Use tool_name as progname for diagnostic. * config/gcn/mkoffload.cc (main): Likewise. libgomp/ChangeLog: * libgomp.texi (Offload-Target Specifics: nvptx): Document that reverse offload requires >= -march=sm_35. * testsuite/libgomp.c-c++-common/requires-4.c: Build for nvptx with -misa=sm_35. * testsuite/libgomp.c-c++-common/requires-5.c: Likewise. * testsuite/libgomp.c-c++-common/requires-6.c: Likewise. * testsuite/libgomp.c-c++-common/reverse-offload-1.c: Likewise. * testsuite/libgomp.fortran/reverse-offload-1.f90: Likewise. * testsuite/libgomp.c/reverse-offload-sm30.c: New test.
85 lines
1.8 KiB
C
85 lines
1.8 KiB
C
/* { dg-do run } */
|
|
/* { dg-additional-options "-foffload-options=nvptx-none=-misa=sm_35" { target { offload_target_nvptx } } } */
|
|
/* { dg-additional-sources reverse-offload-1-aux.c } */
|
|
|
|
/* Check that reverse offload works in particular:
|
|
- no code is generated on the device side (i.e. no
|
|
implicit declare target of called functions and no
|
|
code gen for the target-region body)
|
|
-> would otherwise fail due to 'add_3' symbol
|
|
- Plus the usual (compiles, runs, produces correct result)
|
|
|
|
Note: Running also the non-reverse-offload target regions
|
|
on the host (host fallback) is valid and will pass. */
|
|
|
|
#pragma omp requires reverse_offload
|
|
|
|
extern int add_3 (int);
|
|
|
|
static int global_var = 5;
|
|
|
|
void
|
|
check_offload (int *x, int *y)
|
|
{
|
|
*x = add_3 (*x);
|
|
*y = add_3 (*y);
|
|
}
|
|
|
|
#pragma omp declare target
|
|
void
|
|
tg_fn (int *x, int *y)
|
|
{
|
|
int x2 = *x, y2 = *y;
|
|
if (x2 != 2 || y2 != 3)
|
|
__builtin_abort ();
|
|
x2 = x2 + 2;
|
|
y2 = y2 + 7;
|
|
|
|
#pragma omp target device(ancestor : 1) map(tofrom: x2)
|
|
check_offload(&x2, &y2);
|
|
|
|
if (x2 != 2+2+3 || y2 != 3 + 7)
|
|
__builtin_abort ();
|
|
*x = x2, *y = y2;
|
|
}
|
|
#pragma omp end declare target
|
|
|
|
void
|
|
my_func (int *x, int *y)
|
|
{
|
|
if (global_var != 5)
|
|
__builtin_abort ();
|
|
global_var = 242;
|
|
*x = 2*add_3(*x);
|
|
*y = 3*add_3(*y);
|
|
}
|
|
|
|
int
|
|
main ()
|
|
{
|
|
#pragma omp target
|
|
{
|
|
int x = 2, y = 3;
|
|
tg_fn (&x, &y);
|
|
}
|
|
|
|
#pragma omp target
|
|
{
|
|
int x = -2, y = -1;
|
|
#pragma omp target device ( ancestor:1 ) firstprivate(y) map(tofrom:x)
|
|
{
|
|
if (x != -2 || y != -1)
|
|
__builtin_abort ();
|
|
my_func (&x, &y);
|
|
if (x != 2*(3-2) || y != 3*(3-1))
|
|
__builtin_abort ();
|
|
}
|
|
if (x != 2*(3-2) || y != -1)
|
|
__builtin_abort ();
|
|
}
|
|
|
|
if (global_var != 242)
|
|
__builtin_abort ();
|
|
return 0;
|
|
}
|