Wrap comgr in a separate library, move bindgen output to ext

Cargo.toml

@@ -3,10 +3,12 @@
 resolver = "2"
 
 members = [
+    "ext/hip_runtime-sys",
+    "ext/amd_comgr-sys",
+    "comgr",
     "cuda_base",
     "cuda_types",
     "detours-sys",
-    "hip_runtime-sys",
     "level_zero-sys",
     "level_zero",
     "spirv_tools-sys",

comgr/Cargo.toml

@@ -0,0 +1,10 @@
+[package]
+name = "comgr"
+version = "0.0.0"
+authors = ["Andrzej Janik <vosen@vosen.pl>"]
+edition = "2021"
+
+[lib]
+
+[dependencies]
+amd_comgr-sys = { path = "../ext/amd_comgr-sys" }

comgr/src/lib.rs

@@ -0,0 +1,125 @@
+use amd_comgr_sys::*;
+use std::{ffi::CStr, mem, ptr};
+
+struct Data(amd_comgr_data_t);
+
+impl Data {
+    fn new(
+        kind: amd_comgr_data_kind_t,
+        name: &CStr,
+        content: &[u8],
+    ) -> Result<Self, amd_comgr_status_s> {
+        let mut data = unsafe { mem::zeroed() };
+        unsafe { amd_comgr_create_data(kind, &mut data) }?;
+        unsafe { amd_comgr_set_data_name(data, name.as_ptr()) }?;
+        unsafe { amd_comgr_set_data(data, content.len(), content.as_ptr().cast()) }?;
+        Ok(Self(data))
+    }
+
+    fn get(&self) -> amd_comgr_data_t {
+        self.0
+    }
+
+    fn copy_content(&self) -> Result<Vec<u8>, amd_comgr_status_s> {
+        let mut size = unsafe { mem::zeroed() };
+        unsafe { amd_comgr_get_data(self.get(), &mut size, ptr::null_mut()) }?;
+        let mut result: Vec<u8> = Vec::with_capacity(size);
+        unsafe { result.set_len(size) };
+        unsafe { amd_comgr_get_data(self.get(), &mut size, result.as_mut_ptr().cast()) }?;
+        Ok(result)
+    }
+}
+
+struct DataSet(amd_comgr_data_set_t);
+
+impl DataSet {
+    fn new() -> Result<Self, amd_comgr_status_s> {
+        let mut data_set = unsafe { mem::zeroed() };
+        unsafe { amd_comgr_create_data_set(&mut data_set) }?;
+        Ok(Self(data_set))
+    }
+
+    fn add(&self, data: &Data) -> Result<(), amd_comgr_status_s> {
+        unsafe { amd_comgr_data_set_add(self.get(), data.get()) }
+    }
+
+    fn get(&self) -> amd_comgr_data_set_t {
+        self.0
+    }
+
+    fn get_data(
+        &self,
+        kind: amd_comgr_data_kind_t,
+        index: usize,
+    ) -> Result<Data, amd_comgr_status_s> {
+        let mut data = unsafe { mem::zeroed() };
+        unsafe { amd_comgr_action_data_get_data(self.get(), kind, index, &mut data) }?;
+        Ok(Data(data))
+    }
+}
+
+impl Drop for DataSet {
+    fn drop(&mut self) {
+        unsafe { amd_comgr_destroy_data_set(self.get()).ok() };
+    }
+}
+
+struct ActionInfo(amd_comgr_action_info_t);
+
+impl ActionInfo {
+    fn new() -> Result<Self, amd_comgr_status_s> {
+        let mut action = unsafe { mem::zeroed() };
+        unsafe { amd_comgr_create_action_info(&mut action) }?;
+        Ok(Self(action))
+    }
+
+    fn set_isa_name(&self, isa: &CStr) -> Result<(), amd_comgr_status_s> {
+        let mut full_isa = "amdgcn-amd-amdhsa--".to_string().into_bytes();
+        full_isa.extend(isa.to_bytes_with_nul());
+        unsafe { amd_comgr_action_info_set_isa_name(self.get(), full_isa.as_ptr().cast()) }
+    }
+
+    fn get(&self) -> amd_comgr_action_info_t {
+        self.0
+    }
+}
+
+impl Drop for ActionInfo {
+    fn drop(&mut self) {
+        unsafe { amd_comgr_destroy_action_info(self.get()).ok() };
+    }
+}
+
+pub fn compile_bitcode(gcn_arch: &CStr, buffer: &[u8]) -> Result<Vec<u8>, amd_comgr_status_s> {
+    use amd_comgr_sys::*;
+    let bitcode_data_set = DataSet::new()?;
+    let bitcode_data = Data::new(
+        amd_comgr_data_kind_t::AMD_COMGR_DATA_KIND_BC,
+        c"zluda.bc",
+        buffer,
+    )?;
+    bitcode_data_set.add(&bitcode_data)?;
+    let reloc_data_set = DataSet::new()?;
+    let action_info = ActionInfo::new()?;
+    action_info.set_isa_name(gcn_arch)?;
+    unsafe {
+        amd_comgr_do_action(
+            amd_comgr_action_kind_t::AMD_COMGR_ACTION_CODEGEN_BC_TO_RELOCATABLE,
+            action_info.get(),
+            bitcode_data_set.get(),
+            reloc_data_set.get(),
+        )
+    }?;
+    let exec_data_set = DataSet::new()?;
+    unsafe {
+        amd_comgr_do_action(
+            amd_comgr_action_kind_t::AMD_COMGR_ACTION_LINK_RELOCATABLE_TO_EXECUTABLE,
+            action_info.get(),
+            reloc_data_set.get(),
+            exec_data_set.get(),
+        )
+    }?;
+    let executable =
+        exec_data_set.get_data(amd_comgr_data_kind_t::AMD_COMGR_DATA_KIND_EXECUTABLE, 0)?;
+    executable.copy_content()
+}

ext/amd_comgr-sys/src/amd_comgr.rs

@@ -27,6 +27,10 @@ impl amd_comgr_status_s {
 #[derive(Copy, Clone, Hash, PartialEq, Eq, Debug)]
 pub struct amd_comgr_status_s(pub ::std::num::NonZeroU32);
 type amd_comgr_status_t = Result<(), self::amd_comgr_status_s>;
+// Size check
+const _: fn() = || {
+    let _ = std::mem::transmute::<amd_comgr_status_t, u32>;
+};
 impl amd_comgr_language_s {
     #[doc = " No high level language."]
     pub const AMD_COMGR_LANGUAGE_NONE: amd_comgr_language_s = amd_comgr_language_s(0);

ext/hip_runtime-sys/src/hip_runtime_api.rs

@@ -994,7 +994,10 @@ pub struct hipErrorCode_t(pub ::std::num::NonZeroU32);
 #[doc = " HIP error type\n"]
 #[must_use]
 pub type hipError_t = Result<(), hipErrorCode_t>;
-
+// Size check
+const _: fn() = || {
+    let _ = std::mem::transmute::<hipError_t, u32>;
+};
 impl hipDeviceAttribute_t {
     pub const hipDeviceAttributeCudaCompatibleBegin: hipDeviceAttribute_t = hipDeviceAttribute_t(0);
 }

ptx/Cargo.toml

@@ -27,8 +27,8 @@ version = "0.19.12"
 features = ["lexer"]
 
 [dev-dependencies]
-hip_runtime-sys = { path = "../hip_runtime-sys" }
-amd_comgr-sys = { path = "../amd_comgr-sys" }
+hip_runtime-sys = { path = "../ext/hip_runtime-sys" }
+comgr = { path = "../comgr" }
 spirv_tools-sys = { path = "../spirv_tools-sys" }
 tempfile = "3"
 paste = "1.0"

ptx/src/test/spirv_run/mod.rs

@@ -323,7 +323,11 @@ fn run_hip<Input: From<u8> + Copy + Debug, Output: From<u8> + Copy + Debug + Def
         unsafe { hipStreamCreate(&mut stream) }.unwrap();
         let mut dev_props = unsafe { mem::zeroed() };
         unsafe { hipGetDevicePropertiesR0600(&mut dev_props, dev) }.unwrap();
-        let elf_module = unsafe { compile_amd(&module.llvm_ir) };
+        let elf_module = comgr::compile_bitcode(
+            unsafe { CStr::from_ptr(dev_props.gcnArchName.as_ptr()) },
+            &*module.llvm_ir,
+        )
+        .unwrap();
         let mut module = ptr::null_mut();
         unsafe { hipModuleLoadData(&mut module, elf_module.as_ptr() as _) }.unwrap();
         let mut kernel = ptr::null_mut();
@@ -378,54 +382,6 @@ fn run_hip<Input: From<u8> + Copy + Debug, Output: From<u8> + Copy + Debug + Def
     Ok(result)
 }
 
-unsafe fn compile_amd(buffer: &pass::emit_llvm::MemoryBuffer) -> Vec<u8> {
-    use amd_comgr_sys::*;
-    let mut data_set = mem::zeroed();
-    amd_comgr_create_data_set(&mut data_set).unwrap();
-    let mut data = mem::zeroed();
-    amd_comgr_create_data(amd_comgr_data_kind_t::AMD_COMGR_DATA_KIND_BC, &mut data).unwrap();
-    let buffer = &**buffer;
-    amd_comgr_set_data(data, buffer.len(), buffer.as_ptr().cast()).unwrap();
-    amd_comgr_set_data_name(data, c"zluda.bc".as_ptr()).unwrap();
-    amd_comgr_data_set_add(data_set, data).unwrap();
-    let mut reloc_data = mem::zeroed();
-    amd_comgr_create_data_set(&mut reloc_data).unwrap();
-    let mut action_info = mem::zeroed();
-    amd_comgr_create_action_info(&mut action_info).unwrap();
-    amd_comgr_action_info_set_isa_name(action_info, c"amdgcn-amd-amdhsa--gfx1030".as_ptr())
-        .unwrap();
-    amd_comgr_do_action(
-        amd_comgr_action_kind_t::AMD_COMGR_ACTION_CODEGEN_BC_TO_RELOCATABLE,
-        action_info,
-        data_set,
-        reloc_data,
-    )
-    .unwrap();
-    let mut exec_data = mem::zeroed();
-    amd_comgr_create_data_set(&mut exec_data).unwrap();
-    amd_comgr_do_action(
-        amd_comgr_action_kind_t::AMD_COMGR_ACTION_LINK_RELOCATABLE_TO_EXECUTABLE,
-        action_info,
-        reloc_data,
-        exec_data,
-    )
-    .unwrap();
-    let mut exec = mem::zeroed();
-    amd_comgr_action_data_get_data(
-        exec_data,
-        amd_comgr_data_kind_t::AMD_COMGR_DATA_KIND_EXECUTABLE,
-        0,
-        &mut exec,
-    )
-    .unwrap();
-    let mut size = mem::zeroed();
-    amd_comgr_get_data(exec, &mut size, ptr::null_mut()).unwrap();
-    let mut result: Vec<u8> = Vec::with_capacity(size);
-    result.set_len(size);
-    amd_comgr_get_data(exec, &mut size, result.as_mut_ptr().cast()).unwrap();
-    result
-}
-
 struct EqMap<T>
 where
     T: Eq + Copy + Hash,

zluda/Cargo.toml

@@ -9,7 +9,7 @@ name = "zluda"
 
 [dependencies]
 ptx = { path = "../ptx" }
-hip_runtime-sys = { path = "../hip_runtime-sys" }
+hip_runtime-sys = { path = "../ext/hip_runtime-sys" }
 lazy_static = "1.4"
 num_enum = "0.4"
 lz4-sys = "1.9"