Raw access to the shared memory pointer and length

[Gui-Yom]

I need the publisher shared memory pages to be resident in RAM for a dma copy. How do I get access to a data segment pointer and length to run the necessary syscalls for my needs ? Here are 4 ways I can think of :

1. Run the necessary logic on each loaned sample's pointer. Unfortunately the required process takes a lot of time because it musts zero out memory and the buffers are big. I can't afford to block the pipeline for tens of milliseconds for each previously unseen unpopulated memory region.
2. Run a warmup round (cooperatively with the subscribers) and lock all samples memory ahead of time. The maximum number of samples is :

   iceoryx2/iceoryx2/src/service/static_config/publish_subscribe.rs

   Line 79 in de649db

   pub(crate) fn required_amount_of_samples_per_data_segment(

3. Populate and lock the pages after creating the publisher by getting the base pointer somehow (requires changes to the API)
4. Create my own service variant (from ipc::Service) and customize the shared memory implementation (code duplication ?)

The last option is probably the cleanest and most realistic. Is there a better solution to get raw access the shared memory mapping ?

[elfenpiff]

@Gui-Yom iceoryx2 is able to handle different kinds of memory than just POSIX shared memory. Could you provide me with some more details about your system? Are you using Linux and want to use, for instance, /dev/dma_heap as shared memory? Also, what kind of syscalls would you like to use if you got access to the shared memory mapping? From your question it sounds a bit like you would like to call mlock on that memory region?

[Gui-Yom]

My use case is copying GPU memory directly into an iceoryx sample. My environment is Linux and CUDA. Here is a sample code :

loan = publisher.loan_slice_uninit(size)
cudaMemcpy(loan.payload_ptr, buffer_d, size, cudaMemcpyDeviceToHost)
loan.assume_init().send()

The cuda driver requires for performance that the host memory pages are pinned (fixed and resident in RAM for the duration of the transfer and registered internally) by calling cudaHostRegister. The physical RAM space is only lazily allocated by the kernel on first access. Calling cudaHostRegister has a tremendous cost (100-200~ ms for 250MB) when registering a sample's memory for the first time because it must page fault and zero-clear all associated memory pages.
I'm pretty sure that this is due to the pages lazy allocation because calling madvise(MADV_POPULATE_WRITE) before cudaHostRegister moves the cost to the madvise call and cudaHostRegister then takes only a few milliseconds.

Just to say, I want to run a cudaHostRegister or at least madvise(MADV_POPULATE_WRITE) on the whole shared memory segment before publishing samples as fast as I can.

[elfenpiff]

@Gui-Yom, this is something we tried some years ago and failed. As far as I understand, the memory must be non-swapable and contiguous so that it can be used for CUDA. But the memory the linux kernel provides is never contiguous except when you allocate it with the contiguous memory allocator (which is a kernel driver). Maybe /dev/dma_heap works but I am not sure.

But there are good news for you. We are currently working on cuda so that this is natively supported in iceoryx2. The idea is that you create a cuda service then like that:

let node = NodeBuilder::new().create::<cuda::Service>()?;

let service = node
    .service_builder(&"My/Funk/ServiceName".try_into()?)
    .publish_subscribe::<[u8]>()
    .open_or_create()?;

let publisher = service.publisher_builder().create()?;

//This would then work as you already envisioned it
let loan = publisher.loan_slice_uninit(size)
cudaMemcpy(loan.payload_ptr, buffer_d, size, cudaMemcpyDeviceToHost)
loan.assume_init().send()

[Gui-Yom]

I don't think there are any contiguity requirements, in fact I got it to work using a cudaHostRegister on the payload ptr and size before calling cudaMemcpy (without hugepages nor transparent_hugepages).
I currently settled on a quick workaround by getting the shm pointer from a fake sample (that I don't send) and deducing the shm size using the internal

iceoryx2/iceoryx2/src/service/static_config/publish_subscribe.rs

Line 79 in de649db

pub(crate) fn required_amount_of_samples_per_data_segment(

I then call cudaHostRegister on the memory range just after creating the publisher. (Yes this is dirty and brittle)

I'm trying to write a service variant to accommodate my use case, but I guess an official cuda::Service is great news. I suppose it's able to automatically handle cudaIpcGetMemHandle calls and the like automatically ?

[elfenpiff]

Yes exactly. It handles everything for you, it can even send events via cudaIpcEventHandle.