From 63b785872009f9cc29e53114ff9c86e51477e330 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Jo=C3=ABlle=20van=20Essen?= Date: Mon, 3 Mar 2025 13:00:09 +0100 Subject: [PATCH 01/15] zoc: Working basis --- Cargo.lock | 21 ++++++-- Cargo.toml | 1 + ptx/src/pass/emit_llvm.rs | 4 ++ zoc/Cargo.toml | 15 ++++++ zoc/src/main.rs | 102 ++++++++++++++++++++++++++++++++++++++ 5 files changed, 139 insertions(+), 4 deletions(-) create mode 100644 zoc/Cargo.toml create mode 100644 zoc/src/main.rs 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 From 8e3a15ec3b43ee0cb9976f6951f8ae8b0aa94d54 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Jo=C3=ABlle=20van=20Essen?= Date: Mon, 3 Mar 2025 13:21:40 +0100 Subject: [PATCH 02/15] zoc: Refactoring and fix build --- Cargo.toml | 2 +- zoc/src/main.rs | 12 +++++------- 2 files changed, 6 insertions(+), 8 deletions(-) diff --git a/Cargo.toml b/Cargo.toml index 9842781..04c215b 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -23,7 +23,7 @@ members = [ "zoc", ] -default-members = ["zluda", "zluda_ml", "zluda_inject", "zluda_redirect"] +default-members = ["zluda", "zluda_ml", "zluda_inject", "zluda_redirect", "zoc"] [profile.release-lto] inherits = "release" diff --git a/zoc/src/main.rs b/zoc/src/main.rs index 1742eab..c20730a 100644 --- a/zoc/src/main.rs +++ b/zoc/src/main.rs @@ -2,7 +2,7 @@ use std::error::Error; use std::ffi::CStr; use std::fs::{self, File}; use std::io::{self, Write}; -use std::mem; +use std::mem::MaybeUninit; use std::path::Path; use std::str; @@ -54,12 +54,10 @@ fn ptx_to_llvm(ptx: &str) -> Result> { 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) }?; + let mut dev_props: MaybeUninit = MaybeUninit::uninit(); + unsafe { hipGetDevicePropertiesR0600(dev_props.as_mut_ptr(), 0) }?; + let dev_props = unsafe { dev_props.assume_init() }; + comgr::compile_bitcode( unsafe { CStr::from_ptr(dev_props.gcnArchName.as_ptr()) }, &module.bitcode, From 8c72eae6baee55f3430d727d151f791386589442 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Jo=C3=ABlle=20van=20Essen?= Date: Mon, 3 Mar 2025 13:38:47 +0100 Subject: [PATCH 03/15] zoc: Refactoring --- zoc/src/main.rs | 18 ++++++++++-------- 1 file changed, 10 insertions(+), 8 deletions(-) diff --git a/zoc/src/main.rs b/zoc/src/main.rs index c20730a..9b7d7eb 100644 --- a/zoc/src/main.rs +++ b/zoc/src/main.rs @@ -26,14 +26,14 @@ fn main() -> Result<(), Box> { let ptx = fs::read(ptx_path)?; let ptx = str::from_utf8(&ptx)?; - let llvm_artifacts = ptx_to_llvm(ptx)?; + let llvm = ptx_to_llvm(ptx)?; let ll_path = ptx_path.with_extension("ll"); - write_to_file(&llvm_artifacts.ll, &ll_path)?; + write_to_file(&llvm.llvm_ir, &ll_path)?; - let elf_module = llvm_to_elf(&llvm_artifacts)?; + let elf = llvm_to_elf(&llvm)?; let elf_path = ptx_path.with_extension("elf"); - write_to_file(&elf_module, &elf_path)?; + write_to_file(&elf, &elf_path)?; Ok(()) } @@ -48,18 +48,20 @@ fn ptx_to_llvm(ptx: &str) -> Result> { 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 }) + let llvm_ir = module.llvm_ir.print_module_to_string().to_bytes().to_vec(); + Ok(LLVMArtifacts { bitcode, linked_bitcode, llvm_ir }) } fn llvm_to_elf(module: &LLVMArtifacts) -> Result, ElfError> { use hip_runtime_sys::*; + unsafe { hipInit(0) }?; let mut dev_props: MaybeUninit = MaybeUninit::uninit(); unsafe { hipGetDevicePropertiesR0600(dev_props.as_mut_ptr(), 0) }?; let dev_props = unsafe { dev_props.assume_init() }; + let gcn_arch = unsafe { CStr::from_ptr(dev_props.gcnArchName.as_ptr()) }; comgr::compile_bitcode( - unsafe { CStr::from_ptr(dev_props.gcnArchName.as_ptr()) }, + gcn_arch, &module.bitcode, &module.linked_bitcode, ).map_err(ElfError::from) @@ -76,7 +78,7 @@ fn write_to_file(content: &[u8], path: &Path) -> io::Result<()> { struct LLVMArtifacts { bitcode: Vec, linked_bitcode: Vec, - ll: Vec, + llvm_ir: Vec, } #[derive(Debug, thiserror::Error)] From c6f240e78d72ffae7ae9cdce5aaaeff52fea6535 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Jo=C3=ABlle=20van=20Essen?= Date: Mon, 3 Mar 2025 13:41:19 +0100 Subject: [PATCH 04/15] zoc: Refactoring --- zoc/src/main.rs | 8 ++------ 1 file changed, 2 insertions(+), 6 deletions(-) diff --git a/zoc/src/main.rs b/zoc/src/main.rs index 9b7d7eb..6d7b339 100644 --- a/zoc/src/main.rs +++ b/zoc/src/main.rs @@ -52,7 +52,7 @@ fn ptx_to_llvm(ptx: &str) -> Result> { Ok(LLVMArtifacts { bitcode, linked_bitcode, llvm_ir }) } -fn llvm_to_elf(module: &LLVMArtifacts) -> Result, ElfError> { +fn llvm_to_elf(llvm: &LLVMArtifacts) -> Result, ElfError> { use hip_runtime_sys::*; unsafe { hipInit(0) }?; let mut dev_props: MaybeUninit = MaybeUninit::uninit(); @@ -60,11 +60,7 @@ fn llvm_to_elf(module: &LLVMArtifacts) -> Result, ElfError> { let dev_props = unsafe { dev_props.assume_init() }; let gcn_arch = unsafe { CStr::from_ptr(dev_props.gcnArchName.as_ptr()) }; - comgr::compile_bitcode( - gcn_arch, - &module.bitcode, - &module.linked_bitcode, - ).map_err(ElfError::from) + comgr::compile_bitcode(gcn_arch, &llvm.bitcode, &llvm.linked_bitcode).map_err(ElfError::from) } fn write_to_file(content: &[u8], path: &Path) -> io::Result<()> { From 084339e1414a6e2eba2cf9871b6fc3f33d0903c2 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Jo=C3=ABlle=20van=20Essen?= Date: Thu, 13 Mar 2025 17:12:43 +0100 Subject: [PATCH 05/15] zoc: Specify output type --- Cargo.lock | 18 +++--- zoc/Cargo.toml | 4 +- zoc/src/main.rs | 148 +++++++++++++++++++++++++++++++++++++++--------- 3 files changed, 133 insertions(+), 37 deletions(-) diff --git a/Cargo.lock b/Cargo.lock index c140896..3c809f9 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -98,9 +98,9 @@ checksum = "b048fb63fd8b5923fc5aa7b340d8e156aec7ec02f0c78fa8a6ddc2613f6f71de" [[package]] name = "bpaf" -version = "0.9.17" +version = "0.9.18" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "fd343f000c4472dc4557e093f314c5210375d3e9349ad866fee0ddb7f4ed10bf" +checksum = "4de4d74c5891642753c67ab88f58d971a68dd98673b69689a8c24ce3ec78a412" dependencies = [ "bpaf_derive", ] @@ -151,7 +151,7 @@ dependencies = [ "semver", "serde", "serde_json", - "thiserror 2.0.11", + "thiserror 2.0.12", ] [[package]] @@ -924,11 +924,11 @@ dependencies = [ [[package]] name = "thiserror" -version = "2.0.11" +version = "2.0.12" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "d452f284b73e6d76dd36758a0c8684b1d5be31f92b89d07fd5822175732206fc" +checksum = "567b8a2dae586314f7be2a752ec7474332959c6460e02bde30d702a66d488708" dependencies = [ - "thiserror-impl 2.0.11", + "thiserror-impl 2.0.12", ] [[package]] @@ -944,9 +944,9 @@ dependencies = [ [[package]] name = "thiserror-impl" -version = "2.0.11" +version = "2.0.12" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "26afc1baea8a989337eeb52b6e72a039780ce45c3edfcc9c5b9d112feeb173c2" +checksum = "7f7cf42b4507d8ea322120659672cf1b9dbb93f8f2d4ecfd6e51350ff5b17a1d" dependencies = [ "proc-macro2", "quote", @@ -1229,5 +1229,5 @@ dependencies = [ "hip_runtime-sys", "ptx", "ptx_parser", - "thiserror 1.0.64", + "thiserror 2.0.12", ] diff --git a/zoc/Cargo.toml b/zoc/Cargo.toml index 917d498..146252b 100644 --- a/zoc/Cargo.toml +++ b/zoc/Cargo.toml @@ -7,9 +7,9 @@ edition = "2021" [dependencies] amd_comgr-sys = { path = "../ext/amd_comgr-sys" } -bpaf = { version = "0.9.17", features = ["derive"] } +bpaf = { version = "0.9.18", 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 +thiserror = "2.0.12" \ No newline at end of file diff --git a/zoc/src/main.rs b/zoc/src/main.rs index 6d7b339..a07aadd 100644 --- a/zoc/src/main.rs +++ b/zoc/src/main.rs @@ -1,10 +1,11 @@ +use std::env; use std::error::Error; -use std::ffi::CStr; +use std::ffi::{CStr, OsStr}; use std::fs::{self, File}; use std::io::{self, Write}; use std::mem::MaybeUninit; -use std::path::Path; -use std::str; +use std::path::{Path, PathBuf}; +use std::str::{self, FromStr}; use amd_comgr_sys::amd_comgr_status_s; use bpaf::Bpaf; @@ -14,42 +15,71 @@ use ptx_parser::PtxError; #[derive(Debug, Clone, Bpaf)] #[bpaf(options, version)] pub struct Options { - #[bpaf(positional("file"))] + #[bpaf(external(output_type), optional)] + output_type: Option, + + #[bpaf(positional("filename"))] /// PTX file - file: String, + ptx_path: String, } fn main() -> Result<(), Box> { - let options = options().run(); - let ptx_path = options.file; - let ptx_path = Path::new(&ptx_path); + let opts = options().run(); - let ptx = fs::read(ptx_path)?; + let output_type = match opts.output_type { + Some(t) => t, + None => OutputType::Elf, + }; + + match output_type { + OutputType::LlvmIrLinked | OutputType::Assembly => todo!(), + _ => {} + } + + let ptx_path = Path::new(&opts.ptx_path).to_path_buf(); + check_path(&ptx_path)?; + + let output_path = get_output_path(&ptx_path, &output_type)?; + check_path(&output_path)?; + + let ptx = fs::read(&ptx_path)?; let ptx = str::from_utf8(&ptx)?; let llvm = ptx_to_llvm(ptx)?; - - let ll_path = ptx_path.with_extension("ll"); - write_to_file(&llvm.llvm_ir, &ll_path)?; + + if output_type == OutputType::LlvmIrPreLinked { + write_to_file(&llvm.llvm_ir, &output_path)?; + return Ok(()); + } let elf = llvm_to_elf(&llvm)?; - let elf_path = ptx_path.with_extension("elf"); - write_to_file(&elf, &elf_path)?; + write_to_file(&elf, &output_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 llvm_ir = module.llvm_ir.print_module_to_string().to_bytes().to_vec(); - Ok(LLVMArtifacts { bitcode, linked_bitcode, llvm_ir }) + Ok(LLVMArtifacts { + bitcode, + linked_bitcode, + llvm_ir, + }) +} + +#[derive(Debug)] +struct LLVMArtifacts { + bitcode: Vec, + linked_bitcode: Vec, + llvm_ir: Vec, +} + +fn join_ptx_errors(vector: Vec) -> String { + let errors: Vec = vector.iter().map(PtxError::to_string).collect(); + errors.join("\n") } fn llvm_to_elf(llvm: &LLVMArtifacts) -> Result, ElfError> { @@ -63,20 +93,86 @@ fn llvm_to_elf(llvm: &LLVMArtifacts) -> Result, ElfError> { comgr::compile_bitcode(gcn_arch, &llvm.bitcode, &llvm.linked_bitcode).map_err(ElfError::from) } +fn check_path(path: &Path) -> Result<(), Box> { + if path.try_exists()? && !path.is_file() { + let error = CheckPathError(path.to_path_buf()); + let error = Box::new(error); + return Err(error); + } + Ok(()) +} + +fn get_output_path( + ptx_path: &PathBuf, + output_type: &OutputType, +) -> Result> { + let current_dir = env::current_dir()?; + let output_path = current_dir.join( + ptx_path + .as_path() + .file_stem() + .unwrap_or(OsStr::new("output")), + ); + let output_path = output_path.with_extension(output_type.extension()); + Ok(output_path) +} + fn write_to_file(content: &[u8], path: &Path) -> io::Result<()> { let mut file = File::create(path)?; file.write_all(content)?; file.flush()?; + println!("Wrote to {}", path.to_str().unwrap()); Ok(()) } -#[derive(Debug)] -struct LLVMArtifacts { - bitcode: Vec, - linked_bitcode: Vec, - llvm_ir: Vec, +#[derive(Bpaf, Clone, Copy, Debug, PartialEq)] +enum OutputType { + /// Produce pre-linked LLVM IR + #[bpaf(long("ll"))] + LlvmIrPreLinked, + /// Produce linked LLVM IR + #[bpaf(long("linked-ll"))] + LlvmIrLinked, + /// Produce ELF binary (default) + Elf, + /// Produce assembly + #[bpaf(long("asm"))] + Assembly, } +impl OutputType { + fn extension(self) -> String { + match self { + OutputType::LlvmIrPreLinked | OutputType::LlvmIrLinked => "ll", + OutputType::Assembly => "asm", + OutputType::Elf => "elf", + } + .into() + } +} + +impl FromStr for OutputType { + type Err = ParseOutputTypeError; + + fn from_str(s: &str) -> Result { + match s { + "ll" => Ok(Self::LlvmIrPreLinked), + "ll_linked" => Ok(Self::LlvmIrLinked), + "elf" => Ok(Self::Elf), + "asm" => Ok(Self::Assembly), + _ => Err(ParseOutputTypeError(s.into())), + } + } +} + +#[derive(Debug, thiserror::Error)] +#[error("Not a regular file: {0}")] +struct CheckPathError(PathBuf); + +#[derive(Debug, thiserror::Error)] +#[error("Invalid output type: {0}")] +struct ParseOutputTypeError(String); + #[derive(Debug, thiserror::Error)] enum ElfError { #[error("HIP error: {0:?}")] @@ -95,4 +191,4 @@ impl From for ElfError { fn from(value: amd_comgr_status_s) -> Self { ElfError::AmdComgrError(value) } -} \ No newline at end of file +} From 87b30a660494a92ba616294f14a56c2f33e59b9e Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Jo=C3=ABlle=20van=20Essen?= Date: Thu, 13 Mar 2025 17:27:41 +0100 Subject: [PATCH 06/15] zoc: Default output type --- zoc/src/main.rs | 8 +++----- 1 file changed, 3 insertions(+), 5 deletions(-) diff --git a/zoc/src/main.rs b/zoc/src/main.rs index a07aadd..0f7b91e 100644 --- a/zoc/src/main.rs +++ b/zoc/src/main.rs @@ -26,10 +26,7 @@ pub struct Options { fn main() -> Result<(), Box> { let opts = options().run(); - let output_type = match opts.output_type { - Some(t) => t, - None => OutputType::Elf, - }; + let output_type = opts.output_type.unwrap_or_default(); match output_type { OutputType::LlvmIrLinked | OutputType::Assembly => todo!(), @@ -125,7 +122,7 @@ fn write_to_file(content: &[u8], path: &Path) -> io::Result<()> { Ok(()) } -#[derive(Bpaf, Clone, Copy, Debug, PartialEq)] +#[derive(Bpaf, Clone, Copy, Debug, Default, PartialEq)] enum OutputType { /// Produce pre-linked LLVM IR #[bpaf(long("ll"))] @@ -134,6 +131,7 @@ enum OutputType { #[bpaf(long("linked-ll"))] LlvmIrLinked, /// Produce ELF binary (default) + #[default] Elf, /// Produce assembly #[bpaf(long("asm"))] From 06328891c465299590168c33ced97f7d40d25586 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Jo=C3=ABlle=20van=20Essen?= Date: Wed, 26 Mar 2025 19:53:26 +0100 Subject: [PATCH 07/15] zoc: Linked LLVM IR --- comgr/src/lib.rs | 31 +++++++++----- ptx/src/lib.rs | 4 +- ptx/src/pass/emit_llvm.rs | 48 ++++++++++++++++++++- zoc/src/error.rs | 66 +++++++++++++++++++++++++++++ zoc/src/main.rs | 89 +++++++++++++++------------------------ 5 files changed, 170 insertions(+), 68 deletions(-) create mode 100644 zoc/src/error.rs diff --git a/comgr/src/lib.rs b/comgr/src/lib.rs index ac407ef..e1a872b 100644 --- a/comgr/src/lib.rs +++ b/comgr/src/lib.rs @@ -1,7 +1,9 @@ use amd_comgr_sys::*; -use std::{ffi::CStr, mem, ptr}; +use std::ffi::CStr; +use std::mem; +use std::ptr; -struct Data(amd_comgr_data_t); +pub struct Data(amd_comgr_data_t); impl Data { fn new( @@ -20,7 +22,7 @@ impl Data { self.0 } - fn copy_content(&self) -> Result, amd_comgr_status_s> { + pub fn copy_content(&self) -> Result, 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 = Vec::with_capacity(size); @@ -30,7 +32,7 @@ impl Data { } } -struct DataSet(amd_comgr_data_set_t); +pub struct DataSet(amd_comgr_data_set_t); impl DataSet { fn new() -> Result { @@ -47,7 +49,7 @@ impl DataSet { self.0 } - fn get_data( + pub fn get_data( &self, kind: amd_comgr_data_kind_t, index: usize, @@ -108,11 +110,10 @@ impl Drop for ActionInfo { } } -pub fn compile_bitcode( - gcn_arch: &CStr, +pub fn link_bitcode( main_buffer: &[u8], ptx_impl: &[u8], -) -> Result, amd_comgr_status_s> { +) -> Result { use amd_comgr_sys::*; let bitcode_data_set = DataSet::new()?; let main_bitcode_data = Data::new( @@ -128,11 +129,21 @@ pub fn compile_bitcode( )?; bitcode_data_set.add(&stdlib_bitcode_data)?; let linking_info = ActionInfo::new()?; - let linked_data_set = do_action( + do_action( &bitcode_data_set, &linking_info, amd_comgr_action_kind_t::AMD_COMGR_ACTION_LINK_BC_TO_BC, - )?; + ) +} + +pub fn compile_bitcode( + gcn_arch: &CStr, + main_buffer: &[u8], + ptx_impl: &[u8], +) -> Result, amd_comgr_status_s> { + use amd_comgr_sys::*; + + let linked_data_set = link_bitcode(main_buffer, ptx_impl)?; let compile_to_exec = ActionInfo::new()?; compile_to_exec.set_isa_name(gcn_arch)?; compile_to_exec.set_language(amd_comgr_language_t::AMD_COMGR_LANGUAGE_LLVM_IR)?; diff --git a/ptx/src/lib.rs b/ptx/src/lib.rs index da972f6..aea554e 100644 --- a/ptx/src/lib.rs +++ b/ptx/src/lib.rs @@ -2,5 +2,5 @@ pub(crate) mod pass; #[cfg(test)] mod test; -pub use pass::to_llvm_module; - +pub use pass::{TranslateError, to_llvm_module}; +pub use pass::emit_llvm::bitcode_to_ir; \ No newline at end of file diff --git a/ptx/src/pass/emit_llvm.rs b/ptx/src/pass/emit_llvm.rs index 072903b..2b6b68b 100644 --- a/ptx/src/pass/emit_llvm.rs +++ b/ptx/src/pass/emit_llvm.rs @@ -27,11 +27,13 @@ use std::array::TryFromSliceError; use std::convert::TryInto; use std::ffi::{CStr, NulError}; +use std::mem::MaybeUninit; use std::ops::Deref; -use std::{i8, ptr}; +use std::ptr; use super::*; use llvm_zluda::analysis::{LLVMVerifierFailureAction, LLVMVerifyModule}; +use llvm_zluda::bit_reader::LLVMParseBitcodeInContext2; use llvm_zluda::bit_writer::LLVMWriteBitcodeToMemoryBuffer; use llvm_zluda::{core::*, *}; use llvm_zluda::{prelude::*, LLVMZludaBuildAtomicRMW}; @@ -118,6 +120,24 @@ impl Drop for Module { } } +impl From for Module { + fn from(memory_buffer: MemoryBuffer) -> Self { + let context = Context::new(); + let mut module: MaybeUninit = MaybeUninit::uninit(); + unsafe { + LLVMParseBitcodeInContext2(context.get(), memory_buffer.get(), module.as_mut_ptr()); + } + let module = unsafe { module.assume_init() }; + Self(module, context) + } +} + +pub fn bitcode_to_ir(bitcode: Vec) -> Vec { + let memory_buffer: MemoryBuffer = bitcode.into(); + let module: Module = memory_buffer.into(); + module.print_module_to_string().to_bytes().to_vec() +} + struct Builder(LLVMBuilderRef); impl Builder { @@ -170,6 +190,12 @@ impl Message { pub struct MemoryBuffer(LLVMMemoryBufferRef); +impl MemoryBuffer { + fn get(&self) -> LLVMMemoryBufferRef { + self.0 + } +} + impl Drop for MemoryBuffer { fn drop(&mut self) { unsafe { @@ -188,6 +214,26 @@ impl Deref for MemoryBuffer { } } +impl From> for MemoryBuffer { + fn from(value: Vec) -> Self { + let memory_buffer: LLVMMemoryBufferRef = unsafe { + LLVMCreateMemoryBufferWithMemoryRangeCopy( + value.as_ptr(), + value.len(), + ptr::null() + ) + }; + Self(memory_buffer) + } +} + +impl From> for MemoryBuffer { + fn from(value: Vec) -> Self { + let value: Vec = value.iter().map(|&v| i8::from_ne_bytes([v])).collect(); + value.into() + } +} + pub(super) fn run<'input>( id_defs: GlobalStringIdentResolver2<'input>, directives: Vec, SpirvWord>>, diff --git a/zoc/src/error.rs b/zoc/src/error.rs new file mode 100644 index 0000000..9c268c3 --- /dev/null +++ b/zoc/src/error.rs @@ -0,0 +1,66 @@ +use std::io; +use std::path::PathBuf; +use std::str::Utf8Error; + +use amd_comgr_sys::amd_comgr_status_s; +use hip_runtime_sys::hipErrorCode_t; +use ptx::TranslateError; +use ptx_parser::PtxError; + +#[derive(Debug, thiserror::Error)] +pub enum CompilerError { + #[error("HIP error: {0:?}")] + HipError(hipErrorCode_t), + #[error("amd_comgr error: {0:?}")] + ComgrError(amd_comgr_status_s), + #[error("Not a regular file: {0}")] + CheckPathError(PathBuf), + #[error("Invalid output type: {0}")] + ParseOutputTypeError(String), + #[error("Error parsing PTX: {0}")] + PtxParserError(String), + #[error("Error translating PTX: {0:?}")] + PtxTranslateError(TranslateError), + #[error("IO error: {0:?}")] + IoError(io::Error), + #[error("Error parsing file: {0:?}")] + ParseFileError(Utf8Error), +} + +impl From for CompilerError { + fn from(error_code: hipErrorCode_t) -> Self { + CompilerError::HipError(error_code) + } +} + +impl From for CompilerError { + fn from(error_code: amd_comgr_status_s) -> Self { + CompilerError::ComgrError(error_code) + } +} + +impl From>> for CompilerError { + fn from(causes: Vec) -> Self { + let errors: Vec = causes.iter().map(PtxError::to_string).collect(); + let msg = errors.join("\n"); + CompilerError::PtxParserError(msg) + } +} + +impl From for CompilerError { + fn from(cause: io::Error) -> Self { + CompilerError::IoError(cause) + } +} + +impl From for CompilerError { + fn from(cause: Utf8Error) -> Self { + CompilerError::ParseFileError(cause) + } +} + +impl From for CompilerError { + fn from(cause: TranslateError) -> Self { + CompilerError::PtxTranslateError(cause) + } +} \ No newline at end of file diff --git a/zoc/src/main.rs b/zoc/src/main.rs index 0f7b91e..b34bad1 100644 --- a/zoc/src/main.rs +++ b/zoc/src/main.rs @@ -1,5 +1,4 @@ use std::env; -use std::error::Error; use std::ffi::{CStr, OsStr}; use std::fs::{self, File}; use std::io::{self, Write}; @@ -7,10 +6,11 @@ use std::mem::MaybeUninit; use std::path::{Path, PathBuf}; use std::str::{self, FromStr}; -use amd_comgr_sys::amd_comgr_status_s; +use amd_comgr_sys::amd_comgr_data_kind_s; use bpaf::Bpaf; -use hip_runtime_sys::hipErrorCode_t; -use ptx_parser::PtxError; + +mod error; +use error::CompilerError; #[derive(Debug, Clone, Bpaf)] #[bpaf(options, version)] @@ -23,13 +23,13 @@ pub struct Options { ptx_path: String, } -fn main() -> Result<(), Box> { +fn main() -> Result<(), CompilerError> { let opts = options().run(); let output_type = opts.output_type.unwrap_or_default(); match output_type { - OutputType::LlvmIrLinked | OutputType::Assembly => todo!(), + OutputType::Assembly => todo!(), _ => {} } @@ -39,24 +39,30 @@ fn main() -> Result<(), Box> { let output_path = get_output_path(&ptx_path, &output_type)?; check_path(&output_path)?; - let ptx = fs::read(&ptx_path)?; - let ptx = str::from_utf8(&ptx)?; - let llvm = ptx_to_llvm(ptx)?; + let ptx = fs::read(&ptx_path).map_err(CompilerError::from)?; + let ptx = str::from_utf8(&ptx).map_err(CompilerError::from)?; + let llvm = ptx_to_llvm(ptx).map_err(CompilerError::from)?; if output_type == OutputType::LlvmIrPreLinked { - write_to_file(&llvm.llvm_ir, &output_path)?; + write_to_file(&llvm.llvm_ir, &output_path).map_err(CompilerError::from)?; + return Ok(()); + } + + if output_type == OutputType::LlvmIrLinked { + let linked_llvm = link_llvm(&llvm)?; + write_to_file(&linked_llvm, &output_path).map_err(CompilerError::from)?; return Ok(()); } let elf = llvm_to_elf(&llvm)?; - write_to_file(&elf, &output_path)?; + write_to_file(&elf, &output_path).map_err(CompilerError::from)?; Ok(()) } -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)?; +fn ptx_to_llvm(ptx: &str) -> Result { + let ast = ptx_parser::parse_module_checked(ptx).map_err(CompilerError::from).map_err(CompilerError::from)?; + let module = ptx::to_llvm_module(ast).map_err(CompilerError::from)?; let bitcode = module.llvm_ir.write_bitcode_to_memory().to_vec(); let linked_bitcode = module.linked_bitcode().to_vec(); let llvm_ir = module.llvm_ir.print_module_to_string().to_bytes().to_vec(); @@ -74,12 +80,14 @@ struct LLVMArtifacts { llvm_ir: Vec, } -fn join_ptx_errors(vector: Vec) -> String { - let errors: Vec = vector.iter().map(PtxError::to_string).collect(); - errors.join("\n") +fn link_llvm(llvm: &LLVMArtifacts) -> Result, CompilerError> { + let linked_bitcode = comgr::link_bitcode(&llvm.bitcode, &llvm.linked_bitcode)?; + let data = linked_bitcode.get_data(amd_comgr_data_kind_s::AMD_COMGR_DATA_KIND_BC, 0)?; + let linked_llvm = data.copy_content().map_err(CompilerError::from)?; + Ok(ptx::bitcode_to_ir(linked_llvm)) } -fn llvm_to_elf(llvm: &LLVMArtifacts) -> Result, ElfError> { +fn llvm_to_elf(llvm: &LLVMArtifacts) -> Result, CompilerError> { use hip_runtime_sys::*; unsafe { hipInit(0) }?; let mut dev_props: MaybeUninit = MaybeUninit::uninit(); @@ -87,13 +95,12 @@ fn llvm_to_elf(llvm: &LLVMArtifacts) -> Result, ElfError> { let dev_props = unsafe { dev_props.assume_init() }; let gcn_arch = unsafe { CStr::from_ptr(dev_props.gcnArchName.as_ptr()) }; - comgr::compile_bitcode(gcn_arch, &llvm.bitcode, &llvm.linked_bitcode).map_err(ElfError::from) + comgr::compile_bitcode(gcn_arch, &llvm.bitcode, &llvm.linked_bitcode).map_err(CompilerError::from) } -fn check_path(path: &Path) -> Result<(), Box> { - if path.try_exists()? && !path.is_file() { - let error = CheckPathError(path.to_path_buf()); - let error = Box::new(error); +fn check_path(path: &Path) -> Result<(), CompilerError> { + if path.try_exists().map_err(CompilerError::from)? && !path.is_file() { + let error = CompilerError::CheckPathError(path.to_path_buf()); return Err(error); } Ok(()) @@ -102,8 +109,8 @@ fn check_path(path: &Path) -> Result<(), Box> { fn get_output_path( ptx_path: &PathBuf, output_type: &OutputType, -) -> Result> { - let current_dir = env::current_dir()?; +) -> Result { + let current_dir = env::current_dir().map_err(CompilerError::from)?; let output_path = current_dir.join( ptx_path .as_path() @@ -150,7 +157,7 @@ impl OutputType { } impl FromStr for OutputType { - type Err = ParseOutputTypeError; + type Err = CompilerError; fn from_str(s: &str) -> Result { match s { @@ -158,35 +165,7 @@ impl FromStr for OutputType { "ll_linked" => Ok(Self::LlvmIrLinked), "elf" => Ok(Self::Elf), "asm" => Ok(Self::Assembly), - _ => Err(ParseOutputTypeError(s.into())), + _ => Err(CompilerError::ParseOutputTypeError(s.into())), } } } - -#[derive(Debug, thiserror::Error)] -#[error("Not a regular file: {0}")] -struct CheckPathError(PathBuf); - -#[derive(Debug, thiserror::Error)] -#[error("Invalid output type: {0}")] -struct ParseOutputTypeError(String); - -#[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) - } -} From 6d9ffdc3b0f17d41cafa3838e01ed6d27541524d Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Jo=C3=ABlle=20van=20Essen?= Date: Wed, 26 Mar 2025 19:59:30 +0100 Subject: [PATCH 08/15] zoc: Rename compiler module, executable name stays the same (zoc) --- Cargo.lock | 26 +++++++++++++------------- Cargo.toml | 4 ++-- {zoc => compiler}/Cargo.toml | 6 +++++- {zoc => compiler}/src/error.rs | 0 {zoc => compiler}/src/main.rs | 0 5 files changed, 20 insertions(+), 16 deletions(-) rename {zoc => compiler}/Cargo.toml (90%) rename {zoc => compiler}/src/error.rs (100%) rename {zoc => compiler}/src/main.rs (100%) diff --git a/Cargo.lock b/Cargo.lock index 5f4bf00..b0f2605 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -211,6 +211,19 @@ dependencies = [ "amd_comgr-sys", ] +[[package]] +name = "compiler" +version = "0.0.0" +dependencies = [ + "amd_comgr-sys", + "bpaf", + "comgr", + "hip_runtime-sys", + "ptx", + "ptx_parser", + "thiserror 2.0.12", +] + [[package]] name = "convert_case" version = "0.6.0" @@ -1393,16 +1406,3 @@ dependencies = [ "wchar", "winapi", ] - -[[package]] -name = "zoc" -version = "0.0.0" -dependencies = [ - "amd_comgr-sys", - "bpaf", - "comgr", - "hip_runtime-sys", - "ptx", - "ptx_parser", - "thiserror 2.0.12", -] diff --git a/Cargo.toml b/Cargo.toml index 04c215b..4873fd5 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -20,10 +20,10 @@ members = [ "ptx_parser_macros_impl", "xtask", "zluda_bindgen", - "zoc", + "compiler", ] -default-members = ["zluda", "zluda_ml", "zluda_inject", "zluda_redirect", "zoc"] +default-members = ["zluda", "zluda_ml", "zluda_inject", "zluda_redirect", "compiler"] [profile.release-lto] inherits = "release" diff --git a/zoc/Cargo.toml b/compiler/Cargo.toml similarity index 90% rename from zoc/Cargo.toml rename to compiler/Cargo.toml index 146252b..ebbb314 100644 --- a/zoc/Cargo.toml +++ b/compiler/Cargo.toml @@ -1,10 +1,14 @@ [package] -name = "zoc" +name = "compiler" description = "ZLUDA offline compiler" version = "0.0.0" authors = ["Joëlle van Essen "] edition = "2021" +[[bin]] +name = "zoc" +path = "src/main.rs" + [dependencies] amd_comgr-sys = { path = "../ext/amd_comgr-sys" } bpaf = { version = "0.9.18", features = ["derive"] } diff --git a/zoc/src/error.rs b/compiler/src/error.rs similarity index 100% rename from zoc/src/error.rs rename to compiler/src/error.rs diff --git a/zoc/src/main.rs b/compiler/src/main.rs similarity index 100% rename from zoc/src/main.rs rename to compiler/src/main.rs From 13f1a3490e5f3f447a454cc1837e48667752dd05 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Jo=C3=ABlle=20van=20Essen?= Date: Wed, 26 Mar 2025 21:42:49 +0100 Subject: [PATCH 09/15] zoc: Produce disassembly of executable --- comgr/Cargo.toml | 2 +- comgr/src/lib.rs | 72 +++++++++++++++++++++++++++++------ compiler/src/error.rs | 13 +++++-- compiler/src/main.rs | 57 +++++++++++++-------------- ptx/Cargo.toml | 2 +- ptx/src/test/spirv_run/mod.rs | 2 +- zluda/src/impl/module.rs | 2 +- 7 files changed, 101 insertions(+), 49 deletions(-) diff --git a/comgr/Cargo.toml b/comgr/Cargo.toml index 250bd0a..9598952 100644 --- a/comgr/Cargo.toml +++ b/comgr/Cargo.toml @@ -1,7 +1,7 @@ [package] name = "comgr" version = "0.0.0" -authors = ["Andrzej Janik "] +authors = ["Andrzej Janik ", "Joëlle van Essen "] edition = "2021" [lib] diff --git a/comgr/src/lib.rs b/comgr/src/lib.rs index e1a872b..bb4679b 100644 --- a/comgr/src/lib.rs +++ b/comgr/src/lib.rs @@ -110,10 +110,7 @@ impl Drop for ActionInfo { } } -pub fn link_bitcode( - main_buffer: &[u8], - ptx_impl: &[u8], -) -> Result { +fn link_bitcode(main_buffer: &[u8], ptx_impl: &[u8]) -> Result { use amd_comgr_sys::*; let bitcode_data_set = DataSet::new()?; let main_bitcode_data = Data::new( @@ -136,14 +133,12 @@ pub fn link_bitcode( ) } -pub fn compile_bitcode( +fn compile_bitcode( + linked_data_set: &DataSet, gcn_arch: &CStr, - main_buffer: &[u8], - ptx_impl: &[u8], -) -> Result, amd_comgr_status_s> { +) -> Result { use amd_comgr_sys::*; - let linked_data_set = link_bitcode(main_buffer, ptx_impl)?; let compile_to_exec = ActionInfo::new()?; compile_to_exec.set_isa_name(gcn_arch)?; compile_to_exec.set_language(amd_comgr_language_t::AMD_COMGR_LANGUAGE_LLVM_IR)?; @@ -176,16 +171,69 @@ pub fn compile_bitcode( ] }; compile_to_exec.set_options(common_options.chain(opt_options))?; - let exec_data_set = do_action( + do_action( &linked_data_set, &compile_to_exec, amd_comgr_action_kind_t::AMD_COMGR_ACTION_COMPILE_SOURCE_TO_EXECUTABLE, - )?; + ) +} + +fn disassemble_exec( + exec_data_set: &DataSet, + gcn_arch: &CStr, +) -> Result { + let action_info = ActionInfo::new()?; + action_info.set_isa_name(gcn_arch)?; + do_action( + &exec_data_set, + &action_info, + amd_comgr_action_kind_t::AMD_COMGR_ACTION_DISASSEMBLE_EXECUTABLE_TO_SOURCE, + ) +} + +pub fn get_linked_bitcode_as_bytes( + main_buffer: &[u8], + ptx_impl: &[u8], +) -> Result, amd_comgr_status_s> { + let linked_data_set = link_bitcode(main_buffer, ptx_impl)?; + let linked_bitcode = + linked_data_set.get_data(amd_comgr_data_kind_t::AMD_COMGR_DATA_KIND_BC, 0)?; + linked_bitcode.copy_content() +} + +fn get_executable( + gcn_arch: &CStr, + main_buffer: &[u8], + ptx_impl: &[u8], +) -> Result { + let linked_data_set = link_bitcode(main_buffer, ptx_impl)?; + let exec_data_set = compile_bitcode(&linked_data_set, gcn_arch)?; + Ok(exec_data_set) +} + +pub fn get_executable_as_bytes( + gcn_arch: &CStr, + main_buffer: &[u8], + ptx_impl: &[u8], +) -> Result, amd_comgr_status_s> { + let exec_data_set = get_executable(gcn_arch, main_buffer, ptx_impl)?; let executable = - exec_data_set.get_data(amd_comgr_data_kind_t::AMD_COMGR_DATA_KIND_EXECUTABLE, 0)?; + exec_data_set.get_data(amd_comgr_data_kind_t::AMD_COMGR_DATA_KIND_EXECUTABLE, 0)?; executable.copy_content() } +pub fn get_assembly_as_bytes( + gcn_arch: &CStr, + main_buffer: &[u8], + ptx_impl: &[u8], +) -> Result, amd_comgr_status_s> { + let exec_data_set = get_executable(gcn_arch, main_buffer, ptx_impl)?; + let disassembled_data_set = disassemble_exec(&exec_data_set, gcn_arch)?; + let assembly = + disassembled_data_set.get_data(amd_comgr_data_kind_t::AMD_COMGR_DATA_KIND_SOURCE, 0)?; + assembly.copy_content() +} + fn do_action( data_set: &DataSet, action: &ActionInfo, diff --git a/compiler/src/error.rs b/compiler/src/error.rs index 9c268c3..fc8d004 100644 --- a/compiler/src/error.rs +++ b/compiler/src/error.rs @@ -1,3 +1,4 @@ +use std::ffi::FromBytesUntilNulError; use std::io; use std::path::PathBuf; use std::str::Utf8Error; @@ -17,14 +18,14 @@ pub enum CompilerError { CheckPathError(PathBuf), #[error("Invalid output type: {0}")] ParseOutputTypeError(String), - #[error("Error parsing PTX: {0}")] - PtxParserError(String), #[error("Error translating PTX: {0:?}")] PtxTranslateError(TranslateError), #[error("IO error: {0:?}")] IoError(io::Error), #[error("Error parsing file: {0:?}")] ParseFileError(Utf8Error), + #[error("Error: {0}")] + GenericError(String) } impl From for CompilerError { @@ -43,7 +44,7 @@ impl From>> for CompilerError { fn from(causes: Vec) -> Self { let errors: Vec = causes.iter().map(PtxError::to_string).collect(); let msg = errors.join("\n"); - CompilerError::PtxParserError(msg) + CompilerError::GenericError(msg) } } @@ -63,4 +64,10 @@ impl From for CompilerError { fn from(cause: TranslateError) -> Self { CompilerError::PtxTranslateError(cause) } +} + +impl From for CompilerError { + fn from(cause: FromBytesUntilNulError) -> Self { + CompilerError::GenericError(format!("{}", cause)) + } } \ No newline at end of file diff --git a/compiler/src/main.rs b/compiler/src/main.rs index b34bad1..3df41fb 100644 --- a/compiler/src/main.rs +++ b/compiler/src/main.rs @@ -1,12 +1,11 @@ use std::env; -use std::ffi::{CStr, OsStr}; +use std::ffi::{CStr, CString, OsStr}; use std::fs::{self, File}; use std::io::{self, Write}; use std::mem::MaybeUninit; use std::path::{Path, PathBuf}; use std::str::{self, FromStr}; -use amd_comgr_sys::amd_comgr_data_kind_s; use bpaf::Bpaf; mod error; @@ -28,11 +27,6 @@ fn main() -> Result<(), CompilerError> { let output_type = opts.output_type.unwrap_or_default(); - match output_type { - OutputType::Assembly => todo!(), - _ => {} - } - let ptx_path = Path::new(&opts.ptx_path).to_path_buf(); check_path(&ptx_path)?; @@ -43,20 +37,14 @@ fn main() -> Result<(), CompilerError> { let ptx = str::from_utf8(&ptx).map_err(CompilerError::from)?; let llvm = ptx_to_llvm(ptx).map_err(CompilerError::from)?; - if output_type == OutputType::LlvmIrPreLinked { - write_to_file(&llvm.llvm_ir, &output_path).map_err(CompilerError::from)?; - return Ok(()); - } - - if output_type == OutputType::LlvmIrLinked { - let linked_llvm = link_llvm(&llvm)?; - write_to_file(&linked_llvm, &output_path).map_err(CompilerError::from)?; - return Ok(()); - } - - let elf = llvm_to_elf(&llvm)?; - write_to_file(&elf, &output_path).map_err(CompilerError::from)?; + let output = match output_type { + OutputType::LlvmIrPreLinked => llvm.llvm_ir, + OutputType::LlvmIrLinked => get_linked_bitcode(&llvm)?, + OutputType::Elf => get_elf(&llvm)?, + OutputType::Assembly => get_assembly(&llvm)? + }; + write_to_file(&output, &output_path).map_err(CompilerError::from)?; Ok(()) } @@ -80,22 +68,31 @@ struct LLVMArtifacts { llvm_ir: Vec, } -fn link_llvm(llvm: &LLVMArtifacts) -> Result, CompilerError> { - let linked_bitcode = comgr::link_bitcode(&llvm.bitcode, &llvm.linked_bitcode)?; - let data = linked_bitcode.get_data(amd_comgr_data_kind_s::AMD_COMGR_DATA_KIND_BC, 0)?; - let linked_llvm = data.copy_content().map_err(CompilerError::from)?; - Ok(ptx::bitcode_to_ir(linked_llvm)) -} - -fn llvm_to_elf(llvm: &LLVMArtifacts) -> Result, CompilerError> { +fn get_arch() -> Result { use hip_runtime_sys::*; unsafe { hipInit(0) }?; let mut dev_props: MaybeUninit = MaybeUninit::uninit(); unsafe { hipGetDevicePropertiesR0600(dev_props.as_mut_ptr(), 0) }?; let dev_props = unsafe { dev_props.assume_init() }; - let gcn_arch = unsafe { CStr::from_ptr(dev_props.gcnArchName.as_ptr()) }; + let arch = dev_props.gcnArchName; + let arch: Vec = arch.to_vec().iter().map(|&v| i8::to_ne_bytes(v)[0]).collect(); + let arch = CStr::from_bytes_until_nul(arch.as_slice())?; + Ok(CString::from(arch)) +} - comgr::compile_bitcode(gcn_arch, &llvm.bitcode, &llvm.linked_bitcode).map_err(CompilerError::from) +fn get_linked_bitcode(llvm: &LLVMArtifacts) -> Result, CompilerError> { + let linked_bitcode = comgr::get_linked_bitcode_as_bytes(&llvm.bitcode, &llvm.linked_bitcode)?; + Ok(ptx::bitcode_to_ir(linked_bitcode)) +} + +fn get_elf(llvm: &LLVMArtifacts) -> Result, CompilerError> { + let arch = get_arch()?; + comgr::get_executable_as_bytes(&arch, &llvm.bitcode, &llvm.linked_bitcode).map_err(CompilerError::from) +} + +fn get_assembly(llvm: &LLVMArtifacts) -> Result, CompilerError> { + let arch = get_arch()?; + comgr::get_assembly_as_bytes(&arch, &llvm.bitcode, &llvm.linked_bitcode).map_err(CompilerError::from) } fn check_path(path: &Path) -> Result<(), CompilerError> { diff --git a/ptx/Cargo.toml b/ptx/Cargo.toml index 2876539..59e35f2 100644 --- a/ptx/Cargo.toml +++ b/ptx/Cargo.toml @@ -1,7 +1,7 @@ [package] name = "ptx" version = "0.0.0" -authors = ["Andrzej Janik "] +authors = ["Andrzej Janik ", "Joëlle van Essen "] edition = "2021" [lib] diff --git a/ptx/src/test/spirv_run/mod.rs b/ptx/src/test/spirv_run/mod.rs index cafa480..ddde159 100644 --- a/ptx/src/test/spirv_run/mod.rs +++ b/ptx/src/test/spirv_run/mod.rs @@ -361,7 +361,7 @@ fn run_hip + Copy + Debug, Output: From + 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 = comgr::compile_bitcode( + let elf_module = comgr::get_executable_as_bytes( unsafe { CStr::from_ptr(dev_props.gcnArchName.as_ptr()) }, &*module.llvm_ir.write_bitcode_to_memory(), module.linked_bitcode(), diff --git a/zluda/src/impl/module.rs b/zluda/src/impl/module.rs index a881e16..c74ae03 100644 --- a/zluda/src/impl/module.rs +++ b/zluda/src/impl/module.rs @@ -28,7 +28,7 @@ pub(crate) fn load_data(module: &mut CUmodule, image: *const std::ffi::c_void) - unsafe { hipCtxGetDevice(&mut dev) }?; let mut props = unsafe { mem::zeroed() }; unsafe { hipGetDevicePropertiesR0600(&mut props, dev) }?; - let elf_module = comgr::compile_bitcode( + let elf_module = comgr::get_executable_as_bytes( unsafe { CStr::from_ptr(props.gcnArchName.as_ptr()) }, &*llvm_module.llvm_ir.write_bitcode_to_memory(), llvm_module.linked_bitcode(), From 1ce4a40b0f06819da0d22bddcff4960398fa1c8c Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Jo=C3=ABlle=20van=20Essen?= Date: Wed, 26 Mar 2025 22:06:17 +0100 Subject: [PATCH 10/15] zoc: Revert unnecessary pub keywords --- comgr/src/lib.rs | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/comgr/src/lib.rs b/comgr/src/lib.rs index bb4679b..1c29510 100644 --- a/comgr/src/lib.rs +++ b/comgr/src/lib.rs @@ -3,7 +3,7 @@ use std::ffi::CStr; use std::mem; use std::ptr; -pub struct Data(amd_comgr_data_t); +struct Data(amd_comgr_data_t); impl Data { fn new( @@ -22,7 +22,7 @@ impl Data { self.0 } - pub fn copy_content(&self) -> Result, amd_comgr_status_s> { + fn copy_content(&self) -> Result, 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 = Vec::with_capacity(size); @@ -32,7 +32,7 @@ impl Data { } } -pub struct DataSet(amd_comgr_data_set_t); +struct DataSet(amd_comgr_data_set_t); impl DataSet { fn new() -> Result { @@ -49,7 +49,7 @@ impl DataSet { self.0 } - pub fn get_data( + fn get_data( &self, kind: amd_comgr_data_kind_t, index: usize, From e6e9a79fe710310a1591fc6a06eb0d5880b8f46c Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Jo=C3=ABlle=20van=20Essen?= Date: Thu, 27 Mar 2025 10:18:11 +0100 Subject: [PATCH 11/15] zoc: Specify output path --- Cargo.lock | 4 ++-- compiler/Cargo.toml | 2 +- compiler/src/main.rs | 9 ++++++++- 3 files changed, 11 insertions(+), 4 deletions(-) diff --git a/Cargo.lock b/Cargo.lock index b0f2605..f627511 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -104,9 +104,9 @@ checksum = "b048fb63fd8b5923fc5aa7b340d8e156aec7ec02f0c78fa8a6ddc2613f6f71de" [[package]] name = "bpaf" -version = "0.9.18" +version = "0.9.19" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "4de4d74c5891642753c67ab88f58d971a68dd98673b69689a8c24ce3ec78a412" +checksum = "4848ed5727d39a7573551c205bcb1ccd88c8cad4ed2c80f62e2316f208196b8d" dependencies = [ "bpaf_derive", ] diff --git a/compiler/Cargo.toml b/compiler/Cargo.toml index ebbb314..66d4123 100644 --- a/compiler/Cargo.toml +++ b/compiler/Cargo.toml @@ -11,7 +11,7 @@ path = "src/main.rs" [dependencies] amd_comgr-sys = { path = "../ext/amd_comgr-sys" } -bpaf = { version = "0.9.18", features = ["derive"] } +bpaf = { version = "0.9.19", features = ["derive"] } comgr = { path = "../comgr" } hip_runtime-sys = { path = "../ext/hip_runtime-sys" } ptx = { path = "../ptx" } diff --git a/compiler/src/main.rs b/compiler/src/main.rs index 3df41fb..3a3d7e3 100644 --- a/compiler/src/main.rs +++ b/compiler/src/main.rs @@ -17,6 +17,10 @@ pub struct Options { #[bpaf(external(output_type), optional)] output_type: Option, + #[bpaf(short('o'), argument("file"))] + /// Output path + output_path: Option, + #[bpaf(positional("filename"))] /// PTX file ptx_path: String, @@ -30,7 +34,10 @@ fn main() -> Result<(), CompilerError> { let ptx_path = Path::new(&opts.ptx_path).to_path_buf(); check_path(&ptx_path)?; - let output_path = get_output_path(&ptx_path, &output_type)?; + let output_path = match opts.output_path { + Some(value) => value, + None => get_output_path(&ptx_path, &output_type)? + }; check_path(&output_path)?; let ptx = fs::read(&ptx_path).map_err(CompilerError::from)?; From 69080d2abca9cdad4a4f76b3eb85622326551619 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Jo=C3=ABlle=20van=20Essen?= Date: Thu, 27 Mar 2025 11:29:21 +0100 Subject: [PATCH 12/15] zoc: Improve error messages --- Cargo.lock | 27 +++++++++++++++-- compiler/src/error.rs | 68 ++++++++++++++++++++++--------------------- compiler/src/main.rs | 51 +++++++++++++++++++++++--------- ptx/src/pass/mod.rs | 4 +-- ptx_parser/Cargo.toml | 11 +++---- ptx_parser/src/lib.rs | 2 +- 6 files changed, 107 insertions(+), 56 deletions(-) diff --git a/Cargo.lock b/Cargo.lock index f627511..c049067 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -820,8 +820,8 @@ dependencies = [ "ptx_parser", "quick-error", "rustc-hash 2.0.0", - "strum", - "strum_macros", + "strum 0.26.3", + "strum_macros 0.26.4", "tempfile", "thiserror 1.0.64", "unwrap_or", @@ -836,6 +836,7 @@ dependencies = [ "logos", "ptx_parser_macros", "rustc-hash 2.0.0", + "strum 0.27.1", "thiserror 1.0.64", "winnow", ] @@ -1047,6 +1048,15 @@ version = "0.26.3" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "8fec0f0aef304996cf250b31b5a10dee7980c85da9d759361292b8bca5a18f06" +[[package]] +name = "strum" +version = "0.27.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "f64def088c51c9510a8579e3c5d67c65349dcf755e5479ad3d010aa6454e2c32" +dependencies = [ + "strum_macros 0.27.1", +] + [[package]] name = "strum_macros" version = "0.26.4" @@ -1060,6 +1070,19 @@ dependencies = [ "syn 2.0.89", ] +[[package]] +name = "strum_macros" +version = "0.27.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "c77a8c5abcaf0f9ce05d62342b7d298c346515365c36b673df4ebe3ced01fde8" +dependencies = [ + "heck", + "proc-macro2", + "quote", + "rustversion", + "syn 2.0.89", +] + [[package]] name = "syn" version = "1.0.109" diff --git a/compiler/src/error.rs b/compiler/src/error.rs index fc8d004..396f70d 100644 --- a/compiler/src/error.rs +++ b/compiler/src/error.rs @@ -1,6 +1,5 @@ use std::ffi::FromBytesUntilNulError; use std::io; -use std::path::PathBuf; use std::str::Utf8Error; use amd_comgr_sys::amd_comgr_status_s; @@ -10,22 +9,20 @@ use ptx_parser::PtxError; #[derive(Debug, thiserror::Error)] pub enum CompilerError { - #[error("HIP error: {0:?}")] + #[error("HIP error code: {0:?}")] HipError(hipErrorCode_t), - #[error("amd_comgr error: {0:?}")] + #[error("amd_comgr status code: {0:?}")] ComgrError(amd_comgr_status_s), - #[error("Not a regular file: {0}")] - CheckPathError(PathBuf), - #[error("Invalid output type: {0}")] - ParseOutputTypeError(String), - #[error("Error translating PTX: {0:?}")] - PtxTranslateError(TranslateError), - #[error("IO error: {0:?}")] - IoError(io::Error), - #[error("Error parsing file: {0:?}")] - ParseFileError(Utf8Error), - #[error("Error: {0}")] - GenericError(String) + #[error(transparent)] + IoError(#[from] io::Error), + #[error(transparent)] + Utf8Error(#[from] Utf8Error), + #[error("{message}")] + GenericError { + #[source] + cause: Option>, + message: String, + }, } impl From for CompilerError { @@ -42,32 +39,37 @@ impl From for CompilerError { impl From>> for CompilerError { fn from(causes: Vec) -> Self { - let errors: Vec = causes.iter().map(PtxError::to_string).collect(); - let msg = errors.join("\n"); - CompilerError::GenericError(msg) - } -} - -impl From for CompilerError { - fn from(cause: io::Error) -> Self { - CompilerError::IoError(cause) - } -} - -impl From for CompilerError { - fn from(cause: Utf8Error) -> Self { - CompilerError::ParseFileError(cause) + let errors: Vec = causes + .iter() + .map(|e| { + let msg = match e { + PtxError::UnrecognizedStatement(value) + | PtxError::UnrecognizedDirective(value) => value.unwrap_or("").to_string(), + other => other.to_string(), + }; + format!("PtxError::{}: {}", e.as_ref(), msg) + }) + .collect(); + let message = errors.join("\n"); + CompilerError::GenericError { + cause: None, + message, + } } } impl From for CompilerError { fn from(cause: TranslateError) -> Self { - CompilerError::PtxTranslateError(cause) + let message = format!("PTX TranslateError::{}", cause.as_ref()); + let cause = Some(Box::new(cause) as Box); + CompilerError::GenericError { cause, message } } } impl From for CompilerError { fn from(cause: FromBytesUntilNulError) -> Self { - CompilerError::GenericError(format!("{}", cause)) + let message = format!("{}", cause); + let cause = Some(Box::new(cause) as Box); + CompilerError::GenericError { cause, message } } -} \ No newline at end of file +} diff --git a/compiler/src/main.rs b/compiler/src/main.rs index 3a3d7e3..2d3fa8f 100644 --- a/compiler/src/main.rs +++ b/compiler/src/main.rs @@ -1,9 +1,11 @@ use std::env; +use std::error::Error; use std::ffi::{CStr, CString, OsStr}; use std::fs::{self, File}; use std::io::{self, Write}; use std::mem::MaybeUninit; use std::path::{Path, PathBuf}; +use std::process::ExitCode; use std::str::{self, FromStr}; use bpaf::Bpaf; @@ -26,7 +28,17 @@ pub struct Options { ptx_path: String, } -fn main() -> Result<(), CompilerError> { +fn main() -> ExitCode { + main_core().map_or_else( + |e| { + eprintln!("Error: {}", e); + ExitCode::FAILURE + }, + |_| ExitCode::SUCCESS, + ) +} + +fn main_core() -> Result<(), CompilerError> { let opts = options().run(); let output_type = opts.output_type.unwrap_or_default(); @@ -36,7 +48,7 @@ fn main() -> Result<(), CompilerError> { let output_path = match opts.output_path { Some(value) => value, - None => get_output_path(&ptx_path, &output_type)? + None => get_output_path(&ptx_path, &output_type)?, }; check_path(&output_path)?; @@ -48,7 +60,7 @@ fn main() -> Result<(), CompilerError> { OutputType::LlvmIrPreLinked => llvm.llvm_ir, OutputType::LlvmIrLinked => get_linked_bitcode(&llvm)?, OutputType::Elf => get_elf(&llvm)?, - OutputType::Assembly => get_assembly(&llvm)? + OutputType::Assembly => get_assembly(&llvm)?, }; write_to_file(&output, &output_path).map_err(CompilerError::from)?; @@ -56,7 +68,7 @@ fn main() -> Result<(), CompilerError> { } fn ptx_to_llvm(ptx: &str) -> Result { - let ast = ptx_parser::parse_module_checked(ptx).map_err(CompilerError::from).map_err(CompilerError::from)?; + let ast = ptx_parser::parse_module_checked(ptx).map_err(CompilerError::from)?; let module = ptx::to_llvm_module(ast).map_err(CompilerError::from)?; let bitcode = module.llvm_ir.write_bitcode_to_memory().to_vec(); let linked_bitcode = module.linked_bitcode().to_vec(); @@ -82,7 +94,11 @@ fn get_arch() -> Result { unsafe { hipGetDevicePropertiesR0600(dev_props.as_mut_ptr(), 0) }?; let dev_props = unsafe { dev_props.assume_init() }; let arch = dev_props.gcnArchName; - let arch: Vec = arch.to_vec().iter().map(|&v| i8::to_ne_bytes(v)[0]).collect(); + let arch: Vec = arch + .to_vec() + .iter() + .map(|&v| i8::to_ne_bytes(v)[0]) + .collect(); let arch = CStr::from_bytes_until_nul(arch.as_slice())?; Ok(CString::from(arch)) } @@ -94,26 +110,29 @@ fn get_linked_bitcode(llvm: &LLVMArtifacts) -> Result, CompilerError> { fn get_elf(llvm: &LLVMArtifacts) -> Result, CompilerError> { let arch = get_arch()?; - comgr::get_executable_as_bytes(&arch, &llvm.bitcode, &llvm.linked_bitcode).map_err(CompilerError::from) + comgr::get_executable_as_bytes(&arch, &llvm.bitcode, &llvm.linked_bitcode) + .map_err(CompilerError::from) } fn get_assembly(llvm: &LLVMArtifacts) -> Result, CompilerError> { let arch = get_arch()?; - comgr::get_assembly_as_bytes(&arch, &llvm.bitcode, &llvm.linked_bitcode).map_err(CompilerError::from) + comgr::get_assembly_as_bytes(&arch, &llvm.bitcode, &llvm.linked_bitcode) + .map_err(CompilerError::from) } fn check_path(path: &Path) -> Result<(), CompilerError> { if path.try_exists().map_err(CompilerError::from)? && !path.is_file() { - let error = CompilerError::CheckPathError(path.to_path_buf()); + let message = format!("Not a regular file: {:?}", path.to_path_buf()); + let error = CompilerError::GenericError { + cause: None, + message, + }; return Err(error); } Ok(()) } -fn get_output_path( - ptx_path: &PathBuf, - output_type: &OutputType, -) -> Result { +fn get_output_path(ptx_path: &PathBuf, output_type: &OutputType) -> Result { let current_dir = env::current_dir().map_err(CompilerError::from)?; let output_path = current_dir.join( ptx_path @@ -169,7 +188,13 @@ impl FromStr for OutputType { "ll_linked" => Ok(Self::LlvmIrLinked), "elf" => Ok(Self::Elf), "asm" => Ok(Self::Assembly), - _ => Err(CompilerError::ParseOutputTypeError(s.into())), + _ => { + let message = format!("Not a valid output type: {}", s); + Err(CompilerError::GenericError { + cause: None, + message, + }) + } } } } diff --git a/ptx/src/pass/mod.rs b/ptx/src/pass/mod.rs index 77d7e60..2abdac3 100644 --- a/ptx/src/pass/mod.rs +++ b/ptx/src/pass/mod.rs @@ -17,8 +17,8 @@ mod expand_operands; mod fix_special_registers2; mod hoist_globals; mod insert_explicit_load_store; -mod instruction_mode_to_global_mode; mod insert_implicit_conversions2; +mod instruction_mode_to_global_mode; mod normalize_basic_blocks; mod normalize_identifiers2; mod normalize_predicates2; @@ -31,7 +31,7 @@ static ZLUDA_PTX_IMPL: &'static [u8] = include_bytes!("../../lib/zluda_ptx_impl. const ZLUDA_PTX_PREFIX: &'static str = "__zluda_ptx_impl_"; quick_error! { - #[derive(Debug)] + #[derive(Debug, strum_macros::AsRefStr)] pub enum TranslateError { UnknownSymbol {} UntypedSymbol {} diff --git a/ptx_parser/Cargo.toml b/ptx_parser/Cargo.toml index 9032de5..3b96ac0 100644 --- a/ptx_parser/Cargo.toml +++ b/ptx_parser/Cargo.toml @@ -7,11 +7,12 @@ edition = "2021" [lib] [dependencies] +bitflags = "1.2" +derive_more = { version = "1", features = ["display"] } logos = "0.14" +ptx_parser_macros = { path = "../ptx_parser_macros" } +rustc-hash = "2.0.0" +strum = { version = "0.27.1", features = ["derive"] } +thiserror = "1.0" winnow = { version = "0.6.18" } #winnow = { version = "0.6.18", features = ["debug"] } -ptx_parser_macros = { path = "../ptx_parser_macros" } -thiserror = "1.0" -bitflags = "1.2" -rustc-hash = "2.0.0" -derive_more = { version = "1", features = ["display"] } diff --git a/ptx_parser/src/lib.rs b/ptx_parser/src/lib.rs index da46a8c..aea4dc6 100644 --- a/ptx_parser/src/lib.rs +++ b/ptx_parser/src/lib.rs @@ -1240,7 +1240,7 @@ impl ast::ParsedOperand { } } -#[derive(Debug, thiserror::Error)] +#[derive(Debug, thiserror::Error, strum::AsRefStr)] pub enum PtxError<'input> { #[error("{source}")] ParseInt { From af7139fb9c981c52d9764e6821431987e9af8a57 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Jo=C3=ABlle=20van=20Essen?= Date: Thu, 27 Mar 2025 12:10:03 +0100 Subject: [PATCH 13/15] zoc: Handle FromBytesUntilNulError as transparent --- compiler/src/error.rs | 10 ++-------- 1 file changed, 2 insertions(+), 8 deletions(-) diff --git a/compiler/src/error.rs b/compiler/src/error.rs index 396f70d..0c6bc16 100644 --- a/compiler/src/error.rs +++ b/compiler/src/error.rs @@ -17,6 +17,8 @@ pub enum CompilerError { IoError(#[from] io::Error), #[error(transparent)] Utf8Error(#[from] Utf8Error), + #[error(transparent)] + FromBytesUntilNulError(#[from] FromBytesUntilNulError), #[error("{message}")] GenericError { #[source] @@ -65,11 +67,3 @@ impl From for CompilerError { CompilerError::GenericError { cause, message } } } - -impl From for CompilerError { - fn from(cause: FromBytesUntilNulError) -> Self { - let message = format!("{}", cause); - let cause = Some(Box::new(cause) as Box); - CompilerError::GenericError { cause, message } - } -} From d3b61c9808e882c19798c97e67eb7fffebef1efa Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Jo=C3=ABlle=20van=20Essen?= Date: Thu, 27 Mar 2025 12:28:17 +0100 Subject: [PATCH 14/15] zoc: Treat linked LLVM IR as additional option --linked --- compiler/src/main.rs | 29 +++++++++++++++++++---------- 1 file changed, 19 insertions(+), 10 deletions(-) diff --git a/compiler/src/main.rs b/compiler/src/main.rs index 2d3fa8f..e032f8f 100644 --- a/compiler/src/main.rs +++ b/compiler/src/main.rs @@ -19,6 +19,10 @@ pub struct Options { #[bpaf(external(output_type), optional)] output_type: Option, + #[bpaf(long("linked"))] + /// Produce linked LLVM IR (in combination with --ll) + linked: bool, + #[bpaf(short('o'), argument("file"))] /// Output path output_path: Option, @@ -43,6 +47,10 @@ fn main_core() -> Result<(), CompilerError> { let output_type = opts.output_type.unwrap_or_default(); + if opts.linked && output_type != OutputType::LlvmIr { + println!("Warning: option --linked only makes sense when combined with --ll. Ignoring."); + } + let ptx_path = Path::new(&opts.ptx_path).to_path_buf(); check_path(&ptx_path)?; @@ -57,8 +65,13 @@ fn main_core() -> Result<(), CompilerError> { let llvm = ptx_to_llvm(ptx).map_err(CompilerError::from)?; let output = match output_type { - OutputType::LlvmIrPreLinked => llvm.llvm_ir, - OutputType::LlvmIrLinked => get_linked_bitcode(&llvm)?, + OutputType::LlvmIr => { + if opts.linked { + get_linked_bitcode(&llvm)? + } else { + llvm.llvm_ir + } + } OutputType::Elf => get_elf(&llvm)?, OutputType::Assembly => get_assembly(&llvm)?, }; @@ -154,12 +167,9 @@ fn write_to_file(content: &[u8], path: &Path) -> io::Result<()> { #[derive(Bpaf, Clone, Copy, Debug, Default, PartialEq)] enum OutputType { - /// Produce pre-linked LLVM IR + /// Produce LLVM IR #[bpaf(long("ll"))] - LlvmIrPreLinked, - /// Produce linked LLVM IR - #[bpaf(long("linked-ll"))] - LlvmIrLinked, + LlvmIr, /// Produce ELF binary (default) #[default] Elf, @@ -171,7 +181,7 @@ enum OutputType { impl OutputType { fn extension(self) -> String { match self { - OutputType::LlvmIrPreLinked | OutputType::LlvmIrLinked => "ll", + OutputType::LlvmIr => "ll", OutputType::Assembly => "asm", OutputType::Elf => "elf", } @@ -184,8 +194,7 @@ impl FromStr for OutputType { fn from_str(s: &str) -> Result { match s { - "ll" => Ok(Self::LlvmIrPreLinked), - "ll_linked" => Ok(Self::LlvmIrLinked), + "ll" => Ok(Self::LlvmIr), "elf" => Ok(Self::Elf), "asm" => Ok(Self::Assembly), _ => { From 7b5384e33f514facc01b277deb64d85781dbd1c1 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Jo=C3=ABlle=20van=20Essen?= Date: Sun, 30 Mar 2025 14:42:22 +0200 Subject: [PATCH 15/15] zoc: Remove unused import --- compiler/src/main.rs | 1 - 1 file changed, 1 deletion(-) diff --git a/compiler/src/main.rs b/compiler/src/main.rs index e032f8f..039e104 100644 --- a/compiler/src/main.rs +++ b/compiler/src/main.rs @@ -1,5 +1,4 @@ use std::env; -use std::error::Error; use std::ffi::{CStr, CString, OsStr}; use std::fs::{self, File}; use std::io::{self, Write};