// Tests the amdgpu_dispatch_ptr intrinsic.

//@ compile-flags: --crate-type=rlib --target amdgcn-amd-amdhsa -Ctarget-cpu=gfx900
//@ needs-llvm-components: amdgpu
//@ add-minicore
#![feature(intrinsics, no_core, rustc_attrs)]
#![no_core]

extern crate minicore;

pub struct DispatchPacket {
    pub header: u16,
    pub setup: u16,
    pub workgroup_size_x: u16, // and more
}

#[rustc_intrinsic]
#[rustc_nounwind]
fn amdgpu_dispatch_ptr() -> *const ();

// CHECK-LABEL: @get_dispatch_data
// CHECK: %[[ORIG_PTR:[^ ]+]] = {{(tail )?}}call ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
// CHECK-NEXT: %[[PTR:[^ ]+]] = addrspacecast ptr addrspace(4) %[[ORIG_PTR]] to ptr
#[unsafe(no_mangle)]
pub fn get_dispatch_data() -> &'static DispatchPacket {
    unsafe { &*(amdgpu_dispatch_ptr() as *const _) }
}