Merge 7b5384e33f into 7cdab7abc2

Cargo.lock

@@ -104,18 +104,18 @@ checksum = "b048fb63fd8b5923fc5aa7b340d8e156aec7ec02f0c78fa8a6ddc2613f6f71de"
 
 [[package]]
 name = "bpaf"
-version = "0.9.15"
+version = "0.9.19"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "50fd5174866dc2fa2ddc96e8fb800852d37f064f32a45c7b7c2f8fa2c64c77fa"
+checksum = "4848ed5727d39a7573551c205bcb1ccd88c8cad4ed2c80f62e2316f208196b8d"
 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",
@@ -157,7 +157,7 @@ dependencies = [
  "semver",
  "serde",
  "serde_json",
- "thiserror 2.0.11",
+ "thiserror 2.0.12",
 ]
 
 [[package]]
@@ -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"
@@ -807,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",
@@ -823,6 +836,7 @@ dependencies = [
  "logos",
  "ptx_parser_macros",
  "rustc-hash 2.0.0",
+ "strum 0.27.1",
  "thiserror 1.0.64",
  "winnow",
 ]
@@ -1034,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"
@@ -1047,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"
@@ -1093,11 +1129,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]]
@@ -1113,9 +1149,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",

Cargo.toml

@@ -20,9 +20,10 @@ members = [
     "ptx_parser_macros_impl",
     "xtask",
     "zluda_bindgen",
+    "compiler",
 ]
 
-default-members = ["zluda", "zluda_ml", "zluda_inject", "zluda_redirect"]
+default-members = ["zluda", "zluda_ml", "zluda_inject", "zluda_redirect", "compiler"]
 
 [profile.release-lto]
 inherits = "release"

comgr/Cargo.toml

@@ -1,7 +1,7 @@
 [package]
 name = "comgr"
 version = "0.0.0"
-authors = ["Andrzej Janik <vosen@vosen.pl>"]
+authors = ["Andrzej Janik <vosen@vosen.pl>", "Joëlle van Essen <joelle@v-essen.nl>"]
 edition = "2021"
 
 [lib]

comgr/src/lib.rs

@@ -1,5 +1,7 @@
 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);
 
@@ -108,11 +110,7 @@ impl Drop for ActionInfo {
     }
 }
 
-pub fn compile_bitcode(
-    gcn_arch: &CStr,
-    main_buffer: &[u8],
-    ptx_impl: &[u8],
-) -> Result<Vec<u8>, amd_comgr_status_s> {
+fn link_bitcode(main_buffer: &[u8], ptx_impl: &[u8]) -> Result<DataSet, amd_comgr_status_s> {
     use amd_comgr_sys::*;
     let bitcode_data_set = DataSet::new()?;
     let main_bitcode_data = Data::new(
@@ -128,11 +126,19 @@ 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,
-    )?;
+    )
+}
+
+fn compile_bitcode(
+    linked_data_set: &DataSet,
+    gcn_arch: &CStr,
+) -> Result<DataSet, amd_comgr_status_s> {
+    use amd_comgr_sys::*;
+
     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)?;
