ZLUDA v3.8.6 (#61)

* Fix hipBLASLt on ROCm 6.
Enable hipBLASLt on Windows.

* Clean up.

* Add cublasXt apis.

* Remove warnings.

* Resolve or hide warnings.

* Fix bug.

* Regenerate cublas bindings.

* Build improvements.

* Bump versions.

* Restore Linux build as usual.

* Add missing call to hgemm.

* or_else.

* Fix.

* Fix bug.

* Clean up.

Cargo.toml

@@ -1,5 +1,7 @@
 [workspace]
 
+resolver = "2"
+
 # Remember to also update the project's Cargo.toml
 # if it's a top-level project
 members = [

comgr/src/lib.rs

@@ -300,7 +300,7 @@ impl Comgr {
             }
         };
         if self.2 == 1 {
-            eprintln!("Compiling in progress. Please wait...");
+            eprintln!("Compilation is in progress. Please wait...");
         }
         let relocatable = self.do_action(
             sys::amd_comgr_action_kind_t::AMD_COMGR_ACTION_CODEGEN_BC_TO_RELOCATABLE,

hipblaslt-sys/README

@@ -1 +1 @@
-bindgen /opt/rocm/include/hipblaslt/hipblaslt.h -o src/hipblaslt.rs --no-layout-tests --default-enum-style=newtype --allowlist-function "^hipblasLt.*" --allowlist-type "^hipblasLt.*" --no-derive-debug --must-use-type hiprtError -- -I /opt/rocm/include -D__HIP_PLATFORM_AMD__ -D__HIP_PLATFORM_HCC__ -x c++
+bindgen $Env:HIP_PATH/include/hipblaslt/hipblaslt.h -o src/hipblaslt.rs --no-layout-tests --default-enum-style=newtype --allowlist-function "^hipblasLt.*" --allowlist-type "^hipblasLt.*" --no-derive-debug --must-use-type hiprtError -- -I"$Env:HIP_PATH/include" -D__HIP_PLATFORM_AMD__ -D__HIP_PLATFORM_HCC__ -x c++

hipblaslt-sys/build.rs

@@ -1,7 +1,19 @@
 use std::env::VarError;
+use std::{env, path::PathBuf};
 
 fn main() -> Result<(), VarError> {
     println!("cargo:rustc-link-lib=dylib=hipblaslt");
-    println!("cargo:rustc-link-search=native=/opt/rocm/lib/");
+    if cfg!(windows) {
+        let env = env::var("CARGO_CFG_TARGET_ENV")?;
+        if env == "msvc" {
+            let mut path = PathBuf::from(env::var("HIP_PATH")?);
+            path.push("lib");
+            println!("cargo:rustc-link-search=native={}", path.display());
+        } else {
+            println!("cargo:rustc-link-search=native=C:\\Windows\\System32");
+        };
+    } else {
+        println!("cargo:rustc-link-search=native=/opt/rocm/lib/");
+    }
     Ok(())
 }

hipblaslt-sys/src/lib.rs

@@ -1,3 +1,16 @@
 #[allow(warnings)]
 mod hipblaslt;
-pub use hipblaslt::*;
+pub use hipblaslt::*;
+
+impl hipblasOperation_t {
+    pub const HIPBLAS_OP_N: hipblasOperation_t = hipblasOperation_t(111);
+}
+impl hipblasOperation_t {
+    pub const HIPBLAS_OP_T: hipblasOperation_t = hipblasOperation_t(112);
+}
+impl hipblasOperation_t {
+    pub const HIPBLAS_OP_C: hipblasOperation_t = hipblasOperation_t(113);
+}
+#[repr(transparent)]
+#[derive(Copy, Clone, Hash, PartialEq, Eq)]
+pub struct hipblasOperation_t(pub ::std::os::raw::c_int);

ptx/Cargo.toml

@@ -2,7 +2,7 @@
 name = "ptx"
 version = "0.0.0"
 authors = ["Andrzej Janik <[email protected]>"]
-edition = "2018"
+edition = "2021"
 
 [lib]
 
@@ -12,7 +12,6 @@ rocm5 = ["hip_common/rocm5", "hip_runtime-sys/rocm5", "comgr/rocm5"]
 [dependencies]
 hip_common = { path = "../hip_common" }
 zluda_llvm = { path = "../zluda_llvm" }
-lalrpop-util = "0.20"
 regex = "1"
 thiserror = "1.0"
 num-traits = "0.2"
@@ -26,6 +25,10 @@ either = "1.9"
 version = "1.8"
 features = ["num-traits"]
 
+[dependencies.lalrpop-util]
+version = "0.22"
+features = ["lexer"]
+
 [dev-dependencies]
 cuda_base = { path = "../cuda_base" }
 cuda_types = { path = "../cuda_types" }
@@ -37,5 +40,5 @@ libloading = "0.8"
 lazy_static = "1.4.0"
 
 [build-dependencies.lalrpop]
-version = "0.20"
+version = "0.22"
 features = ["lexer"]

ptx/src/translate.rs

@@ -90,7 +90,7 @@ impl<'input> Module<'input> {
 
     pub fn get_bitcode_all<'a>(
         &'a self,
-    ) -> impl Iterator<Item = (llvm::MemoryBuffer, &'a CStr)> + '_ {
+    ) -> impl Iterator<Item = (llvm::MemoryBuffer, &'a CStr)> + 'a {
         unsafe {
             let main_bc = llvm::MemoryBuffer::from_ffi(LLVMWriteBitcodeToMemoryBuffer(
                 self.llvm_module.get(),
@@ -3278,7 +3278,7 @@ pub(crate) struct DenormSummary {
 pub fn to_llvm_module<'input>(
     compilation_mode: CompilationMode,
     ast: Vec<ast::Module<'input>>,
-) -> Result<Module, TranslateError> {
+) -> Result<Module<'input>, TranslateError> {
     to_llvm_module_impl2(compilation_mode, ast, None)
 }
 

xtask/src/main.rs

@@ -50,6 +50,7 @@ impl Default for Subcommand {
         Subcommand::Build(BuildCommand {
             release: false,
             rocm5: false,
+            nightly: false,
         })
     }
 }
@@ -65,6 +66,10 @@ struct BuildCommand {
     /// build for ROCm 5 (Windows only)
     #[argh(switch)]
     rocm5: bool,
+
+    /// enable unstable features
+    #[argh(switch)]
+    nightly: bool,
 }
 
 #[derive(FromArgs)]
@@ -84,13 +89,17 @@ struct ZipCommand {
 fn main() -> Result<(), DynError> {
     let args: Arguments = argh::from_env();
     std::process::exit(match args.command {
-        Subcommand::Build(BuildCommand { release, rocm5 }) => build(!release, rocm5)?,
+        Subcommand::Build(BuildCommand {
+            release,
+            rocm5,
+            nightly,
+        }) => build(!release, rocm5, nightly)?,
         Subcommand::Zip(ZipCommand { release, rocm5 }) => build_and_zip(!release, rocm5),
     })
 }
 
 fn build_and_zip(is_debug: bool, rocm5: bool) -> i32 {
-    let workspace = build_impl(is_debug, rocm5).unwrap();
+    let workspace = build_impl(is_debug, rocm5, false).unwrap();
     os::zip(workspace)
 }
 
@@ -109,6 +118,8 @@ struct Project {
     #[serde(skip_deserializing)]
     kind: TargetKind,
     #[serde(default)]
+    windows_nightly: bool,
+    #[serde(default)]
     windows_only: bool,
     #[serde(default)]
     linux_only: bool,
@@ -143,7 +154,7 @@ struct Workspace {
 }
 
 impl Workspace {
-    fn open(is_debug: bool) -> Result<Self, DynError> {
+    fn open(is_debug: bool, nightly: bool) -> Result<Self, DynError> {
         let cargo = env::var("CARGO").unwrap_or_else(|_| "cargo".to_string());
         let project_root = Self::project_root()?;
         let mut cmd = cargo_metadata::MetadataCommand::new();
@@ -153,7 +164,7 @@ impl Workspace {
             .packages
             .into_iter()
             .filter_map(Project::new)
-            .filter(|p| !p.skip_build(is_debug))
+            .filter(|p| !p.skip_build(is_debug, nightly))
             .collect::<Vec<_>>();
         let mut target_directory = cargo_metadata.target_directory;
         target_directory.push(if is_debug { "debug" } else { "release" });
@@ -199,7 +210,7 @@ impl Project {
         Some(project)
     }
 
-    fn skip_build(&self, is_debug: bool) -> bool {
+    fn skip_build(&self, is_debug: bool, nightly: bool) -> bool {
         if self.broken {
             return true;
         }
@@ -212,17 +223,20 @@ impl Project {
         if !is_debug && self.debug_only {
             return true;
         }
+        if cfg!(windows) && !nightly && self.windows_nightly {
+            return true;
+        }
         false
     }
 }
 
-fn build(is_debug: bool, rocm5: bool) -> Result<i32, DynError> {
-    build_impl(is_debug, rocm5)?;
+fn build(is_debug: bool, rocm5: bool, nightly: bool) -> Result<i32, DynError> {
+    build_impl(is_debug, rocm5, nightly)?;
     Ok(0)
 }
 
-fn build_impl(is_debug: bool, rocm5: bool) -> Result<Workspace, DynError> {
-    let workspace = Workspace::open(is_debug)?;
+fn build_impl(is_debug: bool, rocm5: bool, nightly: bool) -> Result<Workspace, DynError> {
+    let workspace = Workspace::open(is_debug, nightly)?;
     let mut command = workspace.cargo_command();
     command.arg("build");
     command.arg("--locked");
@@ -243,7 +257,12 @@ fn build_impl(is_debug: bool, rocm5: bool) -> Result<Workspace, DynError> {
         if let Ok(path_default) = env::var("HIP_PATH") {
             env::set_var(
                 "HIP_PATH",
-                env::var(if rocm5 { "HIP_PATH_57" } else { "HIP_PATH_62" }).unwrap_or(path_default),
+                if rocm5 {
+                    env::var("HIP_PATH_57").or_else(|_| env::var("HIP_PATH_55"))
+                } else {
+                    env::var("HIP_PATH_62").or_else(|_| env::var("HIP_PATH_61"))
+                }
+                .unwrap_or(path_default),
             );
         } else {
             return Err(

zluda/Cargo.toml

@@ -2,7 +2,7 @@
 name = "zluda"
 version = "0.0.0"
 authors = ["Andrzej Janik <[email protected]>"]
-edition = "2018"
+edition = "2021"
 
 [lib]
 name = "zluda"

zluda/src/impl/dark_api.rs

@@ -35,8 +35,13 @@ static CUDA_DARK_API_TABLE: CudaDarkApiTable = zluda_dark_api::init_dark_api::<C
 
 struct CudaDarkApiZluda;
 
-static mut TOOLS_RUNTIME_CALLBACK_HOOKS_FN2_SPACE: [usize; 1024] = [0; 1024];
-static mut TOOLS_RUNTIME_CALLBACK_HOOKS_FN6_SPACE: [u8; 14] = [0; 14];
+const TOOLS_RUNTIME_CALLBACK_HOOKS_FN2_SPACE_SIZE: usize = 1024;
+const TOOLS_RUNTIME_CALLBACK_HOOKS_FN6_SPACE_SIZE: usize = 14;
+
+static mut TOOLS_RUNTIME_CALLBACK_HOOKS_FN2_SPACE: [usize;
+    TOOLS_RUNTIME_CALLBACK_HOOKS_FN2_SPACE_SIZE] = [0; TOOLS_RUNTIME_CALLBACK_HOOKS_FN2_SPACE_SIZE];
+static mut TOOLS_RUNTIME_CALLBACK_HOOKS_FN6_SPACE: [u8;
+    TOOLS_RUNTIME_CALLBACK_HOOKS_FN6_SPACE_SIZE] = [0; TOOLS_RUNTIME_CALLBACK_HOOKS_FN6_SPACE_SIZE];
 
 impl CudaDarkApi for CudaDarkApiZluda {
     unsafe extern "system" fn get_module_from_cubin(
@@ -86,9 +91,7 @@ impl CudaDarkApi for CudaDarkApiZluda {
         res
     }
 
-    unsafe extern "system" fn set_device(
-        dev: cuda_types::CUdevice,
-    ) -> CUresult {
+    unsafe extern "system" fn set_device(dev: cuda_types::CUdevice) -> CUresult {
         use hip_runtime_sys::*;
         hipSetDevice(FromCuda::from_cuda(dev)).into_cuda()
     }
@@ -134,16 +137,16 @@ impl CudaDarkApi for CudaDarkApiZluda {
         ptr: *mut *mut usize,
         size: *mut usize,
     ) -> () {
-        *ptr = TOOLS_RUNTIME_CALLBACK_HOOKS_FN2_SPACE.as_mut_ptr();
-        *size = TOOLS_RUNTIME_CALLBACK_HOOKS_FN2_SPACE.len();
+        *ptr = &raw mut TOOLS_RUNTIME_CALLBACK_HOOKS_FN2_SPACE as _;
+        *size = TOOLS_RUNTIME_CALLBACK_HOOKS_FN2_SPACE_SIZE;
     }
 
     unsafe extern "system" fn tools_runtime_callback_hooks_fn6(
         ptr: *mut *mut u8,
         size: *mut usize,
     ) -> () {
-        *ptr = TOOLS_RUNTIME_CALLBACK_HOOKS_FN6_SPACE.as_mut_ptr();
-        *size = TOOLS_RUNTIME_CALLBACK_HOOKS_FN6_SPACE.len();
+        *ptr = &raw mut TOOLS_RUNTIME_CALLBACK_HOOKS_FN6_SPACE as _;
+        *size = TOOLS_RUNTIME_CALLBACK_HOOKS_FN6_SPACE_SIZE;
     }
 
     unsafe extern "system" fn context_local_storage_insert(

zluda/src/impl/mod.rs

@@ -141,7 +141,7 @@ impl<T: ZludaObject> LiveCheck<T> {
         outer_ptr as *mut Self
     }
 
-    pub unsafe fn as_ref_unchecked(&self) -> & T {
+    pub unsafe fn as_ref_unchecked(&self) -> &T {
         &self.data
     }