是否有 OpenCL std::queue 等价物

Question

我有一些代码与此类似：

std::queue<unsigned> q; 
q.push(array[0]); 

while (!q.empty()) 
{ 
    unsigned index = q.front(); 
    q.pop(); 

    if(index >= someval){ 
        index = index - someval;  
        q.push(array[index]); 
    } 
    else{ 
        //do something
    } 
} 
}

我想将其移至 OpenCL 内核，我将如何以最有效的方式复制队列的功能。

目前我用一个固定大小的数组来实现它，它监视队列中元素的头部和计数。请问有没有更优雅的解决方案

谢谢

Answer 1

您应该检查 SYCL：https://www.khronos.org/sycl 其中 C++ 模板函数可以包含主机和设备代码以构建使用 OpenCL 加速的复杂算法。如果你不能使用它，如果你甚至没有 opencl 2.0 的设备：

1.2 版没有任何等效项。但是在评论中，你说：

Not sure that I understand what you mean but this would need to be individual to all threads on the GPU

那么就不需要线程间通信了，为了缓存+内存+计算效率，使用循环缓冲区实现一个FIFO应该很简单，只是不要溢出它（最多64个元素）并且不要't underflow（弹出次数多于推送次数，但易于实现边界检查（性能损失））：

推送

   bool push(__private uint * stack, uint value)
    {
        // pushing from bot, so you can pop it from top later (FIFO)
        // circular buffer for top performance
        uint bufLen=64;
        // zeroth element is counter for newest added element
        // first element is oldest element


        // circular buffer 
        uint nextIndex=(stack[0]%bufLen+2); // +2 because of top-bot headers

        // if overflows, it overwrites oldest elements one by one
        stack[nextIndex]=value;


        // if overflows, it still increments 
        stack[0]++;

        // simple and fast
        return true;
    }

检查是否为空

        bool empty(__private uint * stack)
        {
            // tricky if you overflow both
            return (stack[0]==stack[1]);
        }

前值

        uint front(__private uint * stack)
        {
            uint bufLen=64;

            // oldest element value (top)
            uint ptr=stack[1]%bufLen+2; // circular adr + 2 header


            return stack[ptr];
        }

爆裂

        uint pop(__private uint * stack)
        {
               uint bufLen=64;
               uint ptr=stack[1]%bufLen+2;
               // pop from top (oldest)
               uint returnValue=stack[ptr];
               stack[ptr]=0;

               // this will be new top ctr for ptr
               stack[1]++;

               // if underflows, gets garbage, don't underflow

               return returnValue;
        }

用于基准测试的示例内核：

        __kernel void queue0(__global uint * heap)
        {
            int id=get_global_id(0);
            __private uint q[100];
            for(int i=0;i<256;i++)
                q[i]=0;

            for(int i=0;i<55;i++)    
                push(q,i);

            for(int i=0;i<40;i++)    
                pop(q);

            for(int i=0;i<20;i++)    
                push(q,i);

            for(int i=0;i<35;i++)    
                pop(q);

            for(int i=0;i<35;i++)    
            {
                push(q,i);
                pop(q);
            }
            push(q,'h');
            push(q,'e');
            push(q,'l');
            push(q,'l');
            push(q,'o');
            push(q,' ');
            push(q,'w');
            push(q,'o');
            push(q,'r');
            push(q,'l');
            push(q,'d');
            for(int i=0;i<256;i++)
                heap[id*256+i]=q[i];
        }

缓冲区的输出（显示线程 id = 0 的计算结果）

121 110 0      0 0 0    0 0 0    0 0 0 // 121 pushes total, 110 pops total

0 0 0    0 0 0    0 0 0    0 0 0

0 0 0    0 0 0    0 0 0    0 0 0

0 0 0    0 0 0    0 0 0    0 0 0

104 101 108      108 111 32     119 111 114    108 100 0 
// hello world

在 6.35 毫秒内超过 200k 的 pushes+pops（运行 1024 个线程的内核，每个线程处理 256 个元素，但仅使用 64+2 个元素作为循环缓冲区）对于 1 通道 1600MHz ddr3 RAM 和 Intel具有 12 个计算单元的 HD Graphics 400（总共 96 个内核 @600 MHz）。

如果您使用 64 x 4 元素循环缓冲区构造一个 64 元素循环缓冲区，您也可以在堆栈顶部和底部之间添加元素！

是否有 OpenCL std::queue 等价物

is there an OpenCL std::queue equivalent

c++

opencl