@@ -165,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<DataSet, amd_comgr_status_s> {
+    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<Vec<u8>, 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<DataSet, amd_comgr_status_s> {
+    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<Vec<u8>, 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<Vec<u8>, 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,

compiler/Cargo.toml

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

compiler/src/error.rs

@@ -0,0 +1,69 @@
+use std::ffi::FromBytesUntilNulError;
+use std::io;
+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 code: {0:?}")]
+    HipError(hipErrorCode_t),
+    #[error("amd_comgr status code: {0:?}")]
+    ComgrError(amd_comgr_status_s),
+    #[error(transparent)]
+    IoError(#[from] io::Error),
+    #[error(transparent)]
+    Utf8Error(#[from] Utf8Error),
+    #[error(transparent)]
+    FromBytesUntilNulError(#[from] FromBytesUntilNulError),
+    #[error("{message}")]
+    GenericError {
+        #[source]
+        cause: Option<Box<dyn std::error::Error>>,
+        message: String,
+    },
+}
+
+impl From<hipErrorCode_t> for CompilerError {
+    fn from(error_code: hipErrorCode_t) -> Self {
+        CompilerError::HipError(error_code)
+    }
+}
+
+impl From<amd_comgr_status_s> for CompilerError {
+    fn from(error_code: amd_comgr_status_s) -> Self {
+        CompilerError::ComgrError(error_code)
+    }
+}
+
+impl From<Vec<PtxError<'_>>> for CompilerError {
+    fn from(causes: Vec<PtxError>) -> Self {
+        let errors: Vec<String> = 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<TranslateError> for CompilerError {
+    fn from(cause: TranslateError) -> Self {
+        let message = format!("PTX TranslateError::{}", cause.as_ref());
+        let cause = Some(Box::new(cause) as Box<dyn std::error::Error>);
+        CompilerError::GenericError { cause, message }
+    }
+}

compiler/src/main.rs

@@ -0,0 +1,208 @@
+use std::env;
+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;
+
+mod error;
+use error::CompilerError;
+
+#[derive(Debug, Clone, Bpaf)]
+#[bpaf(options, version)]
+pub struct Options {
+    #[bpaf(external(output_type), optional)]
+    output_type: Option<OutputType>,
+
+    #[bpaf(long("linked"))]
+    /// Produce linked LLVM IR (in combination with --ll)
+    linked: bool,
+
+    #[bpaf(short('o'), argument("file"))]
+    /// Output path
+    output_path: Option<PathBuf>,
+
+    #[bpaf(positional("filename"))]
+    /// PTX file
+    ptx_path: String,
+}
+
+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();
+
+    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)?;
+
+    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)?;
+    let ptx = str::from_utf8(&ptx).map_err(CompilerError::from)?;
+    let llvm = ptx_to_llvm(ptx).map_err(CompilerError::from)?;
+
+    let output = match output_type {
+        OutputType::LlvmIr => {
+            if opts.linked {
+                get_linked_bitcode(&llvm)?
+            } else {
+                llvm.llvm_ir
+            }
+        }
+        OutputType::Elf => get_elf(&llvm)?,
+        OutputType::Assembly => get_assembly(&llvm)?,
+    };
+
+    write_to_file(&output, &output_path).map_err(CompilerError::from)?;
+    Ok(())
+}
+
+fn ptx_to_llvm(ptx: &str) -> Result<LLVMArtifacts, CompilerError> {
+    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();
+    let llvm_ir = module.llvm_ir.print_module_to_string().to_bytes().to_vec();
+    Ok(LLVMArtifacts {
+        bitcode,
+        linked_bitcode,
+        llvm_ir,
+    })
+}
+
+#[derive(Debug)]
+struct LLVMArtifacts {
+    bitcode: Vec<u8>,
+    linked_bitcode: Vec<u8>,
+    llvm_ir: Vec<u8>,
+}
+
+fn get_arch() -> Result<CString, CompilerError> {
+    use hip_runtime_sys::*;
+    unsafe { hipInit(0) }?;
+    let mut dev_props: MaybeUninit<hipDeviceProp_tR0600> = MaybeUninit::uninit();
+    unsafe { hipGetDevicePropertiesR0600(dev_props.as_mut_ptr(), 0) }?;
+    let dev_props = unsafe { dev_props.assume_init() };
+    let arch = dev_props.gcnArchName;
+    let arch: Vec<u8> = 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))
+}
+
+fn get_linked_bitcode(llvm: &LLVMArtifacts) -> Result<Vec<u8>, 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<Vec<u8>, 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<Vec<u8>, 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> {
+    if path.try_exists().map_err(CompilerError::from)? && !path.is_file() {
+        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<PathBuf, CompilerError> {
+    let current_dir = env::current_dir().map_err(CompilerError::from)?;
+    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(Bpaf, Clone, Copy, Debug, Default, PartialEq)]
+enum OutputType {
+    /// Produce LLVM IR
+    #[bpaf(long("ll"))]
+    LlvmIr,
+    /// Produce ELF binary (default)
+    #[default]
+    Elf,
+    /// Produce assembly
+    #[bpaf(long("asm"))]
+    Assembly,
+}
+
+impl OutputType {
+    fn extension(self) -> String {
+        match self {
+            OutputType::LlvmIr => "ll",
+            OutputType::Assembly => "asm",
+            OutputType::Elf => "elf",
+        }
+        .into()
+    }
+}
+
+impl FromStr for OutputType {
+    type Err = CompilerError;
+
+    fn from_str(s: &str) -> Result<Self, Self::Err> {
+        match s {
+            "ll" => Ok(Self::LlvmIr),
+            "elf" => Ok(Self::Elf),
+            "asm" => Ok(Self::Assembly),
+            _ => {
+                let message = format!("Not a valid output type: {}", s);
+                Err(CompilerError::GenericError {
+                    cause: None,
+                    message,
+                })
+            }
+        }
+    }
+}

ptx/Cargo.toml

@@ -1,7 +1,7 @@
 [package]
 name = "ptx"
 version = "0.0.0"
-authors = ["Andrzej Janik <vosen@vosen.pl>"]
+authors = ["Andrzej Janik <vosen@vosen.pl>", "Joëlle van Essen <joelle@v-essen.nl>"]
 edition = "2021"
 
 [lib]

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;

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};
@@ -119,6 +121,24 @@ impl Drop for Module {
     }
 }
 
