Message ID | 0e1a740e-46d5-ebfa-36f4-9a069ddf8620@codesourcery.com |
---|---|
State | New |
Headers | show |
Series | openmp: fix max_vf setting for amdgcn offloading | expand |
On Tue, Jul 12, 2022 at 03:16:35PM +0100, Andrew Stubbs wrote: > --- a/gcc/gimple-loop-versioning.cc > +++ b/gcc/gimple-loop-versioning.cc > @@ -555,7 +555,10 @@ loop_versioning::loop_versioning (function *fn) > unvectorizable code, since it is the largest size that can be > handled efficiently by scalar code. omp_max_vf calculates the > maximum number of bytes in a vector, when such a value is relevant > - to loop optimization. */ > + to loop optimization. > + FIXME: this probably needs to use omp_max_simd_vf when in a target > + region, but how to tell? (And MAX_FIXED_MODE_SIZE is large enough that > + it doesn't actually matter.) */ > m_maximum_scale = estimated_poly_value (omp_max_vf ()); > m_maximum_scale = MAX (m_maximum_scale, MAX_FIXED_MODE_SIZE); I think this shouldn't have the comment added, the use here actually isn't much OpenMP related, it just uses the function because it implements what it wants. > --- a/gcc/omp-general.cc > +++ b/gcc/omp-general.cc > @@ -994,6 +994,24 @@ omp_max_simt_vf (void) > return 0; > } > > +/* Return maximum SIMD width if offloading may target SIMD hardware. */ > + > +int > +omp_max_simd_vf (void) The name is just confusing. omp_max_vf is about the SIMD maximum VF, so if you really want, rename omp_max_vf to omp_max_simd_vf. For the offloading related stuff, IMHO either we put it into that single omp-general.cc function and add a bool argument to it whether it is or might be in offloading region (ordered_maximum from the returned value and the offloading one, but only after the initialy return 1; conditions and adjust callers), or have this separate function, but then IMHO the if (!optimize) return 0; initial test should be if (!optimize || optimize_debug || !flag_tree_loop_optimize || (!flag_tree_loop_vectorize && OPTION_SET_P (flag_tree_loop_vectorize))) return 1; because without that nothing is vectorized, on host nor on offloading targets, and the function should be called omp_max_target_vf or omp_max_target_simd_vf. > +{ > + if (!optimize) > + return 0; > --- a/gcc/omp-low.cc > +++ b/gcc/omp-low.cc > @@ -4646,7 +4646,14 @@ lower_rec_simd_input_clauses (tree new_var, omp_context *ctx, > { > if (known_eq (sctx->max_vf, 0U)) > { > - sctx->max_vf = sctx->is_simt ? omp_max_simt_vf () : omp_max_vf (); > + /* If we are compiling for multiple devices choose the largest VF. */ > + sctx->max_vf = omp_max_vf (); > + if (omp_maybe_offloaded_ctx (ctx)) > + { > + if (sctx->is_simt) > + sctx->max_vf = ordered_max (sctx->max_vf, omp_max_simt_vf ()); > + sctx->max_vf = ordered_max (sctx->max_vf, omp_max_simd_vf ()); > + } This is wrong. If sctx->is_simt, we know it is the SIMT version. So we want to use omp_max_simt_vf (), not maximum of that and something unrelated. Only if !sctx->is_simt, we want to use maximum of omp_max_vf and if omp_maybe_offloaded_ctx also omp_max_target_vf or how it is called (or pass that as argument to omp_max_vf). We have another omp_max_vf () call though, in omp-expand.cc (omp_adjust_chunk_size). That is for schedule (simd: dynamic, 32) and similar, though unlike the omp-low.cc case (where using a larger VF in that case doesn't hurt, it is used for sizing of the maxic arrays that are afterwards resized to the actual size), using too large values in that case is harmful. So dunno if it should take into account offloading vf or not. Maybe if maybe offloading maybe not it should fold to some internal fn call dependent expression that folds to omp_max_vf of the actual target after IPA. Jakub
Hi! In addition to the technical issues pointed out by Jakub for this og12 commit: On 2022-07-12T15:16:35+0100, Andrew Stubbs <ams@codesourcery.com> wrote: > This patch [...] > I will commit a backport to OG12 shortly. > openmp: fix max_vf setting for amdgcn offloading > --- a/gcc/omp-general.h > +++ b/gcc/omp-general.h > extern poly_uint64 omp_max_vf (void); > extern int omp_max_simt_vf (void); > +extern int omp_max_simd_vf (void); > --- a/gcc/omp-low.cc > +++ b/gcc/omp-low.cc > @@ -4646,7 +4646,14 @@ lower_rec_simd_input_clauses (tree new_var, omp_context *ctx, > { > if (known_eq (sctx->max_vf, 0U)) > { > - sctx->max_vf = sctx->is_simt ? omp_max_simt_vf () : omp_max_vf (); > + /* If we are compiling for multiple devices choose the largest VF. */ > + sctx->max_vf = omp_max_vf (); > + if (omp_maybe_offloaded_ctx (ctx)) > + { > + if (sctx->is_simt) > + sctx->max_vf = ordered_max (sctx->max_vf, omp_max_simt_vf ()); > + sctx->max_vf = ordered_max (sctx->max_vf, omp_max_simd_vf ()); > + } > if (maybe_gt (sctx->max_vf, 1U)) > { > tree c = omp_find_clause (gimple_omp_for_clauses (ctx->stmt), ... I've additionally run into a bootstrap error, and have now pushed "Resolve '-Wsign-compare' issue in 'gcc/omp-low.cc:lower_rec_simd_input_clauses'" to devel/omp/gcc-12 in commit 4e32d1582a137d5f34248fdd3e93d35a798f5221, see attached. Grüße Thomas ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955
diff --git a/gcc/gimple-loop-versioning.cc b/gcc/gimple-loop-versioning.cc index 6bcf6eba691..e908c27fc44 100644 --- a/gcc/gimple-loop-versioning.cc +++ b/gcc/gimple-loop-versioning.cc @@ -555,7 +555,10 @@ loop_versioning::loop_versioning (function *fn) unvectorizable code, since it is the largest size that can be handled efficiently by scalar code. omp_max_vf calculates the maximum number of bytes in a vector, when such a value is relevant - to loop optimization. */ + to loop optimization. + FIXME: this probably needs to use omp_max_simd_vf when in a target + region, but how to tell? (And MAX_FIXED_MODE_SIZE is large enough that + it doesn't actually matter.) */ m_maximum_scale = estimated_poly_value (omp_max_vf ()); m_maximum_scale = MAX (m_maximum_scale, MAX_FIXED_MODE_SIZE); } diff --git a/gcc/omp-general.cc b/gcc/omp-general.cc index a406c578f33..8c6fcebc4b3 100644 --- a/gcc/omp-general.cc +++ b/gcc/omp-general.cc @@ -994,6 +994,24 @@ omp_max_simt_vf (void) return 0; } +/* Return maximum SIMD width if offloading may target SIMD hardware. */ + +int +omp_max_simd_vf (void) +{ + if (!optimize) + return 0; + if (ENABLE_OFFLOADING) + for (const char *c = getenv ("OFFLOAD_TARGET_NAMES"); c;) + { + if (startswith (c, "amdgcn")) + return 64; + else if ((c = strchr (c, ':'))) + c++; + } + return 0; +} + /* Store the construct selectors as tree codes from last to first, return their number. */ diff --git a/gcc/omp-general.h b/gcc/omp-general.h index 74e90e1a71a..410343e45fa 100644 --- a/gcc/omp-general.h +++ b/gcc/omp-general.h @@ -104,6 +104,7 @@ extern gimple *omp_build_barrier (tree lhs); extern tree find_combined_omp_for (tree *, int *, void *); extern poly_uint64 omp_max_vf (void); extern int omp_max_simt_vf (void); +extern int omp_max_simd_vf (void); extern int omp_constructor_traits_to_codes (tree, enum tree_code *); extern tree omp_check_context_selector (location_t loc, tree ctx); extern void omp_mark_declare_variant (location_t loc, tree variant, diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc index d73c165f029..1a9a509adb9 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -4646,7 +4646,14 @@ lower_rec_simd_input_clauses (tree new_var, omp_context *ctx, { if (known_eq (sctx->max_vf, 0U)) { - sctx->max_vf = sctx->is_simt ? omp_max_simt_vf () : omp_max_vf (); + /* If we are compiling for multiple devices choose the largest VF. */ + sctx->max_vf = omp_max_vf (); + if (omp_maybe_offloaded_ctx (ctx)) + { + if (sctx->is_simt) + sctx->max_vf = ordered_max (sctx->max_vf, omp_max_simt_vf ()); + sctx->max_vf = ordered_max (sctx->max_vf, omp_max_simd_vf ()); + } if (maybe_gt (sctx->max_vf, 1U)) { tree c = omp_find_clause (gimple_omp_for_clauses (ctx->stmt), diff --git a/gcc/testsuite/gcc.dg/gomp/target-vf.c b/gcc/testsuite/gcc.dg/gomp/target-vf.c new file mode 100644 index 00000000000..14cea45e53c --- /dev/null +++ b/gcc/testsuite/gcc.dg/gomp/target-vf.c @@ -0,0 +1,21 @@ +/* { dg-do compile } */ +/* { dg-options "-fopenmp -O2 -fdump-tree-omplower" } */ + +/* Ensure that the omp_max_vf, omp_max_simt_vf, and omp_max_simd_vf are working + properly to set the OpenMP vectorization factor for the offload target, and + not just for the host. */ + +float +foo (float * __restrict x, float * __restrict y) +{ + float sum = 0.0; + +#pragma omp target teams distribute parallel for simd map(tofrom: sum) reduction(+:sum) + for (int i=0; i<1024; i++) + sum += x[i] * y[i]; + + return sum; +} + +/* { dg-final { scan-tree-dump "safelen\\(64\\)" "omplower" { target amdgcn_offloading_enabled } } } */ +/* { dg-final { scan-tree-dump "safelen\\(32\\)" "omplower" { target { { nvptx_offloading_enabled } && { ! amdgcn_offloading_enabled } } } } } */ diff --git a/gcc/testsuite/lib/target-supports.exp b/gcc/testsuite/lib/target-supports.exp index 4ed7b25b9a4..363354be461 100644 --- a/gcc/testsuite/lib/target-supports.exp +++ b/gcc/testsuite/lib/target-supports.exp @@ -1025,6 +1025,16 @@ proc check_effective_target_offloading_enabled {} { return [check_configured_with "--enable-offload-targets"] } +# Return 1 if compiled with --enable-offload-targets=amdgcn +proc check_effective_target_amdgcn_offloading_enabled {} { + return [check_configured_with {--enable-offload-targets=[^ ]*amdgcn}] +} + +# Return 1 if compiled with --enable-offload-targets=amdgcn +proc check_effective_target_nvptx_offloading_enabled {} { + return [check_configured_with {--enable-offload-targets=[^ ]*nvptx}] +} + # Return 1 if compilation with -fopenacc is error-free for trivial # code, 0 otherwise.