| // 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 _) } |
| } |
