zoc: Working basis

Cargo.lock

@@ -98,18 +98,18 @@ checksum = "b048fb63fd8b5923fc5aa7b340d8e156aec7ec02f0c78fa8a6ddc2613f6f71de"
 
 [[package]]
 name = "bpaf"
-version = "0.9.15"
+version = "0.9.17"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "50fd5174866dc2fa2ddc96e8fb800852d37f064f32a45c7b7c2f8fa2c64c77fa"
+checksum = "fd343f000c4472dc4557e093f314c5210375d3e9349ad866fee0ddb7f4ed10bf"
 dependencies = [
  "bpaf_derive",
 ]
 
 [[package]]
 name = "bpaf_derive"
-version = "0.5.13"
+version = "0.5.17"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "cf95d9c7e6aba67f8fc07761091e93254677f4db9e27197adecebc7039a58722"
+checksum = "fefb4feeec9a091705938922f26081aad77c64cd2e76cd1c4a9ece8e42e1618a"
 dependencies = [
  "proc-macro2",
  "quote",
@@ -1218,3 +1218,16 @@ dependencies = [
  "wchar",
  "winapi",
 ]
+
+[[package]]
+name = "zoc"
+version = "0.0.0"
+dependencies = [
+ "amd_comgr-sys",
+ "bpaf",
+ "comgr",
+ "hip_runtime-sys",
+ "ptx",
+ "ptx_parser",
+ "thiserror 1.0.64",
+]

Cargo.toml

@@ -20,6 +20,7 @@ members = [
     "ptx_parser_macros_impl",
     "xtask",
     "zluda_bindgen",
+    "zoc",
 ]
 
 default-members = ["zluda", "zluda_ml", "zluda_inject", "zluda_redirect"]

ptx/src/pass/emit_llvm.rs

@@ -156,6 +156,10 @@ impl std::fmt::Debug for Message {
 }
 
 impl Message {
+    pub fn to_bytes(&self) -> &[u8] {
+        self.0.to_bytes()
+    }
+
     pub fn to_str(&self) -> &str {
         self.0.to_str().unwrap().trim()
     }

zoc/Cargo.toml

@@ -0,0 +1,15 @@
+[package]
+name = "zoc"
+description = "ZLUDA offline compiler"
+version = "0.0.0"
+authors = ["Joëlle van Essen <joelle@v-essen.nl>"]
+edition = "2021"
+
+[dependencies]
+amd_comgr-sys = { path = "../ext/amd_comgr-sys" }
+bpaf = { version = "0.9.17", features = ["derive"] }
+comgr = { path = "../comgr" }
+hip_runtime-sys = { path = "../ext/hip_runtime-sys" }
+ptx = { path = "../ptx" }
+ptx_parser = { path = "../ptx_parser" }
+thiserror = "1.0"

zoc/src/main.rs

@@ -0,0 +1,102 @@
+use std::error::Error;
+use std::ffi::CStr;
+use std::fs::{self, File};
+use std::io::{self, Write};
+use std::mem;
+use std::path::Path;
+use std::str;
+
+use amd_comgr_sys::amd_comgr_status_s;
+use bpaf::Bpaf;
+use hip_runtime_sys::hipErrorCode_t;
+use ptx_parser::PtxError;
+
+#[derive(Debug, Clone, Bpaf)]
+#[bpaf(options, version)]
+pub struct Options {
+    #[bpaf(positional("file"))]
+    /// PTX file
+    file: String,
+}
+
+fn main() -> Result<(), Box<dyn Error>> {
+    let options = options().run();
+    let ptx_path = options.file;
+    let ptx_path = Path::new(&ptx_path);
+
+    let ptx = fs::read(ptx_path)?;
+    let ptx = str::from_utf8(&ptx)?;
+    let llvm_artifacts = ptx_to_llvm(ptx)?;
+    
+    let ll_path = ptx_path.with_extension("ll");
+    write_to_file(&llvm_artifacts.ll, &ll_path)?;
+
+    let elf_module = llvm_to_elf(&llvm_artifacts)?;
+    let elf_path = ptx_path.with_extension("elf");
+    write_to_file(&elf_module, &elf_path)?;
+
+    Ok(())
+}
+
+fn join_ptx_errors(vector: Vec<PtxError>) -> String {
+    let errors: Vec<String> = vector.iter().map(PtxError::to_string).collect();
+    errors.join("\n")
+}
+
+fn ptx_to_llvm(ptx: &str) -> Result<LLVMArtifacts, Box<dyn Error>> {
+    let ast = ptx_parser::parse_module_checked(ptx).map_err(join_ptx_errors)?;
+    let module = ptx::to_llvm_module(ast)?;
+    let bitcode = module.llvm_ir.write_bitcode_to_memory().to_vec();
+    let linked_bitcode = module.linked_bitcode().to_vec();
+    let ll = module.llvm_ir.print_module_to_string().to_bytes().to_vec();
+    Ok(LLVMArtifacts { bitcode, linked_bitcode, ll })
+}
+
+fn llvm_to_elf(module: &LLVMArtifacts) -> Result<Vec<u8>, ElfError> {
+    use hip_runtime_sys::*;
+    unsafe { hipInit(0) }?;
+    let dev = 0;
+    let mut stream = unsafe { mem::zeroed() };
+    unsafe { hipStreamCreate(&mut stream) }?;
+    let mut dev_props = unsafe { mem::zeroed() };
+    unsafe { hipGetDevicePropertiesR0600(&mut dev_props, dev) }?;
+    comgr::compile_bitcode(
+        unsafe { CStr::from_ptr(dev_props.gcnArchName.as_ptr()) },
+        &module.bitcode,
+        &module.linked_bitcode,
+    ).map_err(ElfError::from)
+}
+
+fn write_to_file(content: &[u8], path: &Path) -> io::Result<()> {
+    let mut file = File::create(path)?;
+    file.write_all(content)?;
+    file.flush()?;
+    Ok(())
+}
+
+#[derive(Debug)]
+struct LLVMArtifacts {
+    bitcode: Vec<u8>,
+    linked_bitcode: Vec<u8>,
+    ll: Vec<u8>,
+}
+
+#[derive(Debug, thiserror::Error)]
+enum ElfError {
+    #[error("HIP error: {0:?}")]
+    HipError(hipErrorCode_t),
+    #[error("amd_comgr error: {0:?}")]
+    AmdComgrError(amd_comgr_status_s),
+}
+
+impl From<hipErrorCode_t> for ElfError {
+    fn from(value: hipErrorCode_t) -> Self {
+        ElfError::HipError(value)
+    }
+}
+
+impl From<amd_comgr_status_s> for ElfError {
+    fn from(value: amd_comgr_status_s) -> Self {
+        ElfError::AmdComgrError(value)
+    }
+}