+impl From<MemoryBuffer> for Module {
+    fn from(memory_buffer: MemoryBuffer) -> Self {
+        let context = Context::new();
+        let mut module: MaybeUninit<LLVMModuleRef> = 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<u8>) -> Vec<u8> {
+    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 {
@@ -160,6 +180,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()
     }
@@ -167,6 +191,12 @@ impl Message {
 
 pub struct MemoryBuffer(LLVMMemoryBufferRef);
 
+impl MemoryBuffer {
+    fn get(&self) -> LLVMMemoryBufferRef {
+        self.0
+    }
+}
+
 impl Drop for MemoryBuffer {
     fn drop(&mut self) {
         unsafe {
@@ -185,6 +215,26 @@ impl Deref for MemoryBuffer {
     }
 }
 
+impl From<Vec<i8>> for MemoryBuffer {
+    fn from(value: Vec<i8>) -> Self {
+        let memory_buffer: LLVMMemoryBufferRef = unsafe {
+            LLVMCreateMemoryBufferWithMemoryRangeCopy(
+                value.as_ptr(),
+                value.len(),
+                ptr::null()
+            )
+        };
+        Self(memory_buffer)
+    }
+}
+
+impl From<Vec<u8>> for MemoryBuffer {
+    fn from(value: Vec<u8>) -> Self {
+        let value: Vec<i8> = value.iter().map(|&v| i8::from_ne_bytes([v])).collect();
+        value.into()
+    }
+}
+
 pub(super) fn run<'input>(
     id_defs: GlobalStringIdentResolver2<'input>,
     directives: Vec<Directive2<ast::Instruction<SpirvWord>, SpirvWord>>,

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 {}

ptx/src/test/spirv_run/mod.rs

@@ -380,7 +380,7 @@ 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 = 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(),

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"] }

ptx_parser/src/lib.rs

@@ -1240,7 +1240,7 @@ impl<Ident> ast::ParsedOperand<Ident> {
     }
 }
 
-#[derive(Debug, thiserror::Error)]
+#[derive(Debug, thiserror::Error, strum::AsRefStr)]
 pub enum PtxError<'input> {
     #[error("{source}")]
     ParseInt {

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