第6章 Images - 图像对象测试

Images 测试套件验证 OpenCL 图像对象的各种操作，包括图像格式支持、1D/2D/3D 图像及数组、图像读写、拷贝和填充操作。图像对象是 OpenCL 中用于高效访问多维纹理数据的专用内存对象。

6.1 图像格式支持测试

6.1.1 图像格式组成

OpenCL 图像格式由通道顺序（Channel Order）和通道数据类型（Channel Data Type）组成：

typedef struct cl_image_format {
    cl_channel_order image_channel_order;
    cl_channel_data_type image_channel_data_type;
} cl_image_format;

常用通道顺序

通道顺序	通道数	说明
CL_R	1	单通道红色
CL_RG	2	双通道红绿
CL_RGB	3	三通道红绿蓝（部分设备支持）
CL_RGBA	4	四通道红绿蓝透明度
CL_BGRA	4	四通道蓝绿红透明度
CL_INTENSITY	1	强度（所有通道相同）
CL_LUMINANCE	1	亮度
CL_DEPTH	1	深度值
CL_sRGB	3	sRGB 色彩空间
CL_sRGBA	4	sRGBA 色彩空间

常用通道数据类型

数据类型	每通道位数	数值范围
CL_UNORM_INT8	8位	[0.0, 1.0] 归一化
CL_UNORM_INT16	16位	[0.0, 1.0] 归一化
CL_SNORM_INT8	8位	[-1.0, 1.0] 归一化
CL_SNORM_INT16	16位	[-1.0, 1.0] 归一化
CL_SIGNED_INT8	8位	[-128, 127] 整数
CL_SIGNED_INT16	16位	[-32768, 32767] 整数
CL_SIGNED_INT32	32位	整数
CL_UNSIGNED_INT8	8位	[0, 255] 整数
CL_UNSIGNED_INT16	16位	[0, 65535] 整数
CL_UNSIGNED_INT32	32位	整数
CL_HALF_FLOAT	16位	半精度浮点
CL_FLOAT	32位	单精度浮点

6.1.2 查询支持的图像格式

int test_supported_image_formats(cl_context context, cl_mem_object_type image_type)
{
    cl_int error;
    cl_uint num_formats;
    
    // 查询支持的格式数量
    error = clGetSupportedImageFormats(context, CL_MEM_READ_WRITE,
                                      image_type, 0, NULL, &num_formats);
    test_error(error, "Unable to query number of supported formats");
    
    log_info("Device supports %u image formats for type %s\n",
             num_formats, GetImageTypeName(image_type));
    
    // 获取支持的格式列表
    cl_image_format *formats = (cl_image_format*)malloc(
        sizeof(cl_image_format) * num_formats);
    
    error = clGetSupportedImageFormats(context, CL_MEM_READ_WRITE,
                                      image_type, num_formats,
                                      formats, NULL);
    test_error(error, "Unable to get supported formats");
    
    // 打印所有支持的格式
    for (cl_uint i = 0; i < num_formats; i++)
    {
        log_info("  Format %u: %s / %s\n", i,
                 GetChannelOrderName(formats[i].image_channel_order),
                 GetChannelTypeName(formats[i].image_channel_data_type));
    }
    
    free(formats);
    return 0;
}

6.1.3 必需图像格式

OpenCL 规范要求设备至少支持以下图像格式：

const cl_image_format required_formats[] = {
    // RGBA 格式（所有设备必须支持）
    { CL_RGBA, CL_UNORM_INT8 },
    { CL_RGBA, CL_UNORM_INT16 },
    { CL_RGBA, CL_SIGNED_INT8 },
    { CL_RGBA, CL_SIGNED_INT16 },
    { CL_RGBA, CL_SIGNED_INT32 },
    { CL_RGBA, CL_UNSIGNED_INT8 },
    { CL_RGBA, CL_UNSIGNED_INT16 },
    { CL_RGBA, CL_UNSIGNED_INT32 },
    { CL_RGBA, CL_HALF_FLOAT },
    { CL_RGBA, CL_FLOAT },
    
    // BGRA 格式（如果支持 CL_BGRA）
    { CL_BGRA, CL_UNORM_INT8 },
};

