33c4e46624
As -foffload={options,targets,targets=options} is very convoluted,
it has been split into -foffload=targets (supporting the old syntax
for backward compatibilty) and -foffload-options={options,target=options}.
Only the new syntax is documented.
Additionally, -foffload=default is supported, which can reset the
devices after -foffload=disable / -foffload=targets to the default,
if needed.
gcc/ChangeLog:
PR other/67300
* common.opt (-foffload=): Update description.
(-foffload-options=): New.
* doc/invoke.texi (C Language Options): Document
-foffload and -foffload-options.
* gcc.c (check_offload_target_name): New, split off from
handle_foffload_option.
(check_foffload_target_names): New.
(handle_foffload_option): Handle -foffload=default.
(driver_handle_option): Update for -foffload-options.
* lto-opts.c (lto_write_options): Use -foffload-options
instead of -foffload.
* lto-wrapper.c (merge_and_complain, append_offload_options):
Likewise.
* opts.c (common_handle_option): Likewise.
libgomp/ChangeLog:
PR other/67300
* testsuite/libgomp.c-c++-common/reduction-16.c: Replace
-foffload=nvptx-none= by -foffload-options=nvptx-none= to
avoid disabling other offload targets.
* testsuite/libgomp.c-c++-common/reduction-5.c: Likewise.
* testsuite/libgomp.c-c++-common/reduction-6.c: Likewise.
* testsuite/libgomp.c/target-44.c: Likewise.
193 lines
4.6 KiB
C
193 lines
4.6 KiB
C
/* { dg-additional-options "-foffload-options=nvptx-none=-latomic" { target { offload_target_nvptx } } } */
|
|
/* C / C++'s logical AND and OR operators take any scalar argument
|
|
which compares (un)equal to 0 - the result 1 or 0 and of type int.
|
|
|
|
In this testcase, the int result is again converted to a floating-poing
|
|
or complex type.
|
|
|
|
While having a floating-point/complex array element with || and && can make
|
|
sense, having a non-integer/non-bool reduction variable is odd but valid.
|
|
|
|
Test: FP reduction variable + FP array - as reduction-1.c but with target */
|
|
|
|
#define N 1024
|
|
_Complex float rcf[N];
|
|
_Complex double rcd[N];
|
|
float rf[N];
|
|
double rd[N];
|
|
|
|
int
|
|
reduction_or ()
|
|
{
|
|
float orf = 0;
|
|
double ord = 0;
|
|
_Complex float orfc = 0;
|
|
_Complex double ordc = 0;
|
|
|
|
#pragma omp target parallel reduction(||: orf) map(orf)
|
|
for (int i=0; i < N; ++i)
|
|
orf = orf || rf[i];
|
|
|
|
#pragma omp target parallel for reduction(||: ord) map(ord)
|
|
for (int i=0; i < N; ++i)
|
|
ord = ord || rcd[i];
|
|
|
|
#pragma omp target parallel for simd reduction(||: orfc) map(orfc)
|
|
for (int i=0; i < N; ++i)
|
|
orfc = orfc || rcf[i];
|
|
|
|
#pragma omp target parallel loop reduction(||: ordc) map(ordc)
|
|
for (int i=0; i < N; ++i)
|
|
ordc = ordc || rcd[i];
|
|
|
|
return orf + ord + __real__ orfc + __real__ ordc;
|
|
}
|
|
|
|
int
|
|
reduction_or_teams ()
|
|
{
|
|
float orf = 0;
|
|
double ord = 0;
|
|
_Complex float orfc = 0;
|
|
_Complex double ordc = 0;
|
|
|
|
#pragma omp target teams distribute parallel for reduction(||: orf) map(orf)
|
|
for (int i=0; i < N; ++i)
|
|
orf = orf || rf[i];
|
|
|
|
#pragma omp target teams distribute parallel for simd reduction(||: ord) map(ord)
|
|
for (int i=0; i < N; ++i)
|
|
ord = ord || rcd[i];
|
|
|
|
#pragma omp target teams distribute parallel for reduction(||: orfc) map(orfc)
|
|
for (int i=0; i < N; ++i)
|
|
orfc = orfc || rcf[i];
|
|
|
|
#pragma omp target teams distribute parallel for simd reduction(||: ordc) map(ordc)
|
|
for (int i=0; i < N; ++i)
|
|
ordc = ordc || rcd[i];
|
|
|
|
return orf + ord + __real__ orfc + __real__ ordc;
|
|
}
|
|
|
|
int
|
|
reduction_and ()
|
|
{
|
|
float andf = 1;
|
|
double andd = 1;
|
|
_Complex float andfc = 1;
|
|
_Complex double anddc = 1;
|
|
|
|
#pragma omp target parallel reduction(&&: andf) map(andf)
|
|
for (int i=0; i < N; ++i)
|
|
andf = andf && rf[i];
|
|
|
|
#pragma omp target parallel for reduction(&&: andd) map(andd)
|
|
for (int i=0; i < N; ++i)
|
|
andd = andd && rcd[i];
|
|
|
|
#pragma omp target parallel for simd reduction(&&: andfc) map(andfc)
|
|
for (int i=0; i < N; ++i)
|
|
andfc = andfc && rcf[i];
|
|
|
|
#pragma omp target parallel loop reduction(&&: anddc) map(anddc)
|
|
for (int i=0; i < N; ++i)
|
|
anddc = anddc && rcd[i];
|
|
|
|
return andf + andd + __real__ andfc + __real__ anddc;
|
|
}
|
|
|
|
int
|
|
reduction_and_teams ()
|
|
{
|
|
float andf = 1;
|
|
double andd = 1;
|
|
_Complex float andfc = 1;
|
|
_Complex double anddc = 1;
|
|
|
|
#pragma omp target teams distribute parallel for reduction(&&: andf) map(andf)
|
|
for (int i=0; i < N; ++i)
|
|
andf = andf && rf[i];
|
|
|
|
#pragma omp target teams distribute parallel for simd reduction(&&: andd) map(andd)
|
|
for (int i=0; i < N; ++i)
|
|
andd = andd && rcd[i];
|
|
|
|
#pragma omp target teams distribute parallel for reduction(&&: andfc) map(andfc)
|
|
for (int i=0; i < N; ++i)
|
|
andfc = andfc && rcf[i];
|
|
|
|
#pragma omp target teams distribute parallel for simd reduction(&&: anddc) map(anddc)
|
|
for (int i=0; i < N; ++i)
|
|
anddc = anddc && rcd[i];
|
|
|
|
return andf + andd + __real__ andfc + __real__ anddc;
|
|
}
|
|
|
|
int
|
|
main ()
|
|
{
|
|
for (int i = 0; i < N; ++i)
|
|
{
|
|
rf[i] = 0;
|
|
rd[i] = 0;
|
|
rcf[i] = 0;
|
|
rcd[i] = 0;
|
|
}
|
|
|
|
if (reduction_or () != 0)
|
|
__builtin_abort ();
|
|
if (reduction_or_teams () != 0)
|
|
__builtin_abort ();
|
|
if (reduction_and () != 0)
|
|
__builtin_abort ();
|
|
if (reduction_and_teams () != 0)
|
|
__builtin_abort ();
|
|
|
|
rf[10] = 1.0;
|
|
rd[15] = 1.0;
|
|
rcf[10] = 1.0;
|
|
rcd[15] = 1.0i;
|
|
|
|
if (reduction_or () != 4)
|
|
__builtin_abort ();
|
|
if (reduction_or_teams () != 4)
|
|
__builtin_abort ();
|
|
if (reduction_and () != 0)
|
|
__builtin_abort ();
|
|
if (reduction_and_teams () != 0)
|
|
__builtin_abort ();
|
|
|
|
for (int i = 0; i < N; ++i)
|
|
{
|
|
rf[i] = 1;
|
|
rd[i] = 1;
|
|
rcf[i] = 1;
|
|
rcd[i] = 1;
|
|
}
|
|
|
|
if (reduction_or () != 4)
|
|
__builtin_abort ();
|
|
if (reduction_or_teams () != 4)
|
|
__builtin_abort ();
|
|
if (reduction_and () != 4)
|
|
__builtin_abort ();
|
|
if (reduction_and_teams () != 4)
|
|
__builtin_abort ();
|
|
|
|
rf[10] = 0.0;
|
|
rd[15] = 0.0;
|
|
rcf[10] = 0.0;
|
|
rcd[15] = 0.0;
|
|
|
|
if (reduction_or () != 4)
|
|
__builtin_abort ();
|
|
if (reduction_or_teams () != 4)
|
|
__builtin_abort ();
|
|
if (reduction_and () != 0)
|
|
__builtin_abort ();
|
|
if (reduction_and_teams () != 0)
|
|
__builtin_abort ();
|
|
|
|
return 0;
|
|
}
|