diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c index bfbb0112e24..d881426ae65 100644 --- a/gcc/omp-offload.c +++ b/gcc/omp-offload.c @@ -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 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) { diff --git a/gcc/passes.def b/gcc/passes.def index e2858368b7d..26d86df2f5a 100644 --- a/gcc/passes.def +++ b/gcc/passes.def @@ -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); diff --git a/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c b/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c index 218f6248062..1d12658790d 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c @@ -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" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/classify-kernels.c b/gcc/testsuite/c-c++-common/goacc/classify-kernels.c index 95a150ca9ac..bdf7b4a0641 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-kernels.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-kernels.c @@ -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" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/classify-parallel.c b/gcc/testsuite/c-c++-common/goacc/classify-parallel.c index 230e70c66cd..9056aa69dad 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-parallel.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-parallel.c @@ -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" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/classify-routine-nohost.c b/gcc/testsuite/c-c++-common/goacc/classify-routine-nohost.c index a58482f7f92..99855822011 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-routine-nohost.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-routine-nohost.c @@ -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" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/classify-routine.c b/gcc/testsuite/c-c++-common/goacc/classify-routine.c index cc0ba2b9a7d..f7f0454009b 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-routine.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-routine.c @@ -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" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/classify-serial.c b/gcc/testsuite/c-c++-common/goacc/classify-serial.c index ae052ae6a1c..f41c141bcd5 100644 --- a/gcc/testsuite/c-c++-common/goacc/classify-serial.c +++ b/gcc/testsuite/c-c++-common/goacc/classify-serial.c @@ -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" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/routine-nohost-1.c b/gcc/testsuite/c-c++-common/goacc/routine-nohost-1.c index c8927416efa..59ebb2bc5a9 100644 --- a/gcc/testsuite/c-c++-common/goacc/routine-nohost-1.c +++ b/gcc/testsuite/c-c++-common/goacc/routine-nohost-1.c @@ -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 } } */ diff --git a/gcc/testsuite/g++.dg/goacc/template.C b/gcc/testsuite/g++.dg/goacc/template.C index f34fcfea52d..10d3f446da7 100644 --- a/gcc/testsuite/g++.dg/goacc/template.C +++ b/gcc/testsuite/g++.dg/goacc/template.C @@ -1,4 +1,4 @@ -/* { dg-additional-options "-fdump-tree-oaccdevlow" } */ +/* { dg-additional-options "-fdump-tree-oaccloops" } */ #pragma acc routine nohost template 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\(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' 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\(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\(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\(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' 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\(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\(int\)double' has 'nohost' clause\.$} 1 oaccloops { target offloading_enabled } } } TODO See PR101551 for 'offloading_enabled' differences. */ diff --git a/gcc/testsuite/gcc.dg/goacc/loop-processing-1.c b/gcc/testsuite/gcc.dg/goacc/loop-processing-1.c index bd4c07e7d81..78b9aed89be 100644 --- a/gcc/testsuite/gcc.dg/goacc/loop-processing-1.c +++ b/gcc/testsuite/gcc.dg/goacc/loop-processing-1.c @@ -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" } } */ diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95 index cb5251a2aeb..3fb48b321f2 100644 --- a/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95 @@ -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" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95 index 07aaf065e1d..6c8d298e236 100644 --- a/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95 @@ -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" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-parallel.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-parallel.f95 index a41e0e68b38..ce4c08ff219 100644 --- a/gcc/testsuite/gfortran.dg/goacc/classify-parallel.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/classify-parallel.f95 @@ -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" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-routine-nohost.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-routine-nohost.f95 index 0e06fb9f0ba..07e2063551f 100644 --- a/gcc/testsuite/gfortran.dg/goacc/classify-routine-nohost.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/classify-routine-nohost.f95 @@ -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. diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-routine.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-routine.f95 index 92d3243cdcf..b065ccadacd 100644 --- a/gcc/testsuite/gfortran.dg/goacc/classify-routine.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/classify-routine.f95 @@ -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. diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-serial.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-serial.f95 index 6dcb1b170f8..f5cb3fe50c5 100644 --- a/gcc/testsuite/gfortran.dg/goacc/classify-serial.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/classify-serial.f95 @@ -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" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-multiple-directives-1.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-multiple-directives-1.f90 index 44ef4533f04..42bcb0e8d63 100644 --- a/gcc/testsuite/gfortran.dg/goacc/routine-multiple-directives-1.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/routine-multiple-directives-1.f90 @@ -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 diff --git a/gcc/tree-parloops.c b/gcc/tree-parloops.c index bb547572653..4d40c96df7d 100644 --- a/gcc/tree-parloops.c +++ b/gcc/tree-parloops.c @@ -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 { diff --git a/gcc/tree-pass.h b/gcc/tree-pass.h index 1f5b1370a95..5484ad5eac7 100644 --- a/gcc/tree-pass.h +++ b/gcc/tree-pass.h @@ -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); diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-2.c index d45326488cd..17cc9bd663e 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-2.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-2.c @@ -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" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-3.c index 33480a4ae68..5d05540ce46 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-3.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-3.c @@ -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" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486.c index 0d98b82f993..f95f2ee3123 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486.c @@ -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" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-nohost-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-nohost-1.c index dc92727d5be..7dc7459e5fe 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-nohost-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-nohost-1.c @@ -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() diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.c index 18d77cc5ecb..5158bb5eb89 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.c @@ -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 @@ -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" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-2.c index 8b5b2a4a92d..a3e44ebfbcb 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-2.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-2.c @@ -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 @@ -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" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-3.c index 59be37a7c27..a85400d09c5 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-3.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-3.c @@ -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" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-4.c index e5d1df09b8a..24c078f377c 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-4.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-4.c @@ -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 @@ -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" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-5.c index e60f1c28db4..fcca9f593bb 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-5.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-5.c @@ -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 @@ -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" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-6.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-6.c index a1f67622f84..0807eab7eee 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-6.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-6.c @@ -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 @@ -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" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-7.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-7.c index c419f6499b5..4a8c1bf549e 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-7.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-7.c @@ -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 @@ -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" } */ diff --git a/libgomp/testsuite/libgomp.oacc-fortran/routine-nohost-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/routine-nohost-1.f90 index cd5bddc8685..b0537b8ff0b 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/routine-nohost-1.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/routine-nohost-1.f90 @@ -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.