int test_min_image_formats(cl_context context)
{
    cl_uint num_formats;
    cl_image_format *formats;
    
    // 获取支持的 2D 图像格式
    clGetSupportedImageFormats(context, CL_MEM_READ_WRITE,
                               CL_MEM_OBJECT_IMAGE2D,
                               0, NULL, &num_formats);
    
    formats = (cl_image_format*)malloc(sizeof(cl_image_format) * num_formats);
    clGetSupportedImageFormats(context, CL_MEM_READ_WRITE,
                               CL_MEM_OBJECT_IMAGE2D,
                               num_formats, formats, NULL);
    
    // 验证必需格式
    for (size_t i = 0; i < ARRAY_SIZE(required_formats); i++)
    {
        bool found = false;
        for (cl_uint j = 0; j < num_formats; j++)
        {
            if (formats[j].image_channel_order == required_formats[i].image_channel_order
                && formats[j].image_channel_data_type == required_formats[i].image_channel_data_type)
            {
                found = true;
                break;
            }
        }
        
        if (!found)
        {
            log_error("ERROR: Required format not supported: %s / %s\n",
                     GetChannelOrderName(required_formats[i].image_channel_order),
                     GetChannelTypeName(required_formats[i].image_channel_data_type));
            free(formats);
            return -1;
        }
    }
    
    log_info("All required image formats are supported\n");
    free(formats);
    return 0;
}

---

6.2 1D/2D/3D 图像测试

6.2.1 1D 图像测试

一维图像用于存储线性数据序列。

创建 1D 图像

int test_1D_image_creation(cl_context context)
{
    cl_int error;
    cl_image_format format;
    cl_image_desc desc;
    
    // 设置图像格式：RGBA，每通道 8 位归一化整数
    format.image_channel_order = CL_RGBA;
    format.image_channel_data_type = CL_UNORM_INT8;
    
    // 设置图像描述符
    memset(&desc, 0, sizeof(desc));
    desc.image_type = CL_MEM_OBJECT_IMAGE1D;
    desc.image_width = 1024;  // 1D 图像宽度
    desc.image_height = 0;    // 1D 图像高度必须为 0
    desc.image_depth = 0;     // 1D 图像深度必须为 0
    desc.image_array_size = 0;
    desc.image_row_pitch = 0;  // 0 表示由实现决定
    desc.image_slice_pitch = 0;
    desc.num_mip_levels = 0;
    desc.num_samples = 0;
    desc.buffer = NULL;
    
    // 创建图像对象
    cl_mem image = clCreateImage(context, CL_MEM_READ_WRITE,
                                &format, &desc, NULL, &error);
    test_error(error, "Unable to create 1D image");
    
    // 查询图像属性
    size_t width;
    error = clGetImageInfo(image, CL_IMAGE_WIDTH, sizeof(width), &width, NULL);
    test_error(error, "Unable to get image width");
    
    if (width != 1024)
    {
        log_error("ERROR: Image width mismatch: expected 1024, got %zu\n", width);
        return -1;
    }
    
    log_info("1D image created successfully: width=%zu\n", width);
    clReleaseMemObject(image);
    return 0;
}

1D 图像读写内核

// 读取 1D 图像
__kernel void test_read_image_1D(__read_only image1d_t src,
                                 sampler_t sampler,
                                 __global float4 *output)
{
    int tid = get_global_id(0);
    
    // 使用整数坐标读取
    int coord = tid;
    float4 pixel = read_imagef(src, sampler, coord);
    
    output[tid] = pixel;
}

// 写入 1D 图像
__kernel void test_write_image_1D(__write_only image1d_t dst,
                                  __global float4 *input)
{
    int tid = get_global_id(0);
    
    int coord = tid;
    float4 pixel = input[tid];
    
    write_imagef(dst, coord, pixel);
}

6.2.2 2D 图像测试

二维图像是最常用的图像类型，用于纹理、图像处理等。

创建 2D 图像

