Skip to content

Commit

Permalink
Make misc fixes (#41)
Browse files Browse the repository at this point in the history 
* Update ze_loader.lib to the newest version
* Export _ptsz/_ptds for which we have a legacy stream implementations
* Stop producing build logs if we are not looking at them anyway
  • Loading branch information
@vosen
vosen committed Feb 22, 2021
1 parent ab690c6 commit a906c35
Show file tree
Hide file tree
Showing 6 changed files with 124 additions and 8 deletions.
Binary file added level_zero-sys/lib/ze_loader.def
View file
Binary file not shown.
Binary file modified level_zero-sys/lib/ze_loader.lib
View file
Binary file not shown.
46 changes: 42 additions & 4 deletions level_zero/src/ze.rs
View file
Original file line number Diff line number Diff line change
Expand Up @@ -270,7 +270,7 @@ impl Module {
};
match ocl_core::get_program_info(&ocl_program, ocl_core::ProgramInfo::Binaries) {
Ok(ocl_core::ProgramInfoResult::Binaries(binaries)) => {
let (module, build_log) = Self::build_native(ctx, d, &binaries[0]);
let (module, build_log) = Self::build_native_logged(ctx, d, &binaries[0]);
(module, Some(build_log))
}
_ => return (Err(sys::ze_result_t::ZE_RESULT_ERROR_UNKNOWN), None),
Expand Down Expand Up @@ -346,12 +346,21 @@ impl Module {
d: &Device,
bin: &[u8],
opts: Option<&CStr>,
) -> (Result<Self>, BuildLog) {
) -> Result<Self> {
Module::new(ctx, true, d, bin, opts)
}

pub fn build_native(ctx: &mut Context, d: &Device, bin: &[u8]) -> (Result<Self>, BuildLog) {
Module::new(ctx, false, d, bin, None)
pub fn build_spirv_logged(
ctx: &mut Context,
d: &Device,
bin: &[u8],
opts: Option<&CStr>,
) -> (Result<Self>, BuildLog) {
Module::new_logged(ctx, true, d, bin, opts)
}

pub fn build_native_logged(ctx: &mut Context, d: &Device, bin: &[u8]) -> (Result<Self>, BuildLog) {
Module::new_logged(ctx, false, d, bin, None)
}

fn new(
Expand All @@ -360,6 +369,35 @@ impl Module {
d: &Device,
bin: &[u8],
opts: Option<&CStr>,
) -> Result<Self> {
let desc = sys::ze_module_desc_t {
stype: sys::ze_structure_type_t::ZE_STRUCTURE_TYPE_MODULE_DESC,
pNext: ptr::null(),
format: if spirv {
sys::ze_module_format_t::ZE_MODULE_FORMAT_IL_SPIRV
} else {
sys::ze_module_format_t::ZE_MODULE_FORMAT_NATIVE
},
inputSize: bin.len(),
pInputModule: bin.as_ptr(),
pBuildFlags: opts.map(|s| s.as_ptr() as *const _).unwrap_or(ptr::null()),
pConstants: ptr::null(),
};
let mut result: sys::ze_module_handle_t = ptr::null_mut();
let err = unsafe { sys::zeModuleCreate(ctx.0, d.0, &desc, &mut result, ptr::null_mut()) };
if err != crate::sys::ze_result_t::ZE_RESULT_SUCCESS {
Result::Err(err)
} else {
Ok(Module(result))
}
}

fn new_logged(
ctx: &mut Context,
spirv: bool,
d: &Device,
bin: &[u8],
opts: Option<&CStr>,
) -> (Result<Self>, BuildLog) {
let desc = sys::ze_module_desc_t {
stype: sys::ze_structure_type_t::ZE_STRUCTURE_TYPE_MODULE_DESC,
Expand Down
4 changes: 2 additions & 2 deletions ptx/src/test/spirv_run/mod.rs
View file
Original file line number Diff line number Diff line change
Expand Up @@ -116,7 +116,7 @@ test_ptx!(cos, [std::f32::consts::PI], [-1f32]);
test_ptx!(lg2, [512f32], [9f32]);
test_ptx!(ex2, [10f32], [1024f32]);
test_ptx!(cvt_rni, [9.5f32, 10.5f32], [10f32, 10f32]);
test_ptx!(cvt_rzi, [-13.8f32, 12.9f32], [-13f32, 13f32]);
test_ptx!(cvt_rzi, [-13.8f32, 12.9f32], [-13f32, 12f32]);
test_ptx!(cvt_s32_f32, [-13.8f32, 12.9f32], [-13i32, 13i32]);
test_ptx!(clz, [0b00000101_00101101_00010011_10101011u32], [5u32]);
test_ptx!(popc, [0b10111100_10010010_01001001_10001010u32], [14u32]);
Expand Down Expand Up @@ -225,7 +225,7 @@ fn run_spirv<
Some(module.build_options.as_c_str()),
),
None => {
let (module, log) = ze::Module::build_spirv(
let (module, log) = ze::Module::build_spirv_logged(
&mut ctx,
&dev,
byte_il,
Expand Down
80 changes: 79 additions & 1 deletion zluda/src/cuda.rs
View file
Original file line number Diff line number Diff line change
Expand Up @@ -2454,6 +2454,7 @@ pub extern "C" fn cuModuleLoadData(
r#impl::module::load_data(module.decuda(), image).encuda()
}

// TODO: parse jit options
#[cfg_attr(not(test), no_mangle)]
pub extern "C" fn cuModuleLoadDataEx(
module: *mut CUmodule,
Expand All @@ -2462,7 +2463,7 @@ pub extern "C" fn cuModuleLoadDataEx(
options: *mut CUjit_option,
optionValues: *mut *mut ::std::os::raw::c_void,
) -> CUresult {
r#impl::unimplemented()
r#impl::module::load_data(module.decuda(), image).encuda()
}

#[cfg_attr(not(test), no_mangle)]
Expand Down Expand Up @@ -2736,6 +2737,16 @@ pub extern "C" fn cuMemcpyHtoD_v2(
r#impl::memory::copy_v2(dstDevice.decuda(), srcHost, ByteCount).encuda()
}

// TODO: implement default stream semantics
#[cfg_attr(not(test), no_mangle)]
pub extern "C" fn cuMemcpyHtoD_v2_ptds(
dstDevice: CUdeviceptr,
srcHost: *const ::std::os::raw::c_void,
ByteCount: usize,
) -> CUresult {
r#impl::memory::copy_v2(dstDevice.decuda(), srcHost, ByteCount).encuda()
}

#[cfg_attr(not(test), no_mangle)]
pub extern "C" fn cuMemcpyDtoH_v2(
dstHost: *mut ::std::os::raw::c_void,
Expand All @@ -2745,6 +2756,16 @@ pub extern "C" fn cuMemcpyDtoH_v2(
r#impl::memory::copy_v2(dstHost, srcDevice.decuda(), ByteCount).encuda()
}

// TODO: implement default stream semantics
#[cfg_attr(not(test), no_mangle)]
pub extern "C" fn cuMemcpyDtoH_v2_ptds(
dstHost: *mut ::std::os::raw::c_void,
srcDevice: CUdeviceptr,
ByteCount: usize,
) -> CUresult {
r#impl::memory::copy_v2(dstHost, srcDevice.decuda(), ByteCount).encuda()
}

#[cfg_attr(not(test), no_mangle)]
pub extern "C" fn cuMemcpyDtoD_v2(
dstDevice: CUdeviceptr,
Expand Down Expand Up @@ -2926,6 +2947,16 @@ pub extern "C" fn cuMemsetD8_v2(
r#impl::memory::set_d8_v2(dstDevice.decuda(), uc, N).encuda()
}

// TODO: implement default stream semantics
#[cfg_attr(not(test), no_mangle)]
pub extern "C" fn cuMemsetD8_v2_ptds(
dstDevice: CUdeviceptr,
uc: ::std::os::raw::c_uchar,
N: usize,
) -> CUresult {
r#impl::memory::set_d8_v2(dstDevice.decuda(), uc, N).encuda()
}

#[cfg_attr(not(test), no_mangle)]
pub extern "C" fn cuMemsetD16_v2(
dstDevice: CUdeviceptr,
Expand All @@ -2944,6 +2975,16 @@ pub extern "C" fn cuMemsetD32_v2(
r#impl::memory::set_d32_v2(dstDevice.decuda(), ui, N).encuda()
}

// TODO: implement default stream semantics
#[cfg_attr(not(test), no_mangle)]
pub extern "C" fn cuMemsetD32_v2_ptds(
dstDevice: CUdeviceptr,
ui: ::std::os::raw::c_uint,
N: usize,
) -> CUresult {
r#impl::memory::set_d32_v2(dstDevice.decuda(), ui, N).encuda()
}

#[cfg_attr(not(test), no_mangle)]
pub extern "C" fn cuMemsetD2D8_v2(
dstDevice: CUdeviceptr,
Expand Down Expand Up @@ -3322,6 +3363,12 @@ pub extern "C" fn cuStreamGetCtx(hStream: CUstream, pctx: *mut CUcontext) -> CUr
r#impl::stream::get_ctx(hStream.decuda(), pctx.decuda()).encuda()
}

// TODO: implement default stream semantics
#[cfg_attr(not(test), no_mangle)]
pub extern "C" fn cuStreamGetCtx_ptsz(hStream: CUstream, pctx: *mut CUcontext) -> CUresult {
r#impl::stream::get_ctx(hStream.decuda(), pctx.decuda()).encuda()
}

#[cfg_attr(not(test), no_mangle)]
pub extern "C" fn cuStreamWaitEvent(
hStream: CUstream,
Expand Down Expand Up @@ -3630,6 +3677,37 @@ pub extern "C" fn cuLaunchKernel(
.encuda()
}

// TODO: implement default stream semantics
#[cfg_attr(not(test), no_mangle)]
pub extern "C" fn cuLaunchKernel_ptsz(
f: CUfunction,
gridDimX: ::std::os::raw::c_uint,
gridDimY: ::std::os::raw::c_uint,
gridDimZ: ::std::os::raw::c_uint,
blockDimX: ::std::os::raw::c_uint,
blockDimY: ::std::os::raw::c_uint,
blockDimZ: ::std::os::raw::c_uint,
sharedMemBytes: ::std::os::raw::c_uint,
hStream: CUstream,
kernelParams: *mut *mut ::std::os::raw::c_void,
extra: *mut *mut ::std::os::raw::c_void,
) -> CUresult {
r#impl::function::launch_kernel(
f.decuda(),
gridDimX,
gridDimY,
gridDimZ,
blockDimX,
blockDimY,
blockDimZ,
sharedMemBytes,
hStream.decuda(),
kernelParams,
extra,
)
.encuda()
}

#[cfg_attr(not(test), no_mangle)]
pub extern "C" fn cuLaunchCooperativeKernel(
f: CUfunction,
Expand Down
2 changes: 1 addition & 1 deletion zluda/src/impl/module.rs
View file
Original file line number Diff line number Diff line change
Expand Up @@ -87,7 +87,7 @@ impl SpirvModule {
};
let l0_module = match self.should_link_ptx_impl {
None => {
l0::Module::build_spirv(ctx, dev, byte_il, Some(self.build_options.as_c_str())).0
l0::Module::build_spirv(ctx, dev, byte_il, Some(self.build_options.as_c_str()))
}
Some(ptx_impl) => {
l0::Module::build_link_spirv(
Expand Down

0 comments on commit a906c35

Please sign in to comment.