灰度内核

让我们构建一个内核来生成灰度图像。我们将使用为每个组件使用 uint 定义的图像数据和 RGBA 订单。

__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
                                CLK_ADDRESS_CLAMP_TO_EDGE |
                                CLK_FILTER_LINEAR;

__kernel void Grayscale(__read_only image2d_t input, __write_only image2d_t output) {
    int2 gid = (int2)(get_global_id(0), get_global_id(1));
    int2 size = get_image_dim(input);

    if(all(gid < size)){
        uint4 pixel = read_imageui(input, sampler, gid);
        float4 color = convert_float4(pixel) / 255;
        color.xyz = 0.2126*color.x + 0.7152*color.y + 0.0722*color.z;
        pixel = convert_uint4_rte(color * 255);
        write_imageui(output, gid, pixel);
    }
}

现在让我们一步一步地浏览该代码。第一行在 sampler_t 类型的__constant 内存区域中创建一个变量。此采样器用于进一步指定对图像数据的访问。有关完整文档,请参阅 Khronos 文档。

在我们调用内核之前,我们将输入分配为 read_only,输出分配为 write_only,因此我们在此处添加这些修饰符。

image2d 和 image3d 总是在全局内存中分配,因此我们可以省略__global 修饰符。

然后我们得到我们的线程 id,它定义了我们要转换为灰度的像素。我们还查询大小以确保我们的线程不访问未分配的内存。如果你忘记了,这将定义你的内核崩溃。

在我们确定我们是合法线程之后,我们从输入图像中读取了像素。然后我们将它转​​换为 float 以避免丢失小数位,进行一些计算,将其转换回来并将其写入输出。