int test_2D_image_creation(cl_context context)
{
    cl_int error;
    cl_image_format format;
    cl_image_desc desc;
    
    // RGBA 8位格式
    format.image_channel_order = CL_RGBA;
    format.image_channel_data_type = CL_UNORM_INT8;
    
    // 设置 2D 图像描述符
    memset(&desc, 0, sizeof(desc));
    desc.image_type = CL_MEM_OBJECT_IMAGE2D;
    desc.image_width = 512;
    desc.image_height = 512;
    desc.image_depth = 0;
    desc.image_array_size = 0;
    desc.image_row_pitch = 0;
    desc.image_slice_pitch = 0;
    desc.num_mip_levels = 0;
    desc.num_samples = 0;
    desc.buffer = NULL;
    
    // 分配主机数据（可选）
    size_t pixel_size = 4;  // RGBA = 4 bytes
    size_t data_size = desc.image_width * desc.image_height * pixel_size;
    unsigned char *host_data = (unsigned char*)malloc(data_size);
    
    // 初始化测试图像（红色渐变）
    for (size_t y = 0; y < desc.image_height; y++)
    {
        for (size_t x = 0; x < desc.image_width; x++)
        {
            size_t idx = (y * desc.image_width + x) * pixel_size;
            host_data[idx + 0] = (unsigned char)(x * 255 / desc.image_width);  // R
            host_data[idx + 1] = (unsigned char)(y * 255 / desc.image_height); // G
            host_data[idx + 2] = 0;                                             // B
            host_data[idx + 3] = 255;                                           // A
        }
    }
    
    // 创建图像并拷贝数据
    cl_mem image = clCreateImage(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
                                &format, &desc, host_data, &error);
    test_error(error, "Unable to create 2D image");
    
    log_info("2D image created: %zux%zu\n", desc.image_width, desc.image_height);
    
    free(host_data);
    clReleaseMemObject(image);
    return 0;
}

2D 图像内核操作

// 2D 图像读取（使用采样器）
__kernel void test_read_image_2D(__read_only image2d_t src,
                                 sampler_t sampler,
                                 __global float4 *output)
{
    int x = get_global_id(0);
    int y = get_global_id(1);
    int width = get_image_width(src);
    
    // 整数坐标读取
    int2 coord = (int2)(x, y);
    float4 pixel = read_imagef(src, sampler, coord);
    
    output[y * width + x] = pixel;
}

// 2D 图像写入
__kernel void test_write_image_2D(__write_only image2d_t dst,
                                  __global float4 *input)
{
    int x = get_global_id(0);
    int y = get_global_id(1);
    int width = get_image_width(dst);
    
    int2 coord = (int2)(x, y);
    float4 pixel = input[y * width + x];
    
    write_imagef(dst, coord, pixel);
}

// 图像处理示例：边缘检测
__kernel void edge_detection(__read_only image2d_t src,
                             __write_only image2d_t dst,
                             sampler_t sampler)
{
    int x = get_global_id(0);
    int y = get_global_id(1);
    
    // Sobel 算子
    float4 pixel00 = read_imagef(src, sampler, (int2)(x-1, y-1));
    float4 pixel01 = read_imagef(src, sampler, (int2)(x, y-1));
    float4 pixel02 = read_imagef(src, sampler, (int2)(x+1, y-1));
    
    float4 pixel10 = read_imagef(src, sampler, (int2)(x-1, y));
    float4 pixel12 = read_imagef(src, sampler, (int2)(x+1, y));
    
    float4 pixel20 = read_imagef(src, sampler, (int2)(x-1, y+1));
    float4 pixel21 = read_imagef(src, sampler, (int2)(x, y+1));
    float4 pixel22 = read_imagef(src, sampler, (int2)(x+1, y+1));
    
    // 水平梯度
    float4 gx = -pixel00 - 2.0f*pixel10 - pixel20
                + pixel02 + 2.0f*pixel12 + pixel22;
    
    // 垂直梯度
    float4 gy = -pixel00 - 2.0f*pixel01 - pixel02
                + pixel20 + 2.0f*pixel21 + pixel22;
    
    // 梯度幅值
    float4 magnitude = sqrt(gx*gx + gy*gy);
    
    write_imagef(dst, (int2)(x, y), magnitude);
}

6.2.3 3D 图像测试

三维图像用于体数据、医学成像等应用。

创建 3D 图像

