std::offload requires mangled names

[ZuseZ4]

I tried this code:

#![feature(asm_experimental_arch)]
#![allow(internal_features)]
#![feature(abi_gpu_kernel)]
#![feature(link_llvm_intrinsics)]
#![feature(rustc_attrs)]
#![feature(core_intrinsics)]
#![cfg_attr(target_os = "amdhsa", feature(stdarch_amdgpu))]
#![no_std]

#[cfg(not(target_os = "linux"))]
use core::arch::amdgpu::*;

#[cfg(target_os = "linux")]
extern crate libc;

#[panic_handler]
fn panic(_: &core::panic::PanicInfo) -> ! {
    loop {}
}

#[cfg(target_os = "linux")]
#[unsafe(no_mangle)]
#[inline(never)]
unsafe fn main() {
    use libc::printf;

    let mut dbg: [usize; 40] = [0; 40];
    core::hint::black_box(&dbg);
    for i in 0..40 {
        core::intrinsics::offload::<_, _, ()>(
            kernel_1,
            [256, 1, 1],
            [32, 1, 1],
            (&mut dbg,),
        );
    }
    unsafe { printf(c"dbg %lu\n".as_ptr(), dbg[1]) };
    unsafe { printf(c"dbg %lu\n".as_ptr(), dbg[2]) };
}

#[cfg(target_os = "linux")]
unsafe extern "C" {
    pub fn kernel_1(dbg: *mut [usize; 40]);
}

//#[unsafe(no_mangle)]
#[cfg(not(target_os = "linux"))]
#[inline(never)]
#[rustc_offload_kernel]
pub unsafe extern "gpu-kernel" fn kernel_1(dbg: &mut [usize; 40]) {
    let ii = workitem_id_x() as usize;
    if ii < 40 {
        dbg[ii] = ii;
    }
}

I expected to see this happen: compiles and runs

Instead, this happened: compiles but fails at runtime.

Meta

rustc --version --verbose:

main branch

Backtrace

omptarget device 0 info: Entering OpenMP data region with being_mapper at unknown:0:0 with 1 arguments:
omptarget device 0 info: tofrom(unknown)[320] 
"PluginInterface" error: Failure to init kernel: "requested object was not found in the binary image" error in hsa_executable_get_symbol_by_name(kernel_1.kd): HSA_STATUS_ERROR_INVALID_SYMBOL_NAME: There is no symbol with the given name.
omptarget error: Failed to load kernel kernel_1
[..]
"PluginInterface" device 0 info: Copying data from host to device, HstPtr=0x00007fffffff9050, TgtPtr=0x0000154d3e801000, Size=16, Name=KernelLaunchEnv
"PluginInterface" device 0 info: ./steps_dec.txt: line 8: 1310132 Segmentation fault      (core dumped) LIBOMPTARGET_INFO=-1 OFFLOAD_TRACK_ALLOCATION_TRACES=true ./bare

Mangled names work for offload_args, but that's trivial since we're only compiling once. Here we have the challenge that we compile twice, for host and device. But then again we're the compiler, so that should be possible to solve. In the future people will want to use generics to instantiate optimized kernels for different dimensions, and then we'd need mangled names.

cc @Sa4dUs