https://gcc.gnu.org/bugzilla/show_bug.cgi?id=81689
--- Comment #2 from Tobias Burnus <burnus at gcc dot gnu.org> --- 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.