int test_3D_image_creation(cl_device_id device, cl_context context)
{
    cl_int error;
    
    // 检查 3D 图像支持
    cl_bool image3d_support;
    error = clGetDeviceInfo(device, CL_DEVICE_IMAGE_SUPPORT,
                           sizeof(image3d_support), &image3d_support, NULL);
    test_error(error, "Unable to query image support");
    
    if (!image3d_support)
    {
        log_info("Device does not support images, skipping test\n");
        return 0;
    }
    
    // 设置 3D 图像格式和描述符
    cl_image_format format;
    format.image_channel_order = CL_RGBA;
    format.image_channel_data_type = CL_FLOAT;
    
    cl_image_desc desc;
    memset(&desc, 0, sizeof(desc));
    desc.image_type = CL_MEM_OBJECT_IMAGE3D;
    desc.image_width = 64;
    desc.image_height = 64;
    desc.image_depth = 64;
    desc.image_array_size = 0;
    desc.image_row_pitch = 0;
    desc.image_slice_pitch = 0;
    desc.num_mip_levels = 0;
    desc.num_samples = 0;
    desc.buffer = NULL;
    
    // 创建 3D 图像
    cl_mem image = clCreateImage(context, CL_MEM_READ_WRITE,
                                &format, &desc, NULL, &error);
    test_error(error, "Unable to create 3D image");
    
    // 查询 3D 图像属性
    size_t width, height, depth;
    clGetImageInfo(image, CL_IMAGE_WIDTH, sizeof(width), &width, NULL);
    clGetImageInfo(image, CL_IMAGE_HEIGHT, sizeof(height), &height, NULL);
    clGetImageInfo(image, CL_IMAGE_DEPTH, sizeof(depth), &depth, NULL);
    
    log_info("3D image created: %zux%zux%zu\n", width, height, depth);
    
    clReleaseMemObject(image);
    return 0;
}

3D 图像内核

// 读取 3D 图像
__kernel void test_read_image_3D(__read_only image3d_t src,
                                 sampler_t sampler,
                                 __global float4 *output)
{
    int x = get_global_id(0);
    int y = get_global_id(1);
    int z = get_global_id(2);
    
    int width = get_image_width(src);
    int height = get_image_height(src);
    
    // 3D 坐标
    int4 coord = (int4)(x, y, z, 0);
    float4 pixel = read_imagef(src, sampler, coord);
    
    int idx = z * height * width + y * width + x;
    output[idx] = pixel;
}

// 写入 3D 图像
__kernel void test_write_image_3D(__write_only image3d_t dst,
                                  __global float4 *input)
{
    int x = get_global_id(0);
    int y = get_global_id(1);
    int z = get_global_id(2);
    
    int width = get_image_width(dst);
    int height = get_image_height(dst);
    
    int4 coord = (int4)(x, y, z, 0);
    int idx = z * height * width + y * width + x;
    
    write_imagef(dst, coord, input[idx]);
}

// 3D 高斯模糊
__kernel void gaussian_blur_3D(__read_only image3d_t src,
                               __write_only image3d_t dst,
                               sampler_t sampler)
{
    int x = get_global_id(0);
    int y = get_global_id(1);
    int z = get_global_id(2);
    
    float4 sum = (float4)(0.0f);
    float weight_sum = 0.0f;
    
    // 3x3x3 高斯核
    for (int dz = -1; dz <= 1; dz++)
    {
        for (int dy = -1; dy <= 1; dy++)
        {
            for (int dx = -1; dx <= 1; dx++)
            {
                int4 coord = (int4)(x + dx, y + dy, z + dz, 0);
                float weight = exp(-(dx*dx + dy*dy + dz*dz) / 2.0f);
                
                float4 pixel = read_imagef(src, sampler, coord);
                sum += pixel * weight;
                weight_sum += weight;
            }
        }
    }
    
    write_imagef(dst, (int4)(x, y, z, 0), sum / weight_sum);
}

---

6.3 图像数组测试

6.3.1 1D 图像数组

1D 图像数组是多个 1D 图像的集合。

int test_1D_image_array(cl_context context)
{
    cl_int error;
    cl_image_format format;
    cl_image_desc desc;
    
    format.image_channel_order = CL_RGBA;
    format.image_channel_data_type = CL_UNORM_INT8;
    
    memset(&desc, 0, sizeof(desc));
    desc.image_type = CL_MEM_OBJECT_IMAGE1D_ARRAY;
    desc.image_width = 256;        // 每个 1D 图像的宽度
    desc.image_height = 0;         // 1D 数组高度为 0
    desc.image_depth = 0;
    desc.image_array_size = 16;    // 数组中 1D 图像的数量
    desc.image_row_pitch = 0;
    desc.image_slice_pitch = 0;
    desc.num_mip_levels = 0;
    desc.num_samples = 0;
    desc.buffer = NULL;
    
    cl_mem image_array = clCreateImage(context, CL_MEM_READ_WRITE,
                                      &format, &desc, NULL, &error);
    test_error(error, "Unable to create 1D image array");
    
    log_info("1D image array created: width=%zu, array_size=%zu\n",
             desc.image_width, desc.image_array_size);
    
    clReleaseMemObject(image_array);
    return 0;
}

