diff --git a/Cargo.lock b/Cargo.lock index e00478e..c140896 100644 --- a/Cargo.lock +++ b/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", +] diff --git a/Cargo.toml b/Cargo.toml index 18fd140..9842781 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -20,6 +20,7 @@ members = [ "ptx_parser_macros_impl", "xtask", "zluda_bindgen", + "zoc", ] default-members = ["zluda", "zluda_ml", "zluda_inject", "zluda_redirect"] diff --git a/ptx/src/pass/emit_llvm.rs b/ptx/src/pass/emit_llvm.rs index 8b43f3e..77c3dcc 100644 --- a/ptx/src/pass/emit_llvm.rs +++ b/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() } diff --git a/zoc/Cargo.toml b/zoc/Cargo.toml new file mode 100644 index 0000000..917d498 --- /dev/null +++ b/zoc/Cargo.toml @@ -0,0 +1,15 @@ +[package] +name = "zoc" +description = "ZLUDA offline compiler" +version = "0.0.0" +authors = ["Joƫlle van Essen "] +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" \ No newline at end of file diff --git a/zoc/src/main.rs b/zoc/src/main.rs new file mode 100644 index 0000000..1742eab --- /dev/null +++ b/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> { + 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) -> String { + let errors: Vec = vector.iter().map(PtxError::to_string).collect(); + errors.join("\n") +} + +fn ptx_to_llvm(ptx: &str) -> Result> { + 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, 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, + linked_bitcode: Vec, + ll: Vec, +} + +#[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 for ElfError { + fn from(value: hipErrorCode_t) -> Self { + ElfError::HipError(value) + } +} + +impl From for ElfError { + fn from(value: amd_comgr_status_s) -> Self { + ElfError::AmdComgrError(value) + } +} \ No newline at end of file