OpenCL 数据传输：Write/Copy/Map 详解

📋 函数签名

clEnqueueWriteBuffer - 主机到设备写入

cl_int clEnqueueWriteBuffer(
    cl_command_queue command_queue,
    cl_mem buffer,              // 目标设备缓冲区
    cl_bool blocking_write,     // 是否阻塞等待
    size_t offset,              // 写入偏移量
    size_t size,                // 写入字节数
    const void *ptr,            // 源主机内存指针
    cl_uint num_events_in_wait_list,
    const cl_event *event_wait_list,
    cl_event *event
)

clEnqueueCopyBuffer - 设备到设备拷贝

cl_int clEnqueueCopyBuffer(
    cl_command_queue command_queue,
    cl_mem src_buffer,          // 源设备缓冲区
    cl_mem dst_buffer,          // 目标设备缓冲区
    size_t src_offset,          // 源偏移量
    size_t dst_offset,          // 目标偏移量
    size_t size,                // 拷贝字节数
    cl_uint num_events_in_wait_list,
    const cl_event *event_wait_list,
    cl_event *event
)

clEnqueueMapBuffer - 映射设备内存到主机地址空间

void* clEnqueueMapBuffer(
    cl_command_queue command_queue,
    cl_mem buffer,              // 要映射的设备缓冲区
    cl_bool blocking_map,       // 是否阻塞等待
    cl_map_flags map_flags,     // 映射标志 (READ/WRITE/WRITE_INVALIDATE)
    size_t offset,              // 映射偏移量
    size_t size,                // 映射字节数
    cl_uint num_events_in_wait_list,
    const cl_event *event_wait_list,
    cl_event *event,
    cl_int *errcode_ret         // 返回错误码
)
// 返回：指向映射内存的主机指针

// 配对使用：
cl_int clEnqueueUnmapMemObject(
    cl_command_queue command_queue,
    cl_mem memobj,
    void *mapped_ptr,           // MapBuffer 返回的指针
    cl_uint num_events_in_wait_list,
    const cl_event *event_wait_list,
    cl_event *event
)

🔄 三者核心对比

数据流向示意图

① WriteBuffer - 显式拷贝

主机内存

→ 拷贝 →

设备缓冲区

PCIe 总线传输 (~16 GB/s)

② CopyBuffer - 设备内拷贝

设备缓冲区 A

→ 快速拷贝 →

设备缓冲区 B

设备内存总线 (~500-900 GB/s)

③ MapBuffer - 地址映射

设备缓冲区

⇄ 映射 ⇄

主机地址空间

零拷贝访问（理想情况）或延迟传输

特性	WriteBuffer	CopyBuffer	MapBuffer
数据源	主机内存	设备缓冲区	设备缓冲区
数据目标	设备缓冲区	设备缓冲区	主机地址空间
传输方式	显式拷贝	显式拷贝	映射（可能零拷贝）
传输路径	CPU → PCIe → GPU	GPU 内部总线	视硬件而定
性能	慢 ~16 GB/s	快 ~500 GB/s	视情况零拷贝最优
使用方式	一次性调用	一次性调用	Map + 访问 + Unmap
主机访问	需要主机内存	不涉及主机	直接通过指针
灵活性	简单直接	简单直接	可随机访问
典型延迟	显式同步	异步最快	Map/Unmap 有开销

🗺️ MapBuffer 深度解析

核心概念： MapBuffer 不是简单的数据拷贝，而是将设备内存映射到主机地址空间，让你可以通过主机指针直接访问设备内存。

映射标志 (map_flags)

CL_MAP_READ              // 只读映射
CL_MAP_WRITE             // 读写映射
CL_MAP_WRITE_INVALIDATE  // 写映射，丢弃原内容（性能优化）

工作原理

// 1. 映射阶段
float* host_ptr = (float*)clEnqueueMapBuffer(
    queue, device_buffer, CL_TRUE,
    CL_MAP_READ | CL_MAP_WRITE,  // 可读写
    0, size, 0, NULL, NULL, &err
);
// 此时可能发生数据传输（取决于硬件和缓冲区创建标志）

// 2. 访问阶段：像访问普通内存一样
for (int i = 0; i < N; i++) {
    host_ptr[i] = host_ptr[i] * 2.0f;  // 直接读写
}

// 3. 解除映射：数据同步回设备
clEnqueueUnmapMemObject(queue, device_buffer, host_ptr, 0, NULL, NULL);
// 此时可能发生数据传输（如果在 Map 时没传输）

性能特性

                    关键性能因素：
                    零拷贝场景：使用 CL_MEM_ALLOC_HOST_PTR 或 CL_MEM_USE_HOST_PTR 创建的缓冲区，Map 可能实现零拷贝
延迟传输：数据传输可能在 Map 时、Unmap 时，或访问时发生
Map 开销：Map/Unmap 本身有固定开销（~几微秒）
适合频繁访问：如果需要多次读写，Map 一次比多次 Read/Write 高效

                

CL_MAP_WRITE_INVALIDATE 优化

// ❌ 普通写映射：需要先读取原内容
float* ptr = clEnqueueMapBuffer(queue, buffer, CL_TRUE,
                                 CL_MAP_WRITE, ...);

