c++ - 是否有 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个元素)并且不要下溢(弹出次数多于推送次数，但易于实现边界检查(具有性能损失)):

推

   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];
        }

buffer的输出(显示thread 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

超过 200k pushes+pops 在 6.35 毫秒内(运行内核 1024 个线程，每个线程处理 256 个元素，但仅使用 64+2 个元素作为循环缓冲区)对于 1 channel 1600MHz ddr3 RAM 和 Intel HD Graphics 400 with 12 compute单位(共 96 个核心 @600 MHz)。

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