On Sun, Aug 11, 2024 at 11:36 PM Prathamesh Kulkarni via Gcc <gcc@gcc.gnu.org> wrote: > > > > > -----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.
Just a quick note, changing it to Save from Optimization will also break gcc.target/aarch64/sve/target_optimization-1.c (which was added with r15-2344-g8a5f528fba788f) . > > FWIW, x86_64 cl_optimization_stream_out also similarly streams out > ix86_unroll_only_small_loops: Note that is also a recent change (little over 3 weeks ago, r15-2430-ga59c4e496fa916) and I suspect nobody tested offloading after that change; it was not backported to the branches either because it was known to need to bump the LTO minor version there. Thanks, Andrew Pinski > > 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
