> -----Original Message-----
> From: Richard Biener <richard.guent...@gmail.com>
> Sent: Monday, July 29, 2024 7:18 PM
> To: Prathamesh Kulkarni <prathame...@nvidia.com>
> Cc: gcc@gcc.gnu.org
> Subject: Re: [RFC] Summary of libgomp failures for offloading to nvptx
> from AArch64
>
> External email: Use caution opening links or attachments
>
>
> On Mon, Jul 29, 2024 at 1:35 PM Prathamesh Kulkarni
> <prathame...@nvidia.com> wrote:
> >
> >
> >
> > > -----Original Message-----
> > > From: Richard Biener <richard.guent...@gmail.com>
> > > Sent: Friday, July 26, 2024 6:51 PM
> > > To: Prathamesh Kulkarni <prathame...@nvidia.com>
> > > Cc: gcc@gcc.gnu.org
> > > Subject: Re: [RFC] Summary of libgomp failures for offloading to
> > > nvptx from AArch64
> > >
> > > External email: Use caution opening links or attachments
> > >
> > >
> > > On Thu, Jul 25, 2024 at 3:36 PM Prathamesh Kulkarni via Gcc
> > > <gcc@gcc.gnu.org> wrote:
> > > >
> > > > Hi,
> > > > I am working on enabling offloading to nvptx from AAarch64 host.
> > > > As mentioned on wiki
> > > > (https://gcc.gnu.org/wiki/Offloading#Running_.27make_check.27),
> > > > I ran make check-target-libgomp on AAarch64 host (and no GPU)
> with
> > > following results:
> > > >
> > > > === libgomp Summary ===
> > > >
> > > > # of expected passes 14568
> > > > # of unexpected failures 1023
> > > > # of expected failures 309
> > > > # of untested testcases 54
> > > > # of unresolved testcases 992
> > > > # of unsupported tests 644
> > > >
> > > > It seems majority of the tests fail due to the following 4
> issues:
> > > >
> > > > * Compiling a minimal test-case:
> > > >
> > > > int main()
> > > > {
> > > > int x;
> > > > #pragma omp target map (to: x)
> > > > {
> > > > x = 0;
> > > > }
> > > > return x;
> > > > }
> > > >
> > > > Compiling with -fopenmp -foffload=nvptx-none results in
> following
> > > issues:
> > > >
> > > > (1) Differing values of NUM_POLY_INT_COEFFS between host and
> > > accelerator, which results in following ICE:
> > > >
> > > > 0x1a6e0a7 pp_quoted_string
> > > > ../../gcc/gcc/pretty-print.cc:2277
> > > > 0x1a6ffb3 pp_format(pretty_printer*, text_info*, urlifier
> const*)
> > > > ../../gcc/gcc/pretty-print.cc:1634
> > > > 0x1a4a3f3
> diagnostic_context::report_diagnostic(diagnostic_info*)
> > > > ../../gcc/gcc/diagnostic.cc:1612
> > > > 0x1a4a727 diagnostic_impl
> > > > ../../gcc/gcc/diagnostic.cc:1775 0x1a4e20b
> > > > fatal_error(unsigned int, char const*, ...)
> > > > ../../gcc/gcc/diagnostic.cc:2218 0xb3088f
> > > > lto_input_mode_table(lto_file_decl_data*)
> > > > ../../gcc/gcc/lto-streamer-in.cc:2121
> > > > 0x6f5cdf lto_file_finalize
> > > > ../../gcc/gcc/lto/lto-common.cc:2285
> > > > 0x6f5cdf lto_create_files_from_ids
> > > > ../../gcc/gcc/lto/lto-common.cc:2309
> > > > 0x6f5cdf lto_file_read
> > > > ../../gcc/gcc/lto/lto-common.cc:2364
> > > > 0x6f5cdf read_cgraph_and_symbols(unsigned int, char const**)
> > > > ../../gcc/gcc/lto/lto-common.cc:2812
> > > > 0x6cfb93 lto_main()
> > > > ../../gcc/gcc/lto/lto.cc:658
> > > >
> > > > This is already tracked in https://gcc.gnu.org/PR96265 (and
> > > > related
> > > > PR's)
> > > >
> > > > Streaming out mode_table:
> > > > mode = SI, mclass = 2, size = 4, prec = 32 mode = DI, mclass =
> 2,
> > > size
> > > > = 8, prec = 64
> > > >
> > > > Streaming in mode_table (in lto_input_mode_table):
> > > > mclass = 2, size = 4, prec = 0
> > > > (and then calculates the correct mode value by iterating over
> all
> > > > modes of mclass starting from narrowest mode)
> > > >
> > > > The issue is that the value for prec is not getting streamed-in
> > > > correctly for SImode as seen above. While streaming out from
> > > > AArch64
> > > host, it is 32, but while streaming in for nvptx, it is 0. This
> > > happens because of differing values of NUM_POLY_INT_COEFFS between
> > > AArch64 and nvptx backend.
> > > >
> > > > Since NUM_POLY_INT_COEFFS is 2 for aarch64, the streamed-out
> > > > values for mode, precision would be <4, 0> and <32, 0>
> > > > respectively (streamed-out in bp_pack_poly_value). Both zeros
> come
> > > > from coeffs[1] of size and prec. While streaming in however,
> > > > NUM_POLY_INT_COEFFS is
> > > 1 for nvptx, and thus it incorrectly treats <4, 0> as size and
> > > precision respectively, which is why precision gets streamed in as
> > > 0, and thus it encounters the above ICE.
> > > >
> > > > Supporting non VLA code with offloading:
> > > >
> > > > In the general case, it's hard to support offloading for
> arbitrary
> > > poly_ints when NUM_POLY_INT_COEFFS differs for host and
> accelerator.
> > > > For example, it's not possible to represent a degree-2 poly_int
> > > > like
> > > 4 + 4x (as-is) on an accelerator with NUM_POLY_INT_COEFFS == 1.
> > > >
> > > > However, IIUC, we can support offloading for restricted set of
> > > > poly_ints whose degree <= accel's NUM_POLY_INT_COEFFS, since
> they
> > > can
> > > > be represented on accelerator ? For a hypothetical example, if
> > > > host
> > > NUM_POLY_INT_COEFFS == 3 and accel NUM_POLY_INT_COEFFS == 2, then
> I
> > > suppose we could represent a degree 2 poly_int on accelerator, but
> > > not a degree 3 poly_int like 3+4x+5x^2 ?
> > > >
> > > > Based on that, I have come up with following approach in
> attached
> > > "quick-and-dirty" patch (p-163-2.diff):
> > > > Stream-out host NUM_POLY_INT_COEFFS, and while streaming-in
> during
> > > lto1, compare it with accelerator's NUM_POLY_INT_COEFFS as
> follows:
> > > >
> > > > Stream in host_num_poly_int_coeffs; if (host_num_poly_int_coeffs
> > > > == NUM_POLY_INT_COEFFS) //
> > > NUM_POLY_INT_COEFFS represents accelerator's value here.
> > > > {
> > > > /* Both are equal, proceed to unpacking NUM_POLY_INT_COEFFS
> > > words
> > > > from bitstream. */ } else if (host_num_poly_int_coeffs <
> > > > NUM_POLY_INT_COEFFS) {
> > > > /* Unpack host_num_poly_int_coeffs words and zero out
> > > > remaining higher coeffs (similar to zero-extension). */ } else
> {
> > > > /* Unpack host_num_poly_int_coeffs words and ensure that
> > > > degree
> > > of
> > > > streamed-out poly_int <= NUM_POLY_INT_COEFFS. */ }
> > > >
> > > > For example, with host NUM_POLY_INT_COEFFS == 2 and accel
> > > > NUM_POLY_INT_COEFFS == 1, this will allow streaming of "degree-
> 1"
> > > poly_ints like 4+0x (which will degenerate to constant 4), but
> give
> > > an error for streaming degree-2 poly_int like 4+4x.
> > > >
> > > > Following this approach, I am assuming we can support
> > > > AArch64/nvptx offloading for non VLA code, since poly_ints used
> > > > for representing various artefacts like mode_size,
> mode_precision, vector length etc.
> > > will be degree-1 poly_int for scalar variables and fixed-length
> > > vectors (and thus degenerate to constants). With the patch
> applied,
> > > it proceeds forward from this point, but fails at a later stage
> (see
> > > below).
> > > >
> > > > Does this approach look reasonable for supporting offloading for
> > > > non
> > > VLA code ?
> > > > Are there any cases I may have overlooked, where offloading will
> > > still fail for non-VLA code due to differing NUM_POLY_INT_COEFFS
> > > issue ?
> > >
> > > I think I'd change how we stream POLY_INTs and make that
> "independent"
> > > of NUM_POLY_INT_COEFFS
> > > in that I'd stream the effective number of coeffs required - as
> you
> > > suggest "trailing" zero coeffs do not need to be represented. We
> > > always stream coeff zero.
> > >
> > > On read in when that N is bigger than NUM_POLY_INT_COEFFS we have
> to
> > > error and we zero-fill extra elements (using the type of coeff
> zero).
> > >
> > > > (2) nvptx mkoffload.cc passes -m64/-m32 to host compiler if
> > > > -foffload-abi=lp64/ilp32 After applying workaround for the above
> > > assertion failure, it hits the following error:
> > > > gcc: error: unrecognized command-line option '-m64'
> > > > nvptx mkoffload: fatal error: ../install/bin/gcc returned 1 exit
> > > > status compilation terminated.
> > > >
> > > > This happens because nvptx/mkoffload.cc:compile_native passes
> > > > -m64/-m32 to host compiler depending on whether offload_abi is
> > > OFFLOAD_ABI_LP64 or OFFLOAD_ABI_ILP32, and aarch64 backend doesn't
> > > recognize these options.
> > > >
> > > > I suppose a simple solution to check if host_compiler supports a
> > > > particular command-line option, would be to create a dummy C
> file
> > > and
> > > > check if the command "host_compiler <opt> dummy_file.c" returns
> > > > zero
> > > exit status. Alternative could be to check exit status for
> > > "host_compiler <opt> --version", once http://gcc.gnu.org/PR116050
> is
> > > fixed, but I am not sure if either is an ideal solution.
> > >
> > > I think an explicit configure when building the offload compiler
> > > would be cleaner, like simply adding an additional
> > > AC_SUBST(host_no_multilibs) var computed from the --enable-as-
> > > accelerator-for target triplet?
> > >
> > > > With workarounds for these 2 issues, the minimal test builds and
> > > runs successfully.
> > > >
> > > > (3) Assertion error in lto_read_decls during lto1:
> > > > There are several failures (~350+) in the testsuite caused due
> to
> > > the
> > > > following assert in lto_read_decls:
> > > > gcc_assert (data_in->reader_cache->nodes.length () == from + 1);
> > > >
> > > > AFAIU, this seems to happen because of presence of LTO_null tag.
> > > > The following workaround avoids hitting the assert, but am not
> > > > sure if
> > > it's the right fix:
> > > >
> > > > t = lto_input_tree_1 (&ib_main, data_in, tag, 0);
> > > > + if (t == NULL_TREE)
> > > > + continue;
> > > > gcc_assert (data_in->reader_cache->nodes.length () ==
> from
> > > > + 1);
> > > >
> > > > (FWIW, this was reproducible with the above minimal test, but
> has
> > > > seemingly gone away for it after updating the sources recently,
> > > > but still reproduces with libgomp tests like for-9.c, baseptrs-
> 4.C
> > > > etc.)
> > >
> > > It looks like a discrepancy between host and target that wasn't
> > > properly analyzed to me. The assert makes sure we read the
> expected
> > > number of tree nodes - what's the tag when you hit this?
> > > Please open a bug.
> > Sure, will do shortly.
> >
> > I checked with for-9.c, and it shows the following state on hitting
> the assert:
> > reader_cache->nodes.length () = 2384, from+1 = 2385, tag: LTO_null
> >
> > From lto_read_decls:
> > unsigned from = data_in->reader_cache->nodes.length ();
> > /* Read and uniquify SCCs as in the input stream. */
> > enum LTO_tags tag = streamer_read_record_start (&ib_main);
> > if (tag == LTO_tree_scc || tag == LTO_trees)
> >
> > In this case, tag is LTO_null, and it skips to the else part:
> > t = lto_input_tree_1 (&ib_main, data_in, tag, 0);
> > gcc_assert (data_in->reader_cache->nodes.length () == from +
> 1);
> >
> > and lto_input_tree_1 simply returns NULL_TREE if tag is LTO_null
> without any additional streaming-in.
> > So IIUC, data_in->reader_cache->nodes.length () won't change since
> > there is no tree node streamed in lto_input_tree_1, and thus
> reader_cache->length() remains equal to from (rather than from + 1),
> which triggers the assert ?
>
> Interesting. The thing is that we don't expect LTO_null when
> streaming in global decls and types - we wouldn't have streamed a NULL
> decl or type. So can you debug this on the writer side and see where
> we stream such?
Hi Richard,
Sorry for late response.
The ICE can also be reproduced with an even simpler test-case:
int main()
{
int x;
#pragma omp target map(x)
x;
}
Compiling with -O3 -fopenmp -foffload=nvptx-none hits the same assert in
lto_read_decls.
I think the issue here possibly is corrupted streaming of optimization_node, due
to presence of AArch64 specific optimization options in LTO bytecode.
From AArch64 cl_optimization_stream_out:
bp_pack_var_len_int (bp, ptr->x_flag_wrapv_pointer);
bp_pack_var_len_int (bp, ptr->x_debug_nonbind_markers_p);
bp_pack_var_len_int (bp, ptr->x_flag_aarch64_early_ldp_fusion);
bp_pack_var_len_int (bp, ptr->x_aarch64_early_ra);
bp_pack_var_len_int (bp, ptr->x_flag_aarch64_late_ldp_fusion);
bp_pack_var_len_int (bp, ptr->x_flag_mlow_precision_div);
bp_pack_var_len_int (bp, ptr->x_flag_mrecip_low_precision_sqrt);
bp_pack_var_len_int (bp, ptr->x_flag_mlow_precision_sqrt);
for (size_t i = 0; i < ARRAY_SIZE (ptr->explicit_mask); i++)
bp_pack_value (bp, ptr->explicit_mask[i], 64);
And nvptx cl_optimization_stream_in:
ptr->x_flag_wrapv_pointer = (signed char ) bp_unpack_var_len_int (bp);
ptr->x_debug_nonbind_markers_p = (signed char ) bp_unpack_var_len_int (bp);
for (size_t i = 0; i < ARRAY_SIZE (ptr->explicit_mask); i++)
ptr->explicit_mask[i] = bp_unpack_value (bp, 64);
While AArch64 host streams out target-specific opts like
flag_aarch64_early_ldp_fusion, aarch64_early_ra etc., there's no corresponding
stream-in for these options for nvptx. And will thus result in invalid
streaming-in for ptr->explicit_mask (and subsequent data structures).
To verify if this was indeed the cause of failure, I (temporarily) changed
options in aarch64.opt marked with Optimization
to Save, and verified that AArch64 cl_optimization_stream_out matches with
nvptx cl_optimization_stream_in,
which prevents the above ICE.
FWIW, x86_64 cl_optimization_stream_out also similarly streams out
ix86_unroll_only_small_loops:
bp_pack_var_len_int (bp, ptr->x_flag_wrapv_pointer);
bp_pack_var_len_int (bp, ptr->x_debug_nonbind_markers_p);
bp_pack_var_len_int (bp, ptr->x_ix86_unroll_only_small_loops);
for (size_t i = 0; i < ARRAY_SIZE (ptr->explicit_mask); i++)
bp_pack_value (bp, ptr->explicit_mask[i], 64);
which I suppose can cause similar streaming issues ?
This change was introduced recently in:
https://gcc.gnu.org/git/?p=gcc.git;a=commit;h=a59c4e496fa916cb9a484a649aa1b4cebd6550f2
Which perhaps explains recent failures I am seeing in libgomp tests for
x86_64->nvptx offloading:
FAIL: libgomp.fortran/target-print-1.f90 -O2 (internal compiler error: in
lto_read_decls, at lto/lto-common.cc:1970)
FAIL: libgomp.oacc-c++/../libgomp.oacc-c-c++-common/private-variables.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O2
(internal compiler error: in lto_read_decls, at lto/lto-common.cc:1970)
FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/private-variables.c
-DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O2
(internal compiler error: in lto_read_decls, at lto/lto-common.cc:1970)
FAIL: libgomp.oacc-fortran/print-1-nvptx.f90 -DACC_DEVICE_TYPE_nvidia=1
-DACC_MEM_SHARED=0 -foffload=nvptx-none -O2 (internal compiler error: in
lto_read_decls, at lto/lto-common.cc:1970)
Thanks,
Prathamesh
>
> Richard.
>
> > >
> > > >
> > > > (4) AAarch64 uses OImode for 256-bit size array, which is not
> > > supported on nvptx:
> > > > This causes ~18 tests to fail.
> > > >
> > > > Can be reproduced with following simple test:
> > > > int main()
> > > > {
> > > > long c[4];
> > > > #pragma omp target map(c)
> > > > c[0] = 0;
> > > > return 0;
> > > > }
> > > >
> > > > Compiling with -O2 -fopenmp -foffload=nvptx-none results in:
> > > > lto1: fatal error: nvptx-none - 256-bit integer numbers
> > > > unsupported (mode 'OI') compilation terminated.
> > > > nvptx mkoffload: fatal error:
> > > > ../install/bin/aarch64-unknown-linux-gnu-accel-nvptx-none-gcc
> > > returned 1 exit status compilation terminated.
> > > >
> > > > This happens with AArch64 host because, it uses OImode (256-bit
> > > > integer mode) for ARRAY_TYPE (long c[4] fits 256-bits), which
> > > > isn't supported on nvptx. This decision is made during
> layout_type
> > > > for
> > > 'c', which calls mode_for_array, and mode_for_array uses target
> > > hooks array_mode and array_mode_supported_p to determine
> > > target-specific modes to use for ARRAY_TYPE.
> > > > For x86_64, AFAIK, it uses BLKmode for ARRAY_TYPE.
> > > >
> > > > I have attached a "quick-and-dirty" patch (p-166-2.diff) which
> > > > falls back to using BLKmode for ARRAY_TYPE if offloading is
> > > > enabled, and avoids streaming-out target-specific int modes in
> > > > lto_write_mode_table. I used default_scalar_mode_supported_p
> check
> > > to test if the int_mode is "generic", but not sure if that's
> > > entirely correct. The test compiles and runs OK with patch
> applied.
> > > I suppose a more general solution would be to somehow "intersect"
> > > available
> > > AArch64 modes with nvptx modes, and use those for offloading ?
> > >
> > > Hmm, I think we shouldn't stream modes for aggregates but instead
> > > let stor-layout re-assign them. They are not usually expected to
> match.
> > > Not sure if we should do that generally or only for offload
> > > streaming though (it does have an overhead to recompute the mode).
> > > We could stream VOIDmode here as indicator it needs recomputation
> > > (splitting out sub-workers from layout_type, there's one already
> for
> > > record type).
> > Thanks, this indeed sounds better. I will try to work on patch
> following this approach.
> >
> > Thanks,
> > Prathamesh
> > >
> > > > With local workarounds for the above 4 issues, running make
> check-
> > > target-libgomp shows following results:
> > > >
> > > > === libgomp Summary ===
> > > >
> > > > # of expected passes 16604
> > > > # of unexpected failures 10
> > > > # of expected failures 309
> > > > # of untested testcases 54
> > > > # of unresolved testcases 3
> > > > # of unsupported tests 643
> > > >
> > > > The remaining issues are:
> > > >
> > > > (5) "error: alias definitions not supported in this
> configuration"
> > > > Fails for pr96390.c, and pr96390.C. This seems to be related to
> > > > https://gcc.gnu.org/PR97102
> > > >
> > > > (6) Execution Failures:
> > > > - libgomp/pr104783.c
> > > > - libgomp/pr104783-2.c
> > > > I haven't investigated these yet.
> > > >
> > > > (7) Several warnings fail for libgomp.oacc-c-c++-
> common/acc_prof-
> > > kernels-1.c and following excess errors:
> > > > acc_prof-kernels-1.c:185:9: optimized: assigned OpenACC seq loop
> > > > parallelism
> > > > acc_prof-kernels-1.c:214:9: optimized: assigned OpenACC seq loop
> > > > parallelism
> > > > acc_prof-kernels-1.c:245:9: optimized: assigned OpenACC seq loop
> > > > parallelism
> > > >
> > > > So far, I have only been testing make check-target-libgomp.
> Should
> > > > I
> > > be testing any additional parts of the testsuite for offloading
> > > changes ?
> > > >
> > > > My initial goals are:
> > > > (a) To get AArch64/nvptx offloading to work for above minimal
> test.
> > > > (b) Testsuite results for libgomp on par with x86_64 for non VLA
> > > code (as far as correctness is concerned).
> > > > (c) After (a) and (b) are in place, try to enable support for
> > > offloading with VLA/SVE.
> > > >
> > > > I am planning to address these issues and will post patches for
> > > > the
> > > same shortly. I will be grateful for any feedback or suggestions
> on
> > > how to proceed forward.
> > > >
> > > > Thanks,
> > > > Prathamesh