1D 图像数组内核

__kernel void test_read_image_1D_array(__read_only image1d_array_t src,
                                       sampler_t sampler,
                                       __global float4 *output)
{
    int x = get_global_id(0);
    int layer = get_global_id(1);
    
    // 1D 数组坐标：(x, layer)
    int2 coord = (int2)(x, layer);
    float4 pixel = read_imagef(src, sampler, coord);
    
    int width = get_image_width(src);
    output[layer * width + x] = pixel;
}

6.3.2 2D 图像数组

2D 图像数组用于纹理图集、视频帧序列等。

int test_2D_image_array(cl_context context)
{
    cl_int error;
    cl_image_format format;
    cl_image_desc desc;
    
    format.image_channel_order = CL_RGBA;
    format.image_channel_data_type = CL_UNORM_INT8;
    
    memset(&desc, 0, sizeof(desc));
    desc.image_type = CL_MEM_OBJECT_IMAGE2D_ARRAY;
    desc.image_width = 128;
    desc.image_height = 128;
    desc.image_depth = 0;
    desc.image_array_size = 8;     // 8 个 2D 图像
    desc.image_row_pitch = 0;
    desc.image_slice_pitch = 0;
    desc.num_mip_levels = 0;
    desc.num_samples = 0;
    desc.buffer = NULL;
    
    cl_mem image_array = clCreateImage(context, CL_MEM_READ_WRITE,
                                      &format, &desc, NULL, &error);
    test_error(error, "Unable to create 2D image array");
    
    log_info("2D image array created: %zux%zu, %zu layers\n",
             desc.image_width, desc.image_height, desc.image_array_size);
    
    clReleaseMemObject(image_array);
    return 0;
}

2D 图像数组内核

__kernel void test_read_image_2D_array(__read_only image2d_array_t src,
                                       sampler_t sampler,
                                       __global float4 *output)
{
    int x = get_global_id(0);
    int y = get_global_id(1);
    int layer = get_global_id(2);
    
    // 2D 数组坐标：(x, y, layer)
    int4 coord = (int4)(x, y, layer, 0);
    float4 pixel = read_imagef(src, sampler, coord);
    
    int width = get_image_width(src);
    int height = get_image_height(src);
    int idx = layer * height * width + y * width + x;
    
    output[idx] = pixel;
}

// 跨层处理：从多个层混合数据
__kernel void blend_layers(__read_only image2d_array_t src,
                          __write_only image2d_t dst,
                          sampler_t sampler,
                          int num_layers)
{
    int x = get_global_id(0);
    int y = get_global_id(1);
    
    float4 sum = (float4)(0.0f);
    
    // 平均所有层的像素
    for (int layer = 0; layer < num_layers; layer++)
    {
        int4 coord = (int4)(x, y, layer, 0);
        sum += read_imagef(src, sampler, coord);
    }
    
    float4 result = sum / (float)num_layers;
    write_imagef(dst, (int2)(x, y), result);
}

---

6.4 图像读写与采样器测试

6.4.1 采样器对象

采样器控制图像读取的行为。

cl_sampler create_test_sampler(cl_context context)
{
    cl_int error;
    cl_sampler_properties properties[] = {
        CL_SAMPLER_NORMALIZED_COORDS, CL_FALSE,          // 非归一化坐标
        CL_SAMPLER_ADDRESSING_MODE, CL_ADDRESS_CLAMP,    // 边界处理：夹紧
        CL_SAMPLER_FILTER_MODE, CL_FILTER_NEAREST,       // 过滤模式：最近邻
        0
    };
    
    cl_sampler sampler = clCreateSamplerWithProperties(context, properties, &error);
    test_error(error, "Unable to create sampler");
    
    return sampler;
}

采样器属性：

属性	值	说明
CL_SAMPLER_NORMALIZED_COORDS	CL_TRUE/CL_FALSE	是否使用归一化坐标 [0,1]
CL_SAMPLER_ADDRESSING_MODE	CL_ADDRESS_CLAMP	边界夹紧到边缘
	CL_ADDRESS_CLAMP_TO_EDGE	夹紧到边缘像素
	CL_ADDRESS_REPEAT	重复平铺
	CL_ADDRESS_MIRRORED_REPEAT	镜像重复
