[OpenACC] Extract 'pass_oacc_loop_designation' out of 'pass_oacc_device_lower'
This really is a separate step -- and another pass to be added between the two, later on. gcc/ * omp-offload.c (oacc_loop_xform_head_tail, oacc_loop_process): 'update_stmt' after modification. (pass_oacc_loop_designation): New function, extracted out of... (pass_oacc_device_lower): ... this. (pass_data_oacc_loop_designation, pass_oacc_loop_designation) (make_pass_oacc_loop_designation): New * passes.def: Add it. * tree-parloops.c (create_parallel_loop): Adjust. * tree-pass.h (make_pass_oacc_loop_designation): New. gcc/testsuite/ * c-c++-common/goacc/classify-kernels-unparallelized.c: 's%oaccdevlow%oaccloops%g'. * c-c++-common/goacc/classify-kernels.c: Likewise. * c-c++-common/goacc/classify-parallel.c: Likewise. * c-c++-common/goacc/classify-routine-nohost.c: Likewise. * c-c++-common/goacc/classify-routine.c: Likewise. * c-c++-common/goacc/classify-serial.c: Likewise. * c-c++-common/goacc/routine-nohost-1.c: Likewise. * g++.dg/goacc/template.C: Likewise. * gcc.dg/goacc/loop-processing-1.c: Likewise. * gfortran.dg/goacc/classify-kernels-unparallelized.f95: Likewise. * gfortran.dg/goacc/classify-kernels.f95: Likewise. * gfortran.dg/goacc/classify-parallel.f95: Likewise. * gfortran.dg/goacc/classify-routine-nohost.f95: Likewise. * gfortran.dg/goacc/classify-routine.f95: Likewise. * gfortran.dg/goacc/classify-serial.f95: Likewise. * gfortran.dg/goacc/routine-multiple-directives-1.f90: Likewise. libgomp/ * testsuite/libgomp.oacc-c-c++-common/pr85486-2.c: 's%oaccdevlow%oaccloops%g'. * testsuite/libgomp.oacc-c-c++-common/pr85486-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/pr85486.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/routine-nohost-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/vector-length-128-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/vector-length-128-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/vector-length-128-4.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/vector-length-128-5.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/vector-length-128-6.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/vector-length-128-7.c: Likewise. * testsuite/libgomp.oacc-fortran/routine-nohost-1.f90: Likewise. Co-Authored-By: Julian Brown <julian@codesourcery.com> Co-Authored-By: Kwok Cheung Yeung <kcy@codesourcery.com>
This commit is contained in:
Thomas Schwinge
parent
f0529d96f5
commit
0829ab79d3
32 changed files with 214 additions and 158 deletions
|
|
@ -1367,6 +1367,7 @@ oacc_loop_xform_head_tail (gcall *from, int level)
|
|||
}
|
||||
else if (gimple_call_internal_p (stmt, IFN_GOACC_REDUCTION))
|
||||
*gimple_call_arg_ptr (stmt, 3) = replacement;
|
||||
update_stmt (stmt);
|
||||
|
||||
gsi_next (&gsi);
|
||||
while (gsi_end_p (gsi))
|
||||
|
|
@ -1392,25 +1393,28 @@ oacc_loop_process (oacc_loop *loop)
|
|||
gcall *call;
|
||||
|
||||
for (ix = 0; loop->ifns.iterate (ix, &call); ix++)
|
||||
switch (gimple_call_internal_fn (call))
|
||||
{
|
||||
case IFN_GOACC_LOOP:
|
||||
{
|
||||
switch (gimple_call_internal_fn (call))
|
||||
{
|
||||
bool is_e = gimple_call_arg (call, 5) == integer_minus_one_node;
|
||||
gimple_call_set_arg (call, 5, is_e ? e_mask_arg : mask_arg);
|
||||
if (!is_e)
|
||||
gimple_call_set_arg (call, 4, chunk_arg);
|
||||
case IFN_GOACC_LOOP:
|
||||
{
|
||||
bool is_e = gimple_call_arg (call, 5) == integer_minus_one_node;
|
||||
gimple_call_set_arg (call, 5, is_e ? e_mask_arg : mask_arg);
|
||||
if (!is_e)
|
||||
gimple_call_set_arg (call, 4, chunk_arg);
|
||||
}
|
||||
break;
|
||||
|
||||
case IFN_GOACC_TILE:
|
||||
gimple_call_set_arg (call, 3, mask_arg);
|
||||
gimple_call_set_arg (call, 4, e_mask_arg);
|
||||
break;
|
||||
|
||||
default:
|
||||
gcc_unreachable ();
|
||||
}
|
||||
break;
|
||||
|
||||
case IFN_GOACC_TILE:
|
||||
gimple_call_set_arg (call, 3, mask_arg);
|
||||
gimple_call_set_arg (call, 4, e_mask_arg);
|
||||
break;
|
||||
|
||||
default:
|
||||
gcc_unreachable ();
|
||||
}
|
||||
update_stmt (call);
|
||||
}
|
||||
|
||||
unsigned dim = GOMP_DIM_GANG;
|
||||
unsigned mask = loop->mask | loop->e_mask;
|
||||
|
|
@ -1912,7 +1916,7 @@ is_sync_builtin_call (gcall *call)
|
|||
point (including the host fallback). */
|
||||
|
||||
static unsigned int
|
||||
execute_oacc_device_lower ()
|
||||
execute_oacc_loop_designation ()
|
||||
{
|
||||
tree attrs = oacc_get_fn_attrib (current_function_decl);
|
||||
|
||||
|
|
@ -1981,6 +1985,8 @@ execute_oacc_device_lower ()
|
|||
gcc_unreachable ();
|
||||
}
|
||||
|
||||
/* This doesn't belong into 'pass_oacc_loop_designation' conceptually, but
|
||||
it's a convenient place, so... */
|
||||
if (is_oacc_routine)
|
||||
{
|
||||
tree attr = lookup_attribute ("omp declare target",
|
||||
|
|
@ -2088,9 +2094,23 @@ execute_oacc_device_lower ()
|
|||
free_oacc_loop (l);
|
||||
}
|
||||
|
||||
/* Offloaded targets may introduce new basic blocks, which require
|
||||
dominance information to update SSA. */
|
||||
calculate_dominance_info (CDI_DOMINATORS);
|
||||
free_oacc_loop (loops);
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
static unsigned int
|
||||
execute_oacc_device_lower ()
|
||||
{
|
||||
tree attrs = oacc_get_fn_attrib (current_function_decl);
|
||||
|
||||
if (!attrs)
|
||||
/* Not an offloaded function. */
|
||||
return 0;
|
||||
|
||||
int dims[GOMP_DIM_MAX];
|
||||
for (unsigned i = 0; i < GOMP_DIM_MAX; i++)
|
||||
dims[i] = oacc_get_fn_dim_size (current_function_decl, i);
|
||||
|
||||
hash_map<tree, tree> adjusted_vars;
|
||||
|
||||
|
|
@ -2355,8 +2375,6 @@ execute_oacc_device_lower ()
|
|||
}
|
||||
}
|
||||
|
||||
free_oacc_loop (loops);
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
|
|
@ -2397,6 +2415,36 @@ default_goacc_dim_limit (int ARG_UNUSED (axis))
|
|||
|
||||
namespace {
|
||||
|
||||
const pass_data pass_data_oacc_loop_designation =
|
||||
{
|
||||
GIMPLE_PASS, /* type */
|
||||
"oaccloops", /* name */
|
||||
OPTGROUP_OMP, /* optinfo_flags */
|
||||
TV_NONE, /* tv_id */
|
||||
PROP_cfg, /* properties_required */
|
||||
0 /* Possibly PROP_gimple_eomp. */, /* properties_provided */
|
||||
0, /* properties_destroyed */
|
||||
0, /* todo_flags_start */
|
||||
TODO_update_ssa | TODO_cleanup_cfg, /* todo_flags_finish */
|
||||
};
|
||||
|
||||
class pass_oacc_loop_designation : public gimple_opt_pass
|
||||
{
|
||||
public:
|
||||
pass_oacc_loop_designation (gcc::context *ctxt)
|
||||
: gimple_opt_pass (pass_data_oacc_loop_designation, ctxt)
|
||||
{}
|
||||
|
||||
/* opt_pass methods: */
|
||||
virtual bool gate (function *) { return flag_openacc; };
|
||||
|
||||
virtual unsigned int execute (function *)
|
||||
{
|
||||
return execute_oacc_loop_designation ();
|
||||
}
|
||||
|
||||
}; // class pass_oacc_loop_designation
|
||||
|
||||
const pass_data pass_data_oacc_device_lower =
|
||||
{
|
||||
GIMPLE_PASS, /* type */
|
||||
|
|
@ -2429,6 +2477,12 @@ public:
|
|||
|
||||
} // anon namespace
|
||||
|
||||
gimple_opt_pass *
|
||||
make_pass_oacc_loop_designation (gcc::context *ctxt)
|
||||
{
|
||||
return new pass_oacc_loop_designation (ctxt);
|
||||
}
|
||||
|
||||
gimple_opt_pass *
|
||||
make_pass_oacc_device_lower (gcc::context *ctxt)
|
||||
{
|
||||
|
|
|
|||
|
|
@ -183,6 +183,7 @@ along with GCC; see the file COPYING3. If not see
|
|||
INSERT_PASSES_AFTER (all_passes)
|
||||
NEXT_PASS (pass_fixup_cfg);
|
||||
NEXT_PASS (pass_lower_eh_dispatch);
|
||||
NEXT_PASS (pass_oacc_loop_designation);
|
||||
NEXT_PASS (pass_oacc_device_lower);
|
||||
NEXT_PASS (pass_omp_device_lower);
|
||||
NEXT_PASS (pass_omp_target_link);
|
||||
|
|
|
|||
|
|
@ -5,7 +5,7 @@
|
|||
{ dg-additional-options "-fopt-info-optimized-omp" }
|
||||
{ dg-additional-options "-fdump-tree-ompexp" }
|
||||
{ dg-additional-options "-fdump-tree-parloops1-all" }
|
||||
{ dg-additional-options "-fdump-tree-oaccdevlow" } */
|
||||
{ dg-additional-options "-fdump-tree-oaccloops" } */
|
||||
|
||||
/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
|
||||
aspects of that functionality. */
|
||||
|
|
@ -38,6 +38,6 @@ void KERNELS ()
|
|||
|
||||
/* Check the offloaded function's classification and compute dimensions (will
|
||||
always be 1 x 1 x 1 for non-offloading compilation).
|
||||
{ dg-final { scan-tree-dump-times "(?n)Function is unparallelized OpenACC kernels offload" 1 "oaccdevlow" } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccdevlow" } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc kernels, omp target entrypoint\\)\\)" 1 "oaccdevlow" } } */
|
||||
{ dg-final { scan-tree-dump-times "(?n)Function is unparallelized OpenACC kernels offload" 1 "oaccloops" } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccloops" } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc kernels, omp target entrypoint\\)\\)" 1 "oaccloops" } } */
|
||||
|
|
|
|||
|
|
@ -5,7 +5,7 @@
|
|||
{ dg-additional-options "-fopt-info-optimized-omp" }
|
||||
{ dg-additional-options "-fdump-tree-ompexp" }
|
||||
{ dg-additional-options "-fdump-tree-parloops1-all" }
|
||||
{ dg-additional-options "-fdump-tree-oaccdevlow" } */
|
||||
{ dg-additional-options "-fdump-tree-oaccloops" } */
|
||||
|
||||
/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
|
||||
aspects of that functionality. */
|
||||
|
|
@ -34,6 +34,6 @@ void KERNELS ()
|
|||
|
||||
/* Check the offloaded function's classification and compute dimensions (will
|
||||
always be 1 x 1 x 1 for non-offloading compilation).
|
||||
{ dg-final { scan-tree-dump-times "(?n)Function is parallelized OpenACC kernels offload" 1 "oaccdevlow" } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccdevlow" } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc kernels parallelized, oacc function \\(, , \\), oacc kernels, omp target entrypoint\\)\\)" 1 "oaccdevlow" } } */
|
||||
{ dg-final { scan-tree-dump-times "(?n)Function is parallelized OpenACC kernels offload" 1 "oaccloops" } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccloops" } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc kernels parallelized, oacc function \\(, , \\), oacc kernels, omp target entrypoint\\)\\)" 1 "oaccloops" } } */
|
||||
|
|
|
|||
|
|
@ -4,7 +4,7 @@
|
|||
/* { dg-additional-options "-O2" }
|
||||
{ dg-additional-options "-fopt-info-optimized-omp" }
|
||||
{ dg-additional-options "-fdump-tree-ompexp" }
|
||||
{ dg-additional-options "-fdump-tree-oaccdevlow" } */
|
||||
{ dg-additional-options "-fdump-tree-oaccloops" } */
|
||||
|
||||
/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
|
||||
aspects of that functionality. */
|
||||
|
|
@ -27,6 +27,6 @@ void PARALLEL ()
|
|||
|
||||
/* Check the offloaded function's classification and compute dimensions (will
|
||||
always be 1 x 1 x 1 for non-offloading compilation).
|
||||
{ dg-final { scan-tree-dump-times "(?n)Function is OpenACC parallel offload" 1 "oaccdevlow" } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccdevlow" } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc parallel, omp target entrypoint\\)\\)" 1 "oaccdevlow" } } */
|
||||
{ dg-final { scan-tree-dump-times "(?n)Function is OpenACC parallel offload" 1 "oaccloops" } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccloops" } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc parallel, omp target entrypoint\\)\\)" 1 "oaccloops" } } */
|
||||
|
|
|
|||
|
|
@ -4,7 +4,7 @@
|
|||
/* { dg-additional-options "-O2" }
|
||||
{ dg-additional-options "-fopt-info-optimized-omp" }
|
||||
{ dg-additional-options "-fdump-tree-ompexp" }
|
||||
{ dg-additional-options "-fdump-tree-oaccdevlow" } */
|
||||
{ dg-additional-options "-fdump-tree-oaccloops" } */
|
||||
|
||||
/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
|
||||
aspects of that functionality. */
|
||||
|
|
@ -28,14 +28,14 @@ void ROUTINE ()
|
|||
{ dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(omp declare target \\(nohost worker\\), oacc function \\(0 1, 1 0, 1 0\\)\\)\\)" 1 "ompexp" } } */
|
||||
|
||||
/* Check the offloaded function's classification.
|
||||
{ dg-final { scan-tree-dump-times "(?n)Function is OpenACC routine level 1" 1 "oaccdevlow" } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE' has 'nohost' clause" 1 "oaccdevlow" { target c } } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)OpenACC routine 'void ROUTINE\\(\\)' has 'nohost' clause" 1 "oaccdevlow" { target { c++ && { ! offloading_enabled } } } } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE\\(\\)' has 'nohost' clause" 1 "oaccdevlow" { target { c++ && offloading_enabled } } } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE' discarded" 1 "oaccdevlow" { target c } } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)OpenACC routine 'void ROUTINE\\(\\)' discarded" 1 "oaccdevlow" { target { c++ && { ! offloading_enabled } } } } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE\\(\\)' discarded" 1 "oaccdevlow" { target { c++ && offloading_enabled } } } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)Function is OpenACC routine level 1" 1 "oaccloops" } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE' has 'nohost' clause" 1 "oaccloops" { target c } } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)OpenACC routine 'void ROUTINE\\(\\)' has 'nohost' clause" 1 "oaccloops" { target { c++ && { ! offloading_enabled } } } } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE\\(\\)' has 'nohost' clause" 1 "oaccloops" { target { c++ && offloading_enabled } } } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE' discarded" 1 "oaccloops" { target c } } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)OpenACC routine 'void ROUTINE\\(\\)' discarded" 1 "oaccloops" { target { c++ && { ! offloading_enabled } } } } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE\\(\\)' discarded" 1 "oaccloops" { target { c++ && offloading_enabled } } } }
|
||||
TODO See PR101551 for 'offloading_enabled' differences.
|
||||
{ dg-final { scan-tree-dump-not "(?n)Compute dimensions" "oaccdevlow" } }
|
||||
{ dg-final { scan-tree-dump-not "(?n)__attribute__\\(.*omp declare target \\(nohost" "oaccdevlow" } }
|
||||
{ dg-final { scan-tree-dump-not "(?n)void ROUTINE \\(\\)" "oaccdevlow" } } */
|
||||
{ dg-final { scan-tree-dump-not "(?n)Compute dimensions" "oaccloops" } }
|
||||
{ dg-final { scan-tree-dump-not "(?n)__attribute__\\(.*omp declare target \\(nohost" "oaccloops" } }
|
||||
{ dg-final { scan-tree-dump-not "(?n)void ROUTINE \\(\\)" "oaccloops" } } */
|
||||
|
|
|
|||
|
|
@ -4,7 +4,7 @@
|
|||
/* { dg-additional-options "-O2" }
|
||||
{ dg-additional-options "-fopt-info-optimized-omp" }
|
||||
{ dg-additional-options "-fdump-tree-ompexp" }
|
||||
{ dg-additional-options "-fdump-tree-oaccdevlow" } */
|
||||
{ dg-additional-options "-fdump-tree-oaccloops" } */
|
||||
|
||||
/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
|
||||
aspects of that functionality. */
|
||||
|
|
@ -29,14 +29,14 @@ void ROUTINE ()
|
|||
|
||||
/* Check the offloaded function's classification and compute dimensions (will
|
||||
always be 1 x 1 x 1 for non-offloading compilation).
|
||||
{ dg-final { scan-tree-dump-times "(?n)Function is OpenACC routine level 1" 1 "oaccdevlow" } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE' doesn't have 'nohost' clause" 1 "oaccdevlow" { target c } } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)OpenACC routine 'void ROUTINE\\(\\)' doesn't have 'nohost' clause" 1 "oaccdevlow" { target { c++ && { ! offloading_enabled } } } } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE\\(\\)' doesn't have 'nohost' clause" 1 "oaccdevlow" { target { c++ && offloading_enabled } } } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE' not discarded" 1 "oaccdevlow" { target c } } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)OpenACC routine 'void ROUTINE\\(\\)' not discarded" 1 "oaccdevlow" { target { c++ && { ! offloading_enabled } } } } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE\\(\\)' not discarded" 1 "oaccdevlow" { target { c++ && offloading_enabled } } } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)Function is OpenACC routine level 1" 1 "oaccloops" } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE' doesn't have 'nohost' clause" 1 "oaccloops" { target c } } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)OpenACC routine 'void ROUTINE\\(\\)' doesn't have 'nohost' clause" 1 "oaccloops" { target { c++ && { ! offloading_enabled } } } } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE\\(\\)' doesn't have 'nohost' clause" 1 "oaccloops" { target { c++ && offloading_enabled } } } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE' not discarded" 1 "oaccloops" { target c } } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)OpenACC routine 'void ROUTINE\\(\\)' not discarded" 1 "oaccloops" { target { c++ && { ! offloading_enabled } } } } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE\\(\\)' not discarded" 1 "oaccloops" { target { c++ && offloading_enabled } } } }
|
||||
TODO See PR101551 for 'offloading_enabled' differences.
|
||||
{ dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccdevlow" } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(0 1, 1 1, 1 1\\), omp declare target \\(worker\\), oacc function \\(0 1, 1 0, 1 0\\)\\)\\)" 1 "oaccdevlow" } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)void ROUTINE \\(\\)" 1 "oaccdevlow" } } */
|
||||
{ dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccloops" } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(0 1, 1 1, 1 1\\), omp declare target \\(worker\\), oacc function \\(0 1, 1 0, 1 0\\)\\)\\)" 1 "oaccloops" } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)void ROUTINE \\(\\)" 1 "oaccloops" } } */
|
||||
|
|
|
|||
|
|
@ -4,7 +4,7 @@
|
|||
/* { dg-additional-options "-O2" }
|
||||
{ dg-additional-options "-fopt-info-optimized-omp" }
|
||||
{ dg-additional-options "-fdump-tree-ompexp" }
|
||||
{ dg-additional-options "-fdump-tree-oaccdevlow" } */
|
||||
{ dg-additional-options "-fdump-tree-oaccloops" } */
|
||||
|
||||
/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
|
||||
aspects of that functionality. */
|
||||
|
|
@ -32,6 +32,6 @@ void SERIAL ()
|
|||
|
||||
/* Check the offloaded function's classification and compute dimensions (will
|
||||
always be 1 x 1 x 1 for non-offloading compilation).
|
||||
{ dg-final { scan-tree-dump-times "(?n)Function is OpenACC serial offload" 1 "oaccdevlow" } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccdevlow" } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc serial, omp target entrypoint\\)\\)" 1 "oaccdevlow" } } */
|
||||
{ dg-final { scan-tree-dump-times "(?n)Function is OpenACC serial offload" 1 "oaccloops" } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccloops" } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc serial, omp target entrypoint\\)\\)" 1 "oaccloops" } } */
|
||||
|
|
|
|||
|
|
@ -1,6 +1,6 @@
|
|||
/* Test OpenACC 'routine' with 'nohost' clause, valid use. */
|
||||
|
||||
/* { dg-additional-options "-fdump-tree-oaccdevlow" } */
|
||||
/* { dg-additional-options "-fdump-tree-oaccloops" } */
|
||||
|
||||
#pragma acc routine nohost
|
||||
int THREE(void)
|
||||
|
|
@ -13,7 +13,7 @@ int THREE(void)
|
|||
#pragma acc routine nohost
|
||||
extern int THREE(void);
|
||||
|
||||
/* { dg-final { scan-tree-dump-times {(?n)^OpenACC routine '[^']*THREE[^']*' has 'nohost' clause\.$} 1 oaccdevlow } } */
|
||||
/* { dg-final { scan-tree-dump-times {(?n)^OpenACC routine '[^']*THREE[^']*' has 'nohost' clause\.$} 1 oaccloops } } */
|
||||
|
||||
|
||||
#pragma acc routine nohost
|
||||
|
|
@ -30,7 +30,7 @@ extern void NOTHING(void);
|
|||
|
||||
#pragma acc routine (NOTHING) nohost
|
||||
|
||||
/* { dg-final { scan-tree-dump-times {(?n)^OpenACC routine '[^']*NOTHING[^']*' has 'nohost' clause\.$} 1 oaccdevlow } } */
|
||||
/* { dg-final { scan-tree-dump-times {(?n)^OpenACC routine '[^']*NOTHING[^']*' has 'nohost' clause\.$} 1 oaccloops } } */
|
||||
|
||||
|
||||
extern float ADD(float, float);
|
||||
|
|
@ -47,4 +47,4 @@ extern float ADD(float, float);
|
|||
|
||||
#pragma acc routine (ADD) nohost
|
||||
|
||||
/* { dg-final { scan-tree-dump-times {(?n)^OpenACC routine '[^']*ADD[^']*' has 'nohost' clause\.$} 1 oaccdevlow } } */
|
||||
/* { dg-final { scan-tree-dump-times {(?n)^OpenACC routine '[^']*ADD[^']*' has 'nohost' clause\.$} 1 oaccloops } } */
|
||||
|
|
|
|||
|
|
@ -1,4 +1,4 @@
|
|||
/* { dg-additional-options "-fdump-tree-oaccdevlow" } */
|
||||
/* { dg-additional-options "-fdump-tree-oaccloops" } */
|
||||
|
||||
#pragma acc routine nohost
|
||||
template <typename T> T
|
||||
|
|
@ -156,13 +156,13 @@ main ()
|
|||
return b + c;
|
||||
}
|
||||
|
||||
/* { dg-final { scan-tree-dump-times {(?n)^OpenACC routine '[^']+' has 'nohost' clause\.$} 4 oaccdevlow } }
|
||||
{ dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'T accDouble\(int\) \[with T = char\]' has 'nohost' clause\.$} 1 oaccdevlow { target { ! offloading_enabled } } } }
|
||||
{ dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'accDouble<char>\(int\)char' has 'nohost' clause\.$} 1 oaccdevlow { target offloading_enabled } } }
|
||||
{ dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'T accDouble\(int\) \[with T = int\]' has 'nohost' clause\.$} 1 oaccdevlow { target { ! offloading_enabled } } } }
|
||||
{ dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'accDouble<int>\(int\)int' has 'nohost' clause\.$} 1 oaccdevlow { target offloading_enabled } } }
|
||||
{ dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'T accDouble\(int\) \[with T = float\]' has 'nohost' clause\.$} 1 oaccdevlow { target { ! offloading_enabled } } } }
|
||||
{ dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'accDouble<float>\(int\)float' has 'nohost' clause\.$} 1 oaccdevlow { target offloading_enabled } } }
|
||||
{ dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'T accDouble\(int\) \[with T = double\]' has 'nohost' clause\.$} 1 oaccdevlow { target { ! offloading_enabled } } } }
|
||||
{ dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'accDouble<double>\(int\)double' has 'nohost' clause\.$} 1 oaccdevlow { target offloading_enabled } } }
|
||||
/* { dg-final { scan-tree-dump-times {(?n)^OpenACC routine '[^']+' has 'nohost' clause\.$} 4 oaccloops } }
|
||||
{ dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'T accDouble\(int\) \[with T = char\]' has 'nohost' clause\.$} 1 oaccloops { target { ! offloading_enabled } } } }
|
||||
{ dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'accDouble<char>\(int\)char' has 'nohost' clause\.$} 1 oaccloops { target offloading_enabled } } }
|
||||
{ dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'T accDouble\(int\) \[with T = int\]' has 'nohost' clause\.$} 1 oaccloops { target { ! offloading_enabled } } } }
|
||||
{ dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'accDouble<int>\(int\)int' has 'nohost' clause\.$} 1 oaccloops { target offloading_enabled } } }
|
||||
{ dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'T accDouble\(int\) \[with T = float\]' has 'nohost' clause\.$} 1 oaccloops { target { ! offloading_enabled } } } }
|
||||
{ dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'accDouble<float>\(int\)float' has 'nohost' clause\.$} 1 oaccloops { target offloading_enabled } } }
|
||||
{ dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'T accDouble\(int\) \[with T = double\]' has 'nohost' clause\.$} 1 oaccloops { target { ! offloading_enabled } } } }
|
||||
{ dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'accDouble<double>\(int\)double' has 'nohost' clause\.$} 1 oaccloops { target offloading_enabled } } }
|
||||
TODO See PR101551 for 'offloading_enabled' differences. */
|
||||
|
|
|
|||
|
|
@ -1,5 +1,5 @@
|
|||
/* Make sure that OpenACC loop processing happens. */
|
||||
/* { dg-additional-options "-O2 -fdump-tree-oaccdevlow" } */
|
||||
/* { dg-additional-options "-O2 -fdump-tree-oaccloops" } */
|
||||
|
||||
extern int place ();
|
||||
|
||||
|
|
@ -15,4 +15,4 @@ void vector_1 (int *ary, int size)
|
|||
}
|
||||
}
|
||||
|
||||
/* { dg-final { scan-tree-dump {OpenACC loops.*Loop 0\(0\).*Loop 24\(1\).*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_HEAD_MARK, 0, 1, 36\);.*Head-0:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_HEAD_MARK, 0, 1, 36\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_FORK, \.data_dep\.[0-9_]+, 0\);.*Tail-0:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_TAIL_MARK, \.data_dep\.[0-9_]+, 1\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_JOIN, \.data_dep\.[0-9_]+, 0\);.*Loop 6\(6\).*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_HEAD_MARK, 0, 2, 6\);.*Head-0:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_HEAD_MARK, 0, 2, 6\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_FORK, \.data_dep\.[0-9_]+, 1\);.*Head-1:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_HEAD_MARK, \.data_dep\.[0-9_]+, 1\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_FORK, \.data_dep\.[0-9_]+, 2\);.*Tail-1:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_TAIL_MARK, \.data_dep\.[0-9_]+, 2\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_JOIN, \.data_dep\.[0-9_]+, 2\);.*Tail-0:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_TAIL_MARK, \.data_dep\.[0-9_]+, 1\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_JOIN, \.data_dep\.[0-9_]+, 1\);} "oaccdevlow" } } */
|
||||
/* { dg-final { scan-tree-dump {OpenACC loops.*Loop 0\(0\).*Loop 24\(1\).*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_HEAD_MARK, 0, 1, 36\);.*Head-0:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_HEAD_MARK, 0, 1, 36\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_FORK, \.data_dep\.[0-9_]+, 0\);.*Tail-0:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_TAIL_MARK, \.data_dep\.[0-9_]+, 1\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_JOIN, \.data_dep\.[0-9_]+, 0\);.*Loop 6\(6\).*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_HEAD_MARK, 0, 2, 6\);.*Head-0:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_HEAD_MARK, 0, 2, 6\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_FORK, \.data_dep\.[0-9_]+, 1\);.*Head-1:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_HEAD_MARK, \.data_dep\.[0-9_]+, 1\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_FORK, \.data_dep\.[0-9_]+, 2\);.*Tail-1:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_TAIL_MARK, \.data_dep\.[0-9_]+, 2\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_JOIN, \.data_dep\.[0-9_]+, 2\);.*Tail-0:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_TAIL_MARK, \.data_dep\.[0-9_]+, 1\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_JOIN, \.data_dep\.[0-9_]+, 1\);} "oaccloops" } } */
|
||||
|
|
|
|||
|
|
@ -5,7 +5,7 @@
|
|||
! { dg-additional-options "-fopt-info-optimized-omp" }
|
||||
! { dg-additional-options "-fdump-tree-ompexp" }
|
||||
! { dg-additional-options "-fdump-tree-parloops1-all" }
|
||||
! { dg-additional-options "-fdump-tree-oaccdevlow" }
|
||||
! { dg-additional-options "-fdump-tree-oaccloops" }
|
||||
|
||||
! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
|
||||
! aspects of that functionality.
|
||||
|
|
@ -40,6 +40,6 @@ end program main
|
|||
|
||||
! Check the offloaded function's classification and compute dimensions (will
|
||||
! always be 1 x 1 x 1 for non-offloading compilation).
|
||||
! { dg-final { scan-tree-dump-times "(?n)Function is unparallelized OpenACC kernels offload" 1 "oaccdevlow" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccdevlow" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc kernels, omp target entrypoint\\)\\)" 1 "oaccdevlow" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)Function is unparallelized OpenACC kernels offload" 1 "oaccloops" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccloops" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc kernels, omp target entrypoint\\)\\)" 1 "oaccloops" } }
|
||||
|
|
|
|||
|
|
@ -5,7 +5,7 @@
|
|||
! { dg-additional-options "-fopt-info-optimized-omp" }
|
||||
! { dg-additional-options "-fdump-tree-ompexp" }
|
||||
! { dg-additional-options "-fdump-tree-parloops1-all" }
|
||||
! { dg-additional-options "-fdump-tree-oaccdevlow" }
|
||||
! { dg-additional-options "-fdump-tree-oaccloops" }
|
||||
|
||||
! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
|
||||
! aspects of that functionality.
|
||||
|
|
@ -36,6 +36,6 @@ end program main
|
|||
|
||||
! Check the offloaded function's classification and compute dimensions (will
|
||||
! always be 1 x 1 x 1 for non-offloading compilation).
|
||||
! { dg-final { scan-tree-dump-times "(?n)Function is parallelized OpenACC kernels offload" 1 "oaccdevlow" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccdevlow" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc kernels parallelized, oacc function \\(, , \\), oacc kernels, omp target entrypoint\\)\\)" 1 "oaccdevlow" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)Function is parallelized OpenACC kernels offload" 1 "oaccloops" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccloops" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc kernels parallelized, oacc function \\(, , \\), oacc kernels, omp target entrypoint\\)\\)" 1 "oaccloops" } }
|
||||
|
|
|
|||
|
|
@ -4,7 +4,7 @@
|
|||
! { dg-additional-options "-O2" }
|
||||
! { dg-additional-options "-fopt-info-optimized-omp" }
|
||||
! { dg-additional-options "-fdump-tree-ompexp" }
|
||||
! { dg-additional-options "-fdump-tree-oaccdevlow" }
|
||||
! { dg-additional-options "-fdump-tree-oaccloops" }
|
||||
|
||||
! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
|
||||
! aspects of that functionality.
|
||||
|
|
@ -29,6 +29,6 @@ end program main
|
|||
|
||||
! Check the offloaded function's classification and compute dimensions (will
|
||||
! always be 1 x 1 x 1 for non-offloading compilation).
|
||||
! { dg-final { scan-tree-dump-times "(?n)Function is OpenACC parallel offload" 1 "oaccdevlow" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccdevlow" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc parallel, omp target entrypoint\\)\\)" 1 "oaccdevlow" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)Function is OpenACC parallel offload" 1 "oaccloops" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccloops" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc parallel, omp target entrypoint\\)\\)" 1 "oaccloops" } }
|
||||
|
|
|
|||
|
|
@ -4,7 +4,7 @@
|
|||
! { dg-additional-options "-O2" }
|
||||
! { dg-additional-options "-fopt-info-optimized-omp" }
|
||||
! { dg-additional-options "-fdump-tree-ompexp" }
|
||||
! { dg-additional-options "-fdump-tree-oaccdevlow" }
|
||||
! { dg-additional-options "-fdump-tree-oaccloops" }
|
||||
|
||||
! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
|
||||
! aspects of that functionality.
|
||||
|
|
@ -27,13 +27,13 @@ end subroutine ROUTINE
|
|||
! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(0 1, 1 0, 1 0\\), omp declare target \\(nohost worker\\)\\)\\)" 1 "ompexp" } }
|
||||
|
||||
! Check the offloaded function's classification.
|
||||
! { dg-final { scan-tree-dump-times "(?n)Function is OpenACC routine level 1" 1 "oaccdevlow" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'routine' has 'nohost' clause" 1 "oaccdevlow" { target { ! offloading_enabled } } } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'routine_' has 'nohost' clause" 1 "oaccdevlow" { target offloading_enabled } } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'routine' discarded" 1 "oaccdevlow" { target { ! offloading_enabled } } } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'routine_' discarded" 1 "oaccdevlow" { target offloading_enabled } } }
|
||||
! { dg-final { scan-tree-dump-not "(?n)Compute dimensions" "oaccdevlow" } }
|
||||
! { dg-final { scan-tree-dump-not "(?n)__attribute__\\(.*omp declare target \\(nohost" "oaccdevlow" } }
|
||||
! { dg-final { scan-tree-dump-not "(?n)void routine \\(\\)" "oaccdevlow" { target { ! offloading_enabled } } } }
|
||||
! { dg-final { scan-tree-dump-not "(?n)void routine_ \\(\\)" "oaccdevlow" { target offloading_enabled } } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)Function is OpenACC routine level 1" 1 "oaccloops" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'routine' has 'nohost' clause" 1 "oaccloops" { target { ! offloading_enabled } } } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'routine_' has 'nohost' clause" 1 "oaccloops" { target offloading_enabled } } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'routine' discarded" 1 "oaccloops" { target { ! offloading_enabled } } } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'routine_' discarded" 1 "oaccloops" { target offloading_enabled } } }
|
||||
! { dg-final { scan-tree-dump-not "(?n)Compute dimensions" "oaccloops" } }
|
||||
! { dg-final { scan-tree-dump-not "(?n)__attribute__\\(.*omp declare target \\(nohost" "oaccloops" } }
|
||||
! { dg-final { scan-tree-dump-not "(?n)void routine \\(\\)" "oaccloops" { target { ! offloading_enabled } } } }
|
||||
! { dg-final { scan-tree-dump-not "(?n)void routine_ \\(\\)" "oaccloops" { target offloading_enabled } } }
|
||||
!TODO See PR101551 for 'offloading_enabled' differences.
|
||||
|
|
|
|||
|
|
@ -4,7 +4,7 @@
|
|||
! { dg-additional-options "-O2" }
|
||||
! { dg-additional-options "-fopt-info-optimized-omp" }
|
||||
! { dg-additional-options "-fdump-tree-ompexp" }
|
||||
! { dg-additional-options "-fdump-tree-oaccdevlow" }
|
||||
! { dg-additional-options "-fdump-tree-oaccloops" }
|
||||
|
||||
! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
|
||||
! aspects of that functionality.
|
||||
|
|
@ -28,13 +28,13 @@ end subroutine ROUTINE
|
|||
|
||||
! Check the offloaded function's classification and compute dimensions (will
|
||||
! always be 1 x 1 x 1 for non-offloading compilation).
|
||||
! { dg-final { scan-tree-dump-times "(?n)Function is OpenACC routine level 1" 1 "oaccdevlow" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'routine' doesn't have 'nohost' clause" 1 "oaccdevlow" { target { ! offloading_enabled } } } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'routine_' doesn't have 'nohost' clause" 1 "oaccdevlow" { target offloading_enabled } } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'routine' not discarded" 1 "oaccdevlow" { target { ! offloading_enabled } } } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'routine_' not discarded" 1 "oaccdevlow" { target offloading_enabled } } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccdevlow" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(0 1, 1 1, 1 1\\), omp declare target \\(worker\\)\\)\\)" 1 "oaccdevlow" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)void routine \\(\\)" 1 "oaccdevlow" { target { ! offloading_enabled } } } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)void routine_ \\(\\)" 1 "oaccdevlow" { target offloading_enabled } } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)Function is OpenACC routine level 1" 1 "oaccloops" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'routine' doesn't have 'nohost' clause" 1 "oaccloops" { target { ! offloading_enabled } } } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'routine_' doesn't have 'nohost' clause" 1 "oaccloops" { target offloading_enabled } } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'routine' not discarded" 1 "oaccloops" { target { ! offloading_enabled } } } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'routine_' not discarded" 1 "oaccloops" { target offloading_enabled } } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccloops" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(0 1, 1 1, 1 1\\), omp declare target \\(worker\\)\\)\\)" 1 "oaccloops" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)void routine \\(\\)" 1 "oaccloops" { target { ! offloading_enabled } } } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)void routine_ \\(\\)" 1 "oaccloops" { target offloading_enabled } } }
|
||||
!TODO See PR101551 for 'offloading_enabled' differences.
|
||||
|
|
|
|||
|
|
@ -4,7 +4,7 @@
|
|||
! { dg-additional-options "-O2" }
|
||||
! { dg-additional-options "-fopt-info-optimized-omp" }
|
||||
! { dg-additional-options "-fdump-tree-ompexp" }
|
||||
! { dg-additional-options "-fdump-tree-oaccdevlow" }
|
||||
! { dg-additional-options "-fdump-tree-oaccloops" }
|
||||
|
||||
! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
|
||||
! aspects of that functionality.
|
||||
|
|
@ -32,6 +32,6 @@ end program main
|
|||
|
||||
! Check the offloaded function's classification and compute dimensions (will
|
||||
! always be 1 x 1 x 1 for non-offloading compilation).
|
||||
! { dg-final { scan-tree-dump-times "(?n)Function is OpenACC serial offload" 1 "oaccdevlow" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccdevlow" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc serial, omp target entrypoint\\)\\)" 1 "oaccdevlow" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)Function is OpenACC serial offload" 1 "oaccloops" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccloops" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc serial, omp target entrypoint\\)\\)" 1 "oaccloops" } }
|
||||
|
|
|
|||
|
|
@ -1,6 +1,6 @@
|
|||
! Check for valid cases of multiple OpenACC 'routine' directives.
|
||||
|
||||
! { dg-additional-options "-fdump-tree-oaccdevlow" }
|
||||
! { dg-additional-options "-fdump-tree-oaccloops" }
|
||||
!TODO See PR101551 for 'offloading_enabled' differences.
|
||||
|
||||
! { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
|
||||
|
|
@ -11,32 +11,32 @@
|
|||
!$ACC ROUTINE(s_1) SEQ
|
||||
!$ACC ROUTINE SEQ
|
||||
END SUBROUTINE s_1
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 's_1' doesn't have 'nohost' clause" 1 "oaccdevlow" { target { ! offloading_enabled } } } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 's_1_' doesn't have 'nohost' clause" 1 "oaccdevlow" { target offloading_enabled } } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 's_1' doesn't have 'nohost' clause" 1 "oaccloops" { target { ! offloading_enabled } } } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 's_1_' doesn't have 'nohost' clause" 1 "oaccloops" { target offloading_enabled } } }
|
||||
|
||||
SUBROUTINE s_1_nh
|
||||
!$ACC ROUTINE(s_1_nh) NOHOST
|
||||
!$ACC ROUTINE(s_1_nh) SEQ NOHOST
|
||||
!$ACC ROUTINE NOHOST SEQ
|
||||
END SUBROUTINE s_1_nh
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 's_1_nh' has 'nohost' clause" 1 "oaccdevlow" { target { ! offloading_enabled } } } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 's_1_nh_' has 'nohost' clause" 1 "oaccdevlow" { target offloading_enabled } } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 's_1_nh' has 'nohost' clause" 1 "oaccloops" { target { ! offloading_enabled } } } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 's_1_nh_' has 'nohost' clause" 1 "oaccloops" { target offloading_enabled } } }
|
||||
|
||||
SUBROUTINE s_2
|
||||
!$ACC ROUTINE
|
||||
!$ACC ROUTINE SEQ
|
||||
!$ACC ROUTINE(s_2)
|
||||
END SUBROUTINE s_2
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 's_2' doesn't have 'nohost' clause" 1 "oaccdevlow" { target { ! offloading_enabled } } } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 's_2_' doesn't have 'nohost' clause" 1 "oaccdevlow" { target offloading_enabled } } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 's_2' doesn't have 'nohost' clause" 1 "oaccloops" { target { ! offloading_enabled } } } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 's_2_' doesn't have 'nohost' clause" 1 "oaccloops" { target offloading_enabled } } }
|
||||
|
||||
SUBROUTINE s_2_nh
|
||||
!$ACC ROUTINE NOHOST
|
||||
!$ACC ROUTINE NOHOST SEQ
|
||||
!$ACC ROUTINE(s_2_nh) NOHOST
|
||||
END SUBROUTINE s_2_nh
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 's_2_nh' has 'nohost' clause" 1 "oaccdevlow" { target { ! offloading_enabled } } } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 's_2_nh_' has 'nohost' clause" 1 "oaccdevlow" { target offloading_enabled } } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 's_2_nh' has 'nohost' clause" 1 "oaccloops" { target { ! offloading_enabled } } } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 's_2_nh_' has 'nohost' clause" 1 "oaccloops" { target offloading_enabled } } }
|
||||
|
||||
SUBROUTINE v_1
|
||||
!$ACC ROUTINE VECTOR
|
||||
|
|
@ -45,8 +45,8 @@
|
|||
!$ACC ROUTINE VECTOR
|
||||
! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-5 }
|
||||
END SUBROUTINE v_1
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'v_1' doesn't have 'nohost' clause" 1 "oaccdevlow" { target { ! offloading_enabled } } } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'v_1_' doesn't have 'nohost' clause" 1 "oaccdevlow" { target offloading_enabled } } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'v_1' doesn't have 'nohost' clause" 1 "oaccloops" { target { ! offloading_enabled } } } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'v_1_' doesn't have 'nohost' clause" 1 "oaccloops" { target offloading_enabled } } }
|
||||
|
||||
SUBROUTINE v_1_nh
|
||||
!$ACC ROUTINE NOHOST VECTOR
|
||||
|
|
@ -55,8 +55,8 @@
|
|||
!$ACC ROUTINE VECTOR NOHOST
|
||||
! { dg-bogus "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-5 }
|
||||
END SUBROUTINE v_1_nh
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'v_1_nh' has 'nohost' clause" 1 "oaccdevlow" { target { ! offloading_enabled } } } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'v_1_nh_' has 'nohost' clause" 1 "oaccdevlow" { target offloading_enabled } } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'v_1_nh' has 'nohost' clause" 1 "oaccloops" { target { ! offloading_enabled } } } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'v_1_nh_' has 'nohost' clause" 1 "oaccloops" { target offloading_enabled } } }
|
||||
|
||||
SUBROUTINE v_2
|
||||
!$ACC ROUTINE(v_2) VECTOR
|
||||
|
|
@ -64,8 +64,8 @@
|
|||
!$ACC ROUTINE(v_2) VECTOR
|
||||
! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-4 }
|
||||
END SUBROUTINE v_2
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'v_2' doesn't have 'nohost' clause" 1 "oaccdevlow" { target { ! offloading_enabled } } } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'v_2_' doesn't have 'nohost' clause" 1 "oaccdevlow" { target offloading_enabled } } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'v_2' doesn't have 'nohost' clause" 1 "oaccloops" { target { ! offloading_enabled } } } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'v_2_' doesn't have 'nohost' clause" 1 "oaccloops" { target offloading_enabled } } }
|
||||
|
||||
SUBROUTINE v_2_nh
|
||||
!$ACC ROUTINE(v_2_nh) VECTOR NOHOST
|
||||
|
|
@ -73,8 +73,8 @@
|
|||
!$ACC ROUTINE(v_2_nh) NOHOST VECTOR
|
||||
! { dg-bogus "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-4 }
|
||||
END SUBROUTINE v_2_nh
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'v_2_nh' has 'nohost' clause" 1 "oaccdevlow" { target { ! offloading_enabled } } } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'v_2_nh_' has 'nohost' clause" 1 "oaccdevlow" { target offloading_enabled } } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'v_2_nh' has 'nohost' clause" 1 "oaccloops" { target { ! offloading_enabled } } } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'v_2_nh_' has 'nohost' clause" 1 "oaccloops" { target offloading_enabled } } }
|
||||
|
||||
SUBROUTINE sub_1
|
||||
IMPLICIT NONE
|
||||
|
|
|
|||
|
|
@ -2867,7 +2867,7 @@ create_parallel_loop (class loop *loop, tree loop_fn, tree data,
|
|||
/* Emit GIMPLE_OMP_FOR. */
|
||||
if (oacc_kernels_p)
|
||||
/* Parallelized OpenACC kernels constructs use gang parallelism. See also
|
||||
omp-offload.c:execute_oacc_device_lower. */
|
||||
omp-offload.c:execute_oacc_loop_designation. */
|
||||
t = build_omp_clause (loc, OMP_CLAUSE_GANG);
|
||||
else
|
||||
{
|
||||
|
|
|
|||
|
|
@ -424,6 +424,7 @@ extern gimple_opt_pass *make_pass_diagnose_omp_blocks (gcc::context *ctxt);
|
|||
extern gimple_opt_pass *make_pass_expand_omp (gcc::context *ctxt);
|
||||
extern gimple_opt_pass *make_pass_expand_omp_ssa (gcc::context *ctxt);
|
||||
extern gimple_opt_pass *make_pass_omp_target_link (gcc::context *ctxt);
|
||||
extern gimple_opt_pass *make_pass_oacc_loop_designation (gcc::context *ctxt);
|
||||
extern gimple_opt_pass *make_pass_oacc_device_lower (gcc::context *ctxt);
|
||||
extern gimple_opt_pass *make_pass_omp_device_lower (gcc::context *ctxt);
|
||||
extern gimple_opt_pass *make_pass_object_sizes (gcc::context *ctxt);
|
||||
|
|
|
|||
|
|
@ -2,10 +2,10 @@
|
|||
/* { dg-additional-options "-DVECTOR_LENGTH=" } */
|
||||
/* { dg-additional-options "-fopenacc-dim=::128" } */
|
||||
|
||||
/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */
|
||||
/* { dg-additional-options "-foffload=-fdump-tree-oaccloops" } */
|
||||
/* { dg-set-target-env-var "GOMP_DEBUG" "1" } */
|
||||
|
||||
#include "pr85486.c"
|
||||
|
||||
/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 32\\)" "oaccdevlow" } } */
|
||||
/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 32\\)" "oaccloops" } } */
|
||||
/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=1, vectors=32" } */
|
||||
|
|
|
|||
|
|
@ -2,10 +2,10 @@
|
|||
/* { dg-additional-options "-DVECTOR_LENGTH=" } */
|
||||
/* { dg-set-target-env-var "GOMP_OPENACC_DIM" "::128" } */
|
||||
|
||||
/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */
|
||||
/* { dg-additional-options "-foffload=-fdump-tree-oaccloops" } */
|
||||
/* { dg-set-target-env-var "GOMP_DEBUG" "1" } */
|
||||
|
||||
#include "pr85486.c"
|
||||
|
||||
/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 32\\)" "oaccdevlow" } } */
|
||||
/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 32\\)" "oaccloops" } } */
|
||||
/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=1, vectors=32" } */
|
||||
|
|
|
|||
|
|
@ -1,7 +1,7 @@
|
|||
/* { dg-do run { target openacc_nvidia_accel_selected } } */
|
||||
/* { dg-additional-options "-DVECTOR_LENGTH=vector_length(128)" } */
|
||||
|
||||
/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */
|
||||
/* { dg-additional-options "-foffload=-fdump-tree-oaccloops" } */
|
||||
/* { dg-set-target-env-var "GOMP_DEBUG" "1" } */
|
||||
|
||||
/* Minimized from ref-1.C. */
|
||||
|
|
@ -54,5 +54,5 @@ main (void)
|
|||
return 0;
|
||||
}
|
||||
|
||||
/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 32\\)" "oaccdevlow" } } */
|
||||
/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 32\\)" "oaccloops" } } */
|
||||
/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=1, vectors=32" } */
|
||||
|
|
|
|||
|
|
@ -4,7 +4,7 @@
|
|||
{ dg-skip-if "TODO PR82391" { *-*-* } { "-O0" } }
|
||||
*/
|
||||
|
||||
/* { dg-additional-options "-fdump-tree-oaccdevlow" } */
|
||||
/* { dg-additional-options "-fdump-tree-oaccloops" } */
|
||||
|
||||
/* { dg-additional-options "-fno-inline" } for stable results regarding OpenACC 'routine'. */
|
||||
|
||||
|
|
@ -36,9 +36,9 @@ static int fact_nohost(int n)
|
|||
|
||||
return fact(n);
|
||||
}
|
||||
/* { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'fact_nohost' has 'nohost' clause\.$} 1 oaccdevlow { target c } } }
|
||||
{ dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'int fact_nohost\(int\)' has 'nohost' clause\.$} 1 oaccdevlow { target { c++ && { ! offloading_enabled } } } } }
|
||||
{ dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'fact_nohost\(int\)' has 'nohost' clause\.$} 1 oaccdevlow { target { c++ && offloading_enabled } } } }
|
||||
/* { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'fact_nohost' has 'nohost' clause\.$} 1 oaccloops { target c } } }
|
||||
{ dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'int fact_nohost\(int\)' has 'nohost' clause\.$} 1 oaccloops { target { c++ && { ! offloading_enabled } } } } }
|
||||
{ dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'fact_nohost\(int\)' has 'nohost' clause\.$} 1 oaccloops { target { c++ && offloading_enabled } } } }
|
||||
TODO See PR101551 for 'offloading_enabled' differences. */
|
||||
|
||||
int main()
|
||||
|
|
|
|||
|
|
@ -1,5 +1,5 @@
|
|||
/* { dg-do run { target openacc_nvidia_accel_selected } } */
|
||||
/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */
|
||||
/* { dg-additional-options "-foffload=-fdump-tree-oaccloops" } */
|
||||
/* { dg-set-target-env-var "GOMP_DEBUG" "1" } */
|
||||
|
||||
#include <stdlib.h>
|
||||
|
|
@ -34,5 +34,5 @@ main (void)
|
|||
return 0;
|
||||
}
|
||||
|
||||
/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 128\\)" "oaccdevlow" } } */
|
||||
/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 128\\)" "oaccloops" } } */
|
||||
/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=1, vectors=128" } */
|
||||
|
|
|
|||
|
|
@ -1,6 +1,6 @@
|
|||
/* { dg-do run { target openacc_nvidia_accel_selected } } */
|
||||
/* { dg-additional-options "-fopenacc-dim=::128" } */
|
||||
/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */
|
||||
/* { dg-additional-options "-foffload=-fdump-tree-oaccloops" } */
|
||||
/* { dg-set-target-env-var "GOMP_DEBUG" "1" } */
|
||||
|
||||
#include <stdlib.h>
|
||||
|
|
@ -35,5 +35,5 @@ main (void)
|
|||
return 0;
|
||||
}
|
||||
|
||||
/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 128\\)" "oaccdevlow" } } */
|
||||
/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 128\\)" "oaccloops" } } */
|
||||
/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=1, vectors=128" } */
|
||||
|
|
|
|||
|
|
@ -1,5 +1,5 @@
|
|||
/* { dg-do run { target openacc_nvidia_accel_selected } } */
|
||||
/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */
|
||||
/* { dg-additional-options "-foffload=-fdump-tree-oaccloops" } */
|
||||
/* We default to warp size 32 for the vector length, so the GOMP_OPENACC_DIM has
|
||||
no effect. */
|
||||
/* { dg-set-target-env-var "GOMP_OPENACC_DIM" "::128" } */
|
||||
|
|
@ -38,5 +38,5 @@ main (void)
|
|||
return 0;
|
||||
}
|
||||
|
||||
/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 32\\)" "oaccdevlow" } } */
|
||||
/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 32\\)" "oaccloops" } } */
|
||||
/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=1, vectors=32" } */
|
||||
|
|
|
|||
|
|
@ -1,5 +1,5 @@
|
|||
/* { dg-do run { target openacc_nvidia_accel_selected } } */
|
||||
/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */
|
||||
/* { dg-additional-options "-foffload=-fdump-tree-oaccloops" } */
|
||||
/* { dg-set-target-env-var "GOMP_DEBUG" "1" } */
|
||||
|
||||
#include <stdlib.h>
|
||||
|
|
@ -36,5 +36,5 @@ main (void)
|
|||
return 0;
|
||||
}
|
||||
|
||||
/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 2, 128\\)" "oaccdevlow" } } */
|
||||
/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 2, 128\\)" "oaccloops" } } */
|
||||
/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=2, vectors=128" } */
|
||||
|
|
|
|||
|
|
@ -1,6 +1,6 @@
|
|||
/* { dg-do run { target openacc_nvidia_accel_selected } } */
|
||||
/* { dg-additional-options "-fopenacc-dim=:2:128" } */
|
||||
/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */
|
||||
/* { dg-additional-options "-foffload=-fdump-tree-oaccloops" } */
|
||||
/* { dg-set-target-env-var "GOMP_DEBUG" "1" } */
|
||||
|
||||
#include <stdlib.h>
|
||||
|
|
@ -37,5 +37,5 @@ main (void)
|
|||
return 0;
|
||||
}
|
||||
|
||||
/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 2, 128\\)" "oaccdevlow" } } */
|
||||
/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 2, 128\\)" "oaccloops" } } */
|
||||
/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=2, vectors=128" } */
|
||||
|
|
|
|||
|
|
@ -1,6 +1,6 @@
|
|||
/* { dg-do run { target openacc_nvidia_accel_selected } } */
|
||||
/* { dg-set-target-env-var "GOMP_OPENACC_DIM" ":2:" } */
|
||||
/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */
|
||||
/* { dg-additional-options "-foffload=-fdump-tree-oaccloops" } */
|
||||
/* { dg-set-target-env-var "GOMP_DEBUG" "1" } */
|
||||
|
||||
#include <stdlib.h>
|
||||
|
|
@ -37,5 +37,5 @@ main (void)
|
|||
return 0;
|
||||
}
|
||||
|
||||
/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 0, 128\\)" "oaccdevlow" } } */
|
||||
/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 0, 128\\)" "oaccloops" } } */
|
||||
/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=2, vectors=128" } */
|
||||
|
|
|
|||
|
|
@ -1,5 +1,5 @@
|
|||
/* { dg-do run { target openacc_nvidia_accel_selected } } */
|
||||
/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */
|
||||
/* { dg-additional-options "-foffload=-fdump-tree-oaccloops" } */
|
||||
/* { dg-set-target-env-var "GOMP_DEBUG" "1" } */
|
||||
|
||||
#include <stdlib.h>
|
||||
|
|
@ -36,5 +36,5 @@ main (void)
|
|||
return 0;
|
||||
}
|
||||
|
||||
/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 0, 128\\)" "oaccdevlow" } } */
|
||||
/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 0, 128\\)" "oaccloops" } } */
|
||||
/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=8, vectors=128" } */
|
||||
|
|
|
|||
|
|
@ -5,7 +5,7 @@
|
|||
! With optimizations disabled, we currently don't expect that 'acc_on_device' "evaluates at compile time to a constant".
|
||||
! { dg-skip-if "TODO PR82391" { *-*-* } { "-O0" } }
|
||||
|
||||
! { dg-additional-options "-fdump-tree-oaccdevlow" }
|
||||
! { dg-additional-options "-fdump-tree-oaccloops" }
|
||||
|
||||
program main
|
||||
use openacc
|
||||
|
|
@ -58,6 +58,6 @@ function fact_nohost(x) result(res)
|
|||
|
||||
res = fact(x)
|
||||
end function fact_nohost
|
||||
! { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'fact_nohost' has 'nohost' clause\.$} 1 oaccdevlow { target { ! offloading_enabled } } } }
|
||||
! { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'fact_nohost_' has 'nohost' clause\.$} 1 oaccdevlow { target offloading_enabled } } }
|
||||
! { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'fact_nohost' has 'nohost' clause\.$} 1 oaccloops { target { ! offloading_enabled } } } }
|
||||
! { dg-final { scan-tree-dump-times {(?n)^OpenACC routine 'fact_nohost_' has 'nohost' clause\.$} 1 oaccloops { target offloading_enabled } } }
|
||||
!TODO See PR101551 for 'offloading_enabled' differences.
|
||||
|
|
|
|||
Loading…
Add table
Reference in a new issue