From 1f1b770a511e4dd9116167bf1b02aae0df07ee3b Mon Sep 17 00:00:00 2001
From: Andrzej Janik <vosen@vosen.pl>
Date: Sun, 1 Dec 2024 03:36:16 +0000
Subject: Fix more missing stuff in the host code

---
 ptx/lib/zluda_ptx_impl.bc  | Bin 4624 -> 4816 bytes
 ptx/lib/zluda_ptx_impl.cpp |   8 ++++----
 zluda/src/impl/memory.rs   |  12 +++++++++++-
 zluda/src/impl/mod.rs      |   2 +-
 zluda/src/lib.rs           |   1 +
 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
Binary files a/ptx/lib/zluda_ptx_impl.bc and b/ptx/lib/zluda_ptx_impl.bc differ
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,
-- 
cgit v1.2.3