CL_SAMPLER_FILTER_MODE	CL_FILTER_NEAREST	最近邻插值
	CL_FILTER_LINEAR	线性插值

6.4.2 归一化坐标 vs 整数坐标

// 归一化坐标读取（0.0 到 1.0）
__kernel void read_normalized(__read_only image2d_t src,
                              sampler_t sampler,
                              __global float4 *output)
{
    int x = get_global_id(0);
    int y = get_global_id(1);
    int width = get_image_width(src);
    int height = get_image_height(src);
    
    // 归一化坐标
    float2 norm_coord = (float2)((float)x / width, (float)y / height);
    float4 pixel = read_imagef(src, sampler, norm_coord);
    
    output[y * width + x] = pixel;
}

// 整数坐标读取
__kernel void read_integer(__read_only image2d_t src,
                          sampler_t sampler,
                          __global float4 *output)
{
    int x = get_global_id(0);
    int y = get_global_id(1);
    int width = get_image_width(src);
    
    // 整数坐标
    int2 coord = (int2)(x, y);
    float4 pixel = read_imagef(src, sampler, coord);
    
    output[y * width + x] = pixel;
}

6.4.3 图像过滤模式

// 最近邻过滤
__constant sampler_t nearest_sampler = CLK_NORMALIZED_COORDS_FALSE |
                                       CLK_ADDRESS_CLAMP |
                                       CLK_FILTER_NEAREST;

// 线性过滤
__constant sampler_t linear_sampler = CLK_NORMALIZED_COORDS_FALSE |
                                      CLK_ADDRESS_CLAMP |
                                      CLK_FILTER_LINEAR;

__kernel void compare_filtering(__read_only image2d_t src,
                               __write_only image2d_t dst_nearest,
                               __write_only image2d_t dst_linear)
{
    int x = get_global_id(0);
    int y = get_global_id(1);
    
    // 使用浮点坐标（介于像素中心之间）
    float2 coord = (float2)((float)x + 0.5f, (float)y + 0.5f);
    
    // 最近邻插值
    float4 pixel_nearest = read_imagef(src, nearest_sampler, coord);
    write_imagef(dst_nearest, (int2)(x, y), pixel_nearest);
    
    // 线性插值
    float4 pixel_linear = read_imagef(src, linear_sampler, coord);
    write_imagef(dst_linear, (int2)(x, y), pixel_linear);
}

---

6.5 图像拷贝与填充

6.5.1 图像拷贝

使用 clEnqueueCopyImage 在设备端拷贝图像。

int test_copy_image_2D(cl_command_queue queue, cl_context context)
{
    cl_int error;
    cl_image_format format;
    cl_image_desc desc;
    
    format.image_channel_order = CL_RGBA;
    format.image_channel_data_type = CL_UNORM_INT8;
    
    memset(&desc, 0, sizeof(desc));
    desc.image_type = CL_MEM_OBJECT_IMAGE2D;
    desc.image_width = 256;
    desc.image_height = 256;
    
    // 创建源图像和目标图像
    cl_mem src_image = clCreateImage(context, CL_MEM_READ_ONLY,
                                     &format, &desc, NULL, &error);
    test_error(error, "Unable to create source image");
    
    cl_mem dst_image = clCreateImage(context, CL_MEM_WRITE_ONLY,
                                     &format, &desc, NULL, &error);
    test_error(error, "Unable to create dest image");
    
    // 生成测试数据并写入源图像
    size_t pixel_size = 4;  // RGBA
    size_t data_size = desc.image_width * desc.image_height * pixel_size;
    unsigned char *test_data = (unsigned char*)malloc(data_size);
    
    for (size_t i = 0; i < data_size; i++)
        test_data[i] = (unsigned char)(i % 256);
    
    size_t origin[3] = {0, 0, 0};
    size_t region[3] = {desc.image_width, desc.image_height, 1};
    
    error = clEnqueueWriteImage(queue, src_image, CL_TRUE,
                                origin, region, 0, 0,
                                test_data, 0, NULL, NULL);
    test_error(error, "Unable to write source image");
    
    // 拷贝图像
    error = clEnqueueCopyImage(queue, src_image, dst_image,
                              origin, origin, region,
                              0, NULL, NULL);
    test_error(error, "Unable to copy image");
    
    // 读取并验证
    unsigned char *result = (unsigned char*)malloc(data_size);
    error = clEnqueueReadImage(queue, dst_image, CL_TRUE,
                               origin, region, 0, 0,
                               result, 0, NULL, NULL);
    test_error(error, "Unable to read dest image");
    
    // 验证数据
    if (memcmp(test_data, result, data_size) != 0)
    {
        log_error("ERROR: Image copy data mismatch\n");
        free(test_data);
        free(result);
        return -1;
    }
    
    log_info("Image copy test passed\n");
    
    free(test_data);
    free(result);
    clReleaseMemObject(src_image);
    clReleaseMemObject(dst_image);
    return 0;
}

