Archive/ZLUDA
0
0
Fork 0
mirror of https://github.com/vosen/ZLUDA.git synced 2025-04-28 21:47:57 +03:00
Code Issues Projects Releases Packages Wiki Activity Actions

Zero out buffer on creation

Browse source
This commit is contained in:
Andrzej Janik 2021-09-13 23:43:50 +00:00
parent dbb6f09ffa
commit 986fa49097
3 changed files with 19 additions and 1 deletions
Download patch file Download diff file Expand all files Collapse all files

1
ptx/src/test/spirv_run/mod.rs
View file

@ -637,6 +637,7 @@ fn compile_amd(
let mut compiler_cmd = Command::new(&clang_exe); let mut compiler_cmd = Command::new(&clang_exe);
compiler_cmd compiler_cmd
.arg(format!("-mcpu={}", device_name)) .arg(format!("-mcpu={}", device_name))
.arg("-ffp-contract=off")
.arg("-nogpulib") .arg("-nogpulib")
.arg("-mno-wavefrontsize64") .arg("-mno-wavefrontsize64")
.arg("-O3") .arg("-O3")

18
zluda/src/cuda.rs
View file

@ -2588,7 +2588,23 @@ pub extern "system" fn cuMemGetInfo_v2(free: *mut usize, total: *mut usize) -> C
#[cfg_attr(not(test), no_mangle)] #[cfg_attr(not(test), no_mangle)]
pub extern "system" fn cuMemAlloc_v2(dptr: *mut CUdeviceptr, bytesize: usize) -> CUresult { pub extern "system" fn cuMemAlloc_v2(dptr: *mut CUdeviceptr, bytesize: usize) -> CUresult {
unsafe { hipMalloc(dptr as _, bytesize).into() } let mut dev_ptr = std::ptr::null_mut();
let err = unsafe { hipMalloc(&mut dev_ptr, bytesize) }.into();
if err != CUresult::CUDA_SUCCESS {
return err;
}
// HACK ALERT: GeekBench is buggy and sometimes assumes that buffers are zeroed-out on creation
let err = unsafe { hipMemsetD8(dev_ptr, 0, bytesize) }.into();
/*
let bytesize_rounded_down = bytesize & !3usize;
let bytes = usize::min(bytesize_rounded_down, 4096);
let err = unsafe { hipMemsetD32(dev_ptr, 0, bytes / 1024).into() };
*/
if err != CUresult::CUDA_SUCCESS {
return err;
}
unsafe { *dptr = CUdeviceptr(dev_ptr as usize) };
CUresult::CUDA_SUCCESS
} }
#[cfg_attr(not(test), no_mangle)] #[cfg_attr(not(test), no_mangle)]

1
zluda/src/impl/module.rs
View file

@ -185,6 +185,7 @@ fn compile_amd(
let mut compiler_cmd = Command::new(&clang_exe); let mut compiler_cmd = Command::new(&clang_exe);
compiler_cmd compiler_cmd
.arg(format!("-mcpu={}", device_name)) .arg(format!("-mcpu={}", device_name))
.arg("-ffp-contract=off")
.arg("-nogpulib") .arg("-nogpulib")
.arg("-mno-wavefrontsize64") .arg("-mno-wavefrontsize64")
.arg("-O3") .arg("-O3")