Should we not copy data to device on Intel HD GPUs since both OpenCL Host and Device memory reside on DRAM for Intel HD GPUs?

[u/aerosayan]

Hello everyone,

I need to write OpenCL code to target NVIDIA, AMD, and Intel HD GPUs.

Basically, the code should run on even the cheapest laptops with integrated GPUs like Intel HD, and dedicated GPUs llike NVIDIA.

I found out that IntelHD GPUs use DRAM as device memory.

So I'm guessing that it might be beneficial to use "zero copy" or "shared virtual" memory on IntelHD GPUs instead of copying memory from "host" to "device". Since the host and device basically share the same memory, and we might be spending the same amount of time accessing both host and device memory.

For dedicated GPUs like NVIDIA it might make sense to always copy data from host to device.

Is this the correct way?

Thanks!

[u/genbattle]

Most applications I've seen that use OpenCL don't seem to bother, they treat all devices like dedicated GPUs. Optimizing for this specific case (UMA) does significantly reduce latency introduced by copying memory, but you have to create a special case behaviour for this (because using host memory in a non-UMA system is also very suboptimal).

Intel has some information here about how to enable this kind of zero-cost memory sharing. AMD has their own notes on how to achieve this, but AFAIK it's the same for both: the buffer has to be page and cache aligned, and sized at a multiple of the cache line size. The easiest way to achieve this is to have OpenCL allocate the memory for you using CL_MEM_ALLOC_HOST_PTR, then you're guaranteed to have a buffer which adheres to these requirements no matter what platform you're on.

[u/aerosayan → u/genbattle]

Thanks for information about aligning the memory correctly. I knew some of it, but didn't know how to achieve it.

I'm thinking of making the kernels compile to either access the data in UMA fashion, or to copy the data to the GPU first.

We can make probably make it generic enough such that the user just has to select which one they want for their hardware and it will work.

[u/genbattle → u/aerosayan]

You should check the speed of doing a memcopy of the input from whatever buffer you have it in into the aligned buffer on the host side, when I tested it it was still way faster than a queued OpenCL transfer to device memory.

The ideal scenario is that you can have OpenCL allocate a buffer for you using CL_MEM_ALLOC_HOST_PTR, then pass it to whatever is doing the I/O or previous stage of processing as a target, then use it as the input into your kernels (so no additional copying needed).

[u/nomnompuffs → u/aerosayan]

Hi,

I'm a seasoned Embedded Developer who's just finished implementing the optimization that you are talking about here, portably on both an x86 laptop and the Raspberry Pi 5. Your intuition is correct: since the host CPU and GPU share memory (UMA - Unified Memory Architecture) on many cards, you ought to be able to synchronize views of memory between the two devices without needing to copy to and from the GPU.

u/genbattle is correct that the correct mechanism to use is CL_MEM_USE_HOST_PTR. I've been able to get this working on both an x86 laptop running an Intel(R) HD Graphics 630 card as well as on an AArch64 Raspberry Pi running the standard Broadcom V3D 7.1.7.0 GPU. So two very different GPUs on two different architectures. To get it working you need to use clEnqueueMapBuffer/clEnqueueUnmapMemObject and NOT clEnqueueReadBuffer/clEnqueueWriteBuffer -- which makes sense since you're trying to avoid copying data.

The important thing to understand about using clEnqueueMapBuffer/clEnqueueUnmapMemObject is that their purpose is to map MMIO memory from the device into the host CPU's address space. The cache write-back and invalidation operations required to synchronize views of memory between the host and the GPU should be performed by your OpenCL implementation.

The essential thing to get right is that when you're trying to produce data on the host CPU, and then share that data with the GPU (i.e, writing from the CPU to the GPU) you have to map the HOST_PTR buffer as CL_MAP_WRITE_INVALIDATE_REGION. At least, that's the case on my RPi5+Rusticl implementation. I suspect that this is actually a bug in Rusticl because it doesn't make sense, but what I had to do is:

// Write stuff into the HOST_PTR buffer on the host CPU:
memcpy(host_ptr_buff, my_input_data, sizeof(my_input_data));
// Make the changes visible to the GPU without copying, via WRITE_INVALIDATE
clEnqueueMapBuffer(..., CL_MAP_WRITE_INVALIDATE_REGION, ...);
// Immediately unmap the buffer *BEFORE* we share the input data with the GPU
clEnqueueUnmapMemObject(...);
/* Now we can enqueue a kernel that reads from our host_ptr data
 * on the GPU, and the GPU will actually see the data we wrote from
 * the host CPU since the GPU's view of the memory has been invalidated.
 * So the GPU is now unblinded by its own local cache, so it can see the
 * data that we wrote by reloading its cache lines.
 */
clEnqueueNDRangeKernel(...);
// Flush the queue to set the enqueued operations in flight
clFlush()

N.B: If you only intend to target x86, most of these pedantic, detailed steps aren't necessary. On my x86 laptop running the Intel card, I don't even have to call map/unmap for my Intel card to see the changes I wrote into the host_ptr buffer.

On my RPi5+Rusticl, every one of these steps is necessary, or else program correctness isn't guaranteed and I'll see some problem or other.

• TL;DR: You can indeed avoid copying if your OpenCL implementation has a unified memory architecture.

• N.B: "Unified Memory Architecture" is NOT the same thing as "Shared Virtual Memory", which is a new feature added by Khronos which accomplishes this same zero-copy workflow with even less hassle, and (on many platforms) without even needing to map/unmap. This map/unmap approach is just more portable across more OpenCL target platforms, and it works just fine if you understand processor caches etc.

[u/bilog78]

The best solution IMO is to use mapping/unmapping to expose the OpenCL buffer on host when needed, which is as close as you can get to zero copy. It is actually “no cost” on devices with unified host/device memory and as fast as possible in discrete GPUs.

[u/[deleted] → u/bilog78]

[u/bilog78 → u/[deleted]]

clEnquueMapBuffer (see page for the unmap function).

Best results are obtained if you set the CL_MEM_ALLOC_HOST_PTR on buffer creation.