Fix more missing stuff in the host code

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));

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) }
+}

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 {

zluda/src/lib.rs

@@ -71,6 +71,7 @@ cuda_base::cuda_function_declarations!(
         cuModuleLoadData,
         cuModuleUnload,
         cuPointerGetAttribute,
+        cuMemGetAddressRange_v2,
     ],
     implemented_in_function <= [
         cuLaunchKernel,