The error looks like: ... ptxas /tmp/ccOmk7fK.o, line 102; error : State space mismatch between instruction and address in instruction 'ld'^M ptxas /tmp/ccOmk7fK.o, line 121; error : State space mismatch between instruction and address in instruction 'ld'^M ptxas /tmp/ccOmk7fK.o, line 126; error : State space mismatch between instruction and address in instruction 'ld'^M ptxas /tmp/ccOmk7fK.o, line 231; error : State space mismatch between instruction and address in instruction 'ld'^M ptxas /tmp/ccOmk7fK.o, line 252; error : State space mismatch between instruction and address in instruction 'ld'^M ptxas /tmp/ccOmk7fK.o, line 289; error : State space mismatch between instruction and address in instruction 'ld'^M ptxas /tmp/ccOmk7fK.o, line 294; error : State space mismatch between instruction and address in instruction 'ld'^M ptxas /tmp/ccOmk7fK.o, line 301; error : State space mismatch between instruction and address in instruction 'ld'^M ptxas /tmp/ccOmk7fK.o, line 102; error : Unknown symbol 'a_linkptr'^M ptxas /tmp/ccOmk7fK.o, line 102; fatal : Label expected for forward reference of 'a_linkptr'^M ptxas fatal : Ptx assembly aborted due to errors^M nvptx-as: ptxas returned 255 exit status^M mkoffload: fatal error: x86_64-none-linux-gnu-accel-nvptx-none-gcc returned 1 exit status^M compilation terminated.^M ... The problem is that #pragma omp target link is not implemented for nvptx, as mentioned before here ( https://gcc.gnu.org/ml/gcc-patches/2016-11/msg00972.html ): ... With OpenMP/PTX offloading there are 5 additional failures in check-target-libgomp: <SNIP> One with 'target link' (not implemented) FAIL: libgomp.c/target-link-1.c (test for excess errors) ...
I believe the compile-time issue has been fixed by PR 94233. There is a run-time issue which is tracked by PR 94251. NOTE: This test is currently disabled (since r10-6542-gfd789c816b06235b04698636db69e302b24c83ba ) due to the fails.
While the two PRs have fixed offloading for AMDGCN, nvptx still fails with the same error. The generated assembler is: "ld.global.u64 %r27,[a$linkptr];\n" But there is NO associated global-var declaration similar to: "// BEGIN GLOBAL VAR DECL: __nvptx_uni\n" ".extern .shared .u32 __nvptx_uni[32];\n" In principle, it should be generated by nvptx_output_aligned_decl – but this never called. It is currently called via ASM_OUTPUT_ALIGNED_DECL_COMMON and ASM_OUTPUT_ALIGNED_DECL_LOCAL. Without proper debugging, I have the feeling that ASM_OUTPUT_ALIGNED_BSS is called for this – but that's not defined for nvptx. Nor is TARGET_ASM_SELECT_SECTION.
The master branch has been updated by Tobias Burnus <burnus@gcc.gnu.org>: https://gcc.gnu.org/g:c2211a60ff05b7a0289d3e287e72c181bb4d5d8b commit r10-7354-gc2211a60ff05b7a0289d3e287e72c181bb4d5d8b Author: Tobias Burnus <tobias@codesourcery.com> Date: Tue Mar 24 15:13:56 2020 +0100 Fix OpenMP offload handling for target-link variables for nvptx (PR81689) PR libgomp/81689 * lto.c (offload_handle_link_vars): Propagate TREE_PUBLIC state. PR libgomp/81689 * omp-offload.c (omp_finish_file): Fix target-link handling if targetm_common.have_named_sections is false. PR libgomp/81689 * testsuite/libgomp.c/target-link-1.c: Remove xfail.
FIXED.