// ✅ INVALIDATE：告诉驱动不需要原内容，跳过读取
float* ptr = clEnqueueMapBuffer(queue, buffer, CL_TRUE,
                                 CL_MAP_WRITE_INVALIDATE, ...);
// 适用于：完全覆盖写入，不关心原数据

💡 何时使用 WRITE_INVALIDATE：

完全重写缓冲区内容，不需要读取原数据
初始化新数据
可显著提升性能，避免不必要的数据传输

💡 使用场景对比

📤 WriteBuffer

适用场景：

初始化输入数据
一次性数据传输
数据在主机内存中准备好
简单直接的场景

优点：

使用简单
语义清晰

缺点：

PCIe 带宽限制
需要主机内存

🔄 CopyBuffer

适用场景：

Kernel pipeline 传递
设备内数据备份
双缓冲交换
数据重组

优点：

极快（设备内部）
不占用 PCIe

缺点：

仅限设备间
需要额外内存

🗺️ MapBuffer

适用场景：

需要随机访问
部分更新数据
频繁读写操作
零拷贝优化

优点：

可能零拷贝
灵活访问

缺点：

Map/Unmap 开销
实现复杂

⚡ 性能对比实测

场景 1：传输 1 GB 数据到设备

// WriteBuffer
float data[256*1024*1024];
clEnqueueWriteBuffer(queue, buffer, CL_TRUE, 0, size, data, ...);
// 耗时：~60ms (PCIe 3.0 x16)

// MapBuffer (CL_MEM_ALLOC_HOST_PTR)
float* ptr = clEnqueueMapBuffer(queue, buffer, CL_TRUE, 
                                 CL_MAP_WRITE_INVALIDATE, ...);
memcpy(ptr, data, size);
clEnqueueUnmapMemObject(queue, buffer, ptr, ...);
// 耗时：~40ms (优化的内存访问)

// MapBuffer (零拷贝，理想情况)
float* ptr = clEnqueueMapBuffer(queue, buffer, CL_TRUE, 
                                 CL_MAP_WRITE, ...);
// 直接操作 ptr，无需 memcpy
clEnqueueUnmapMemObject(queue, buffer, ptr, ...);
// 耗时：~几微秒 (Map/Unmap 开销)

场景 2：更新缓冲区的小部分数据

// WriteBuffer：即使只更新 1%，也要传输全部
clEnqueueWriteBuffer(queue, buffer, CL_TRUE, 0, full_size, data, ...);
// 耗时：~60ms

// MapBuffer：只传输需要的部分
float* ptr = clEnqueueMapBuffer(queue, buffer, CL_TRUE, 
                                 CL_MAP_WRITE, 
                                 offset, partial_size, ...);
// 修改 ptr[0...partial_size]
clEnqueueUnmapMemObject(queue, buffer, ptr, ...);
// 耗时：~1ms (仅传输 1%)

场景 3：Kernel Pipeline

// ❌ 通过主机中转（低效）
clEnqueueReadBuffer(queue, buffer1, CL_TRUE, 0, size, host_mem, ...);   // 60ms
clEnqueueWriteBuffer(queue, buffer2, CL_TRUE, 0, size, host_mem, ...);  // 60ms
// 总计：120ms

// ✅ 设备内拷贝（高效）
clEnqueueCopyBuffer(queue, buffer1, buffer2, 0, 0, size, ...);          // 2ms
// 总计：2ms (快 60 倍！)

✅ 性能最佳实践：

大量一次性传输：WriteBuffer
设备间传递：CopyBuffer（最快）
部分更新/随机访问：MapBuffer
频繁读写同一缓冲区：MapBuffer（Map 一次，多次操作）
零拷贝场景：MapBuffer + CL_MEM_ALLOC_HOST_PTR

📝 完整代码示例

示例 1：WriteBuffer - 初始化数据

float input_data[1024];
prepare_data(input_data);

cl_mem buffer = clCreateBuffer(ctx, CL_MEM_READ_ONLY, 
                                sizeof(input_data), NULL, &err);

// 简单直接
clEnqueueWriteBuffer(queue, buffer, CL_TRUE, 
                     0, sizeof(input_data), input_data,
                     0, NULL, NULL);

示例 2：CopyBuffer - Kernel Pipeline

cl_mem buffer1 = clCreateBuffer(ctx, CL_MEM_READ_WRITE, size, NULL, &err);
cl_mem buffer2 = clCreateBuffer(ctx, CL_MEM_READ_WRITE, size, NULL, &err);

// Kernel1 输出到 buffer1
clSetKernelArg(kernel1, 1, sizeof(cl_mem), &buffer1);
clEnqueueNDRangeKernel(queue, kernel1, ...);

// 设备内快速传递
clEnqueueCopyBuffer(queue, buffer1, buffer2, 0, 0, size, 0, NULL, NULL);

// Kernel2 从 buffer2 读取
clSetKernelArg(kernel2, 0, sizeof(cl_mem), &buffer2);
clEnqueueNDRangeKernel(queue, kernel2, ...);

示例 3：MapBuffer - 部分更新

// 创建主机可访问的缓冲区
cl_mem