cuIpcGetMemHandle results in CUDA_ERROR_INVALID_VALUE

[genevanmeter]

Trying to get the ipc memory handle for a CudaSlice results in a CUDA_ERROR_INVALID_VALUE.

use cudarc::driver::{CudaContext, CudaSlice, DevicePtr, DriverError, sys::CUipcMemHandle};
use cudarc::driver::sys::cuIpcGetMemHandle;
use std::mem::MaybeUninit;

fn main() -> Result<(), DriverError> {
    let ctx = CudaContext::new(0)?;
    let stream = ctx.default_stream();

    let a: CudaSlice<f64> = stream.alloc_zeros::<f64>(10)?;

    let mut handle = MaybeUninit::<CUipcMemHandle>::uninit();

    let (src, _record_src) = a.device_ptr(&stream);

    unsafe {
        let r = cuIpcGetMemHandle(handle.as_mut_ptr(), src);
        println!("cuIpcGetMemHandle result: {:?}", r);
    }

    Ok(())
}

Result:

cuIpcGetMemHandle result: CUDA_ERROR_INVALID_VALUE

Using the unsafe runtime does work

use cudarc::runtime::sys::{cudaIpcGetMemHandle, cudaIpcMemHandle_st, cudaMalloc};
use std::{io::Error, os::raw::c_void};

fn main() -> Result<(), Error> {

    let mut devPtr: *mut c_void = std::ptr::null_mut();
    unsafe {
        cudaMalloc(&mut devPtr, 10);
    }

    let mut handle = cudaIpcMemHandle_st { reserved: [0;64]};

    unsafe {
        let r = cudaIpcGetMemHandle(&mut handle, devPtr);
        println!("cuIpcGetMemHandle result: {:?}", r);
    }

    println!("Handle: {:?}", handle);

    Ok(())
}

cuIpcGetMemHandle result: cudaSuccess
Handle: cudaIpcMemHandle_st { reserved: [48, 117, -81, 4, 41, 86, 0, 0, 76, 108, 48, 0, 0, 0, 0, 0, 10, 0, 0, 0, 0, 0, 0, 0, 0, 2, 0, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 45, 0, 0, 0, 0, 0, 0, 0, -108, -43, -36, -63, -126, 0, 0, 92, 0, 0, 0, 0, 0, 0, 0, 0] }

[chelsea0x3b]

Okay so the reason this fails is because the stream.alloc_zeros() calls sys::cuMemAllocAsync under the hood. If you call sys::cuMemAlloc_v2, the call to ipc get mem handle works:

use cudarc::driver::{CudaContext, CudaSlice, DevicePtr, DriverError, sys::CUipcMemHandle};
use cudarc::driver::sys::{cuIpcGetMemHandle, self};
use std::mem::MaybeUninit;

fn main() -> Result<(), DriverError> {
    let ctx = CudaContext::new(0)?;
    let stream = ctx.default_stream();

    let mut src = 0;

    // note: WORKS
    unsafe { sys::cuMemAlloc_v2(&mut src, 10) }.result()?;

    // note: DOESN'T WORK
    // unsafe { sys::cuMemAllocAsync(&mut src, 10, stream.cu_stream()) }.result()?;

    let mut handle = CUipcMemHandle { reserved: [0;64]};

    unsafe {
        let r = cuIpcGetMemHandle(&mut handle, src);
        println!("cuIpcGetMemHandle result: {:?}", r);
    }

    Ok(())
}

[chelsea0x3b]

A workaround for this currently is to allocate using cuMemAlloc_v2 and then upgrade it to a cuda slice:

    let mut src = 0;
    unsafe { sys::cuMemAlloc_v2(&mut src, 10 * std::mem::size_of::<f32>()) }.result()?;
    let a = unsafe { stream.upgrade_device_ptr::<f32>(src, 10) };

@genevanmeter is that sufficient for you? Otherwise we could potentially add something like stream.blocking_alloc to call the cuMemAlloc_v2 under the hood.

[genevanmeter]

Yes I can work with that.

Thank you for looking into this and the explanation.