6.5.2 部分图像拷贝

拷贝图像的矩形区域：

int test_copy_image_region(cl_command_queue queue, 
                          cl_mem src_image, cl_mem dst_image)
{
    cl_int error;
    
    // 源图像区域：从 (32, 32) 开始，大小 64x64
    size_t src_origin[3] = {32, 32, 0};
    
    // 目标图像位置：放在 (64, 64)
    size_t dst_origin[3] = {64, 64, 0};
    
    // 拷贝区域大小
    size_t region[3] = {64, 64, 1};
    
    error = clEnqueueCopyImage(queue, src_image, dst_image,
                              src_origin, dst_origin, region,
                              0, NULL, NULL);
    test_error(error, "Unable to copy image region");
    
    return 0;
}

6.5.3 图像填充

使用 clEnqueueFillImage 填充图像（OpenCL 1.2+）。

int test_fill_image(cl_command_queue queue, cl_context context)
{
    cl_int error;
    cl_image_format format;
    cl_image_desc desc;
    
    format.image_channel_order = CL_RGBA;
    format.image_channel_data_type = CL_FLOAT;
    
    memset(&desc, 0, sizeof(desc));
    desc.image_type = CL_MEM_OBJECT_IMAGE2D;
    desc.image_width = 128;
    desc.image_height = 128;
    
    cl_mem image = clCreateImage(context, CL_MEM_READ_WRITE,
                                &format, &desc, NULL, &error);
    test_error(error, "Unable to create image");
    
    // 填充颜色：红色 (1.0, 0.0, 0.0, 1.0)
    cl_float4 fill_color = {{1.0f, 0.0f, 0.0f, 1.0f}};
    
    size_t origin[3] = {0, 0, 0};
    size_t region[3] = {desc.image_width, desc.image_height, 1};
    
    error = clEnqueueFillImage(queue, image, &fill_color,
                              origin, region, 0, NULL, NULL);
    test_error(error, "Unable to fill image");
    
    // 读取验证
    size_t pixel_count = desc.image_width * desc.image_height;
    cl_float4 *result = (cl_float4*)malloc(sizeof(cl_float4) * pixel_count);
    
    error = clEnqueueReadImage(queue, image, CL_TRUE,
                               origin, region, 0, 0,
                               result, 0, NULL, NULL);
    test_error(error, "Unable to read image");
    
    // 验证所有像素是否为红色
    for (size_t i = 0; i < pixel_count; i++)
    {
        if (result[i].s[0] != 1.0f || result[i].s[1] != 0.0f ||
            result[i].s[2] != 0.0f || result[i].s[3] != 1.0f)
        {
            log_error("ERROR: Fill color mismatch at pixel %zu\n", i);
            free(result);
            clReleaseMemObject(image);
            return -1;
        }
    }
    
    log_info("Image fill test passed\n");
    
    free(result);
    clReleaseMemObject(image);
    return 0;
}

填充不同数据类型：

// 填充整数图像
cl_int4 fill_int = {{255, 128, 64, 32}};
clEnqueueFillImage(queue, int_image, &fill_int, origin, region, 0, NULL, NULL);

// 填充无符号整数图像
cl_uint4 fill_uint = {{1000, 2000, 3000, 4000}};
clEnqueueFillImage(queue, uint_image, &fill_uint, origin, region, 0, NULL, NULL);

---

本章小结

第6章详细介绍了 Images 测试套件的核心内容：

• 图像格式支持：通道顺序、数据类型、必需格式验证
• 1D/2D/3D 图像：创建、读写、图像处理算法
• 图像数组：1D/2D 图像数组、跨层操作
• 采样器：归一化坐标、过滤模式、边界处理
• 拷贝填充：完整/部分图像拷贝、模式填充

这些测试确保 OpenCL 图像对象在各种场景下正确工作，对图形和图像处理应用至关重要。下一章将介绍 Allocations 内存分配测试。