Fix more missing stuff in the host codecuda-12-4

5 files changed, 17 insertions, 6 deletions

diff --git a/ptx/lib/zluda_ptx_impl.bc b/ptx/lib/zluda_ptx_impl.bc
index 6651430..4b5a5d8 100644
--- a/ptx/lib/zluda_ptx_impl.bc
+++ b/ptx/lib/zluda_ptx_impl.bc
diff --git a/ptx/lib/zluda_ptx_impl.cpp b/ptx/lib/zluda_ptx_impl.cpp
index d0ec853..f86a7fd 100644
--- a/ptx/lib/zluda_ptx_impl.cpp
+++ b/ptx/lib/zluda_ptx_impl.cpp
@@ -25,16 +25,16 @@ extern "C"
         return (uint32_t)__ockl_get_local_size(member);

     }

 

-    size_t __ockl_get_global_id(uint32_t) __device__;

+    size_t __ockl_get_group_id(uint32_t) __device__;

     uint32_t FUNC(sreg_ctaid)(uint8_t member)

     {

-        return (uint32_t)__ockl_get_global_id(member);

+        return (uint32_t)__ockl_get_group_id(member);

     }

 

-    size_t __ockl_get_global_size(uint32_t) __device__;

+    size_t __ockl_get_num_groups(uint32_t) __device__;

     uint32_t FUNC(sreg_nctaid)(uint8_t member)

     {

-        return (uint32_t)__ockl_get_global_size(member);

+        return (uint32_t)__ockl_get_num_groups(member);

     }

 

     uint32_t __ockl_bfe_u32(uint32_t, uint32_t, uint32_t) __attribute__((device));

diff --git a/zluda/src/impl/memory.rs b/zluda/src/impl/memory.rs
index b23afa9..3843776 100644
--- a/zluda/src/impl/memory.rs
+++ b/zluda/src/impl/memory.rs
@@ -1,7 +1,9 @@
 use hip_runtime_sys::*;
 
 pub(crate) fn alloc_v2(dptr: *mut hipDeviceptr_t, bytesize: usize) -> hipError_t {
-    unsafe { hipMalloc(dptr.cast(), bytesize) }
+    unsafe { hipMalloc(dptr.cast(), bytesize) }?;
+    // TODO: parametrize for non-Geekbench
+    unsafe { hipMemsetD8(*dptr, 0, bytesize) }
 }
 
 pub(crate) fn free_v2(dptr: hipDeviceptr_t) -> hipError_t {
@@ -23,3 +25,11 @@ pub(crate) fn copy_hto_d_v2(
 ) -> hipError_t {
     unsafe { hipMemcpyHtoD(dst_device, src_host.cast_mut(), byte_count) }
 }
+
+pub(crate) fn get_address_range_v2(
+    pbase: *mut hipDeviceptr_t,
+    psize: *mut usize,
+    dptr: hipDeviceptr_t,
+) -> hipError_t {
+    unsafe { hipMemGetAddressRange(pbase, psize, dptr) }
+}
diff --git a/zluda/src/impl/mod.rs b/zluda/src/impl/mod.rs
index 7b4afc5..766b4a5 100644
--- a/zluda/src/impl/mod.rs
+++ b/zluda/src/impl/mod.rs
@@ -162,7 +162,7 @@ impl<T: ZludaObject> LiveCheck<T> {
     }
 
     fn as_handle(&self) -> T::CudaHandle {
-        unsafe { mem::transmute_copy(self) }
+        unsafe { mem::transmute_copy(&self) }
     }
 
     fn wrap(data: T) -> *mut Self {
diff --git a/zluda/src/lib.rs b/zluda/src/lib.rs
index bda67e1..1568f47 100644
--- a/zluda/src/lib.rs
+++ b/zluda/src/lib.rs
@@ -71,6 +71,7 @@ cuda_base::cuda_function_declarations!(
         cuModuleLoadData,
         cuModuleUnload,
         cuPointerGetAttribute,
+        cuMemGetAddressRange_v2,
     ],
     implemented_in_function <= [
         cuLaunchKernel,