共享内存GPU无需CPU-GPU数据拷贝的技术咨询及实例、API支持求证

Great question—this is a super common point of confusion when working with discrete vs. integrated GPUs, so let’s break it down clearly:

• Discrete GPUs: These have their own dedicated VRAM, which is physically separate from your system’s RAM. Since CPU and discrete GPU can’t directly access each other’s memory (they communicate over PCIe), you must explicitly copy data from RAM to VRAM before the GPU can use it, and copy it back afterward if you need the results on the CPU.
• Integrated GPUs (like AMD Ryzen iGPUs): These use a Unified Memory Architecture (UMA), meaning they share the same physical system RAM as the CPU. There’s no separate VRAM pool—your CPU and GPU are accessing the exact same memory region. This means no explicit data copy is required for the GPU to use data the CPU has written, or vice versa.

Absolutely! In UMA setups, the CPU and integrated GPU share a unified virtual address space. Any memory you allocate on the CPU is immediately accessible to the GPU via the same virtual address, and modifications from either side are visible to the other without any intermediate copy steps. This is a huge advantage for low-latency tasks or small data workloads, as it eliminates the overhead of copying data back and forth.

Let’s look at two practical examples—one with SYCL (modern heterogeneous programming standard) and one with OpenCL:

SYCL Example (Zero-Copy Shared Memory) #

#include <sycl/sycl.hpp>
#include <cstdio>

int main() {
    // Create a queue targeting the integrated GPU
    sycl::queue q(sycl::gpu_selector_v);
    printf("Using device: %s\n", q.get_device().get_info<sycl::info::device::name>().c_str());

    // Allocate shared memory (accessible by both CPU and GPU, no copy needed)
    constexpr size_t array_size = 1024;
    auto shared_data = sycl::malloc_shared<int>(array_size, q);

    // CPU writes to the shared memory
    for (size_t i = 0; i < array_size; ++i) {
        shared_data[i] = static_cast<int>(i);
    }

    // GPU kernel modifies the same memory directly
    q.parallel_for(array_size, [=](sycl::id<1> idx) {
        shared_data[idx] *= 2;
    }).wait();

    // CPU reads the modified data directly from the same memory
    printf("First 10 modified values:\n");
    for (size_t i = 0; i < 10; ++i) {
        printf("shared_data[%zu] = %d\n", i, shared_data[i]);
    }

    // Cleanup
    sycl::free(shared_data, q);
    return 0;
}

Here, sycl::malloc_shared allocates memory mapped to both CPU and GPU address spaces. The CPU writes to it, the GPU modifies it, and the CPU reads the results—no sycl::copy calls required at all.

OpenCL Example (Zero-Copy Shared Memory) #

#include <CL/cl.h>
#include <cstdio>

int main() {
    cl_platform_id platform;
    cl_device_id device;
    cl_context context;
    cl_command_queue queue;

    // Grab the first integrated GPU device
    clGetPlatformIDs(1, &platform, nullptr);
    clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, nullptr);
    
    char device_name[256];
    clGetDeviceInfo(device, CL_DEVICE_NAME, sizeof(device_name), device_name, nullptr);
    printf("Using device: %s\n", device_name);

    // Create context and command queue
    context = clCreateContext(nullptr, 1, &device, nullptr, nullptr, nullptr);
    queue = clCreateCommandQueue(context, device, 0, nullptr);

    // Allocate host memory (will be shared with GPU)
    constexpr size_t array_size = 1024;
    int* host_data = static_cast<int*>(malloc(array_size * sizeof(int)));

    // Initialize data on CPU
    for (size_t i = 0; i < array_size; ++i) {
        host_data[i] = static_cast<int>(i);
    }

    // Create OpenCL buffer that uses the existing host memory (no copy)
    cl_mem shared_buffer = clCreateBuffer(
        context, 
        CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR, 
        array_size * sizeof(int), 
        host_data, 
        nullptr
    );

    // Kernel source: multiply each element by 2
    const char* kernel_source = R"(
        __kernel void multiply_by_two(__global int* data) {
            int idx = get_global_id(0);
            data[idx] *= 2;
        }
    )";

    // Build program and kernel
    cl_program program = clCreateProgramWithSource(context, 1, &kernel_source, nullptr, nullptr);
    clBuildProgram(program, 1, &device, nullptr, nullptr, nullptr);
    cl_kernel kernel = clCreateKernel(program, "multiply_by_two", nullptr);

    // Set kernel argument
    clSetKernelArg(kernel, 0, sizeof(cl_mem), &shared_buffer);

    // Run the kernel
    size_t global_work_size = array_size;
    clEnqueueNDRangeKernel(queue, kernel, 1, nullptr, &global_work_size, nullptr, 0, nullptr, nullptr);
    clFinish(queue);

    // CPU reads modified data directly (no clEnqueueReadBuffer needed)
    printf("First 10 modified values:\n");
    for (size_t i = 0; i < 10; ++i) {
        printf("host_data[%zu] = %d\n", i, host_data[i]);
    }

    // Cleanup
    clReleaseKernel(kernel);
    clReleaseProgram(program);
    clReleaseMemObject(shared_buffer);
    clReleaseCommandQueue(queue);
    clReleaseContext(context);
    free(host_data);
    return 0;
}

Using CL_MEM_USE_HOST_PTR tells OpenCL to reuse the existing host memory directly for the GPU buffer. No clEnqueueWriteBuffer or clEnqueueReadBuffer calls are needed—changes from the GPU are immediately visible to the CPU, and vice versa.

Both standards fully support zero-copy shared memory on UMA devices like AMD Ryzen iGPUs:

• OpenCL: You use memory flags like CL_MEM_USE_HOST_PTR (reuse existing host memory) or CL_MEM_ALLOC_HOST_PTR (allocate memory accessible by both CPU and GPU) to create zero-copy buffers. On UMA devices, these flags bypass any data copying—on discrete GPUs, they may still trigger copies (though some discrete GPUs support unified memory too).
• SYCL: The standard provides sycl::malloc_shared (cross-device shared memory) and sycl::malloc_host (host-accessible memory that devices can directly access) to enable zero-copy operations. SYCL’s unified memory model is designed to abstract away these details, making it easy to write code that works efficiently on both integrated and discrete GPUs.

内容的提问来源于stack exchange，提问作者Pietro