用于处理 3 字节块的 OpenCL 内核

Question

我有一个 OpenCL 内核，它改变了 24 位宽块内的位顺序。

我第一次尝试实现这一点是创建 buffersize/3 线程，但不知何故，我的 OpenCl 内核与 CPU（2 GHz Intel Core i7）上的相同算法相比运行速度较慢。

typedef unsigned char uint8_t;

__kernel void decode(
  __global uint8_t* in,
  __global uint8_t* out,
  const unsigned int count) {
    int i = get_global_id(0)*3;

    if(i<count){
        out[i]=in[i];
        out[i+1]=(in[i+1]&0b00001111)<<4|(in[i+2]&0b11110000)>>4;
        out[i+2]=(in[i+2]&0b00001111)<<4|(in[i+1]&0b11110000)>>4;
    }
}

这个内核是这样调用的：

count = buffersize/3; // buffersize is approx. 6 to 8 MB
error = clEnqueueNDRangeKernel(
        commands, 
        koDecode, 
        1, 
        NULL, 
        &count, 
        NULL, 
        0, 
        NULL, 
        NULL);

有更好的方法吗？

Answer 1

我不确定我是否理解您要在您的算法中实现的目标，但您应该使用 CL 向量。以下行内的内容：

__kernel void decode(
            __global uchar3* in,
            __global uchar3* out,
            const unsigned int count)
{
   int i = get_global_id(0);
   out[i] = out[i].zyx;
}

Answer 2

这可能会更快，因为 out "may" 指向 in 大多数编译器不缓存全局变量：

__kernel void decode(
  __global uint8_t* in,
  __global uint8_t* out,
  const unsigned int count) {
    int i = get_global_id(0)*3;

    if(i<count){
        out[i]=in[i];
        uint8_t t[2];
        t[0] = in[i+1];
        t[1] = in[i+2];
        out[i+1]=(t[0]&0b00001111)<<4|(t[1]&0b11110000)>>4;
        out[i+2]=(t[1]&0b00001111)<<4|(t[0]&0b11110000)>>4;
    }
}

此外，如果您始终启动所需数量的工作项，您实际上并不需要 i<count 检查。

用于处理 3 字节块的 OpenCL 内核

OpenCL kernel for working with 3 Bytes blocks

opencl