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(),