c – OpenCL AES并行化

印度阿三17 2019-09-02

展开全文

我正在尝试编写一些为SSL服务器进行AES解密的代码.为了加快速度,我试图将多个数据包组合在一起,一次在GPU上解密.

如果我只是遍历每个数据包,我将每个内核提交给gpu,然后是一个使用内核事件进行等待的读取.然后我收集所有读取的事件并同时等待它们,但它似乎只是一次运行一个块然后执行下一个块.这不是我所期望的.我希望如果我将所有内核排队,那么我希望驱动程序能够尝试尽可能多地并行工作.

我错过了什么吗？我是否必须将全局工作量指定为所有数据包块的大小,并将内核本地大小指定为每个数据包块的大小？

这是我的OpenCL内核的代码.

__kernel void decryptCBC( __global const uchar *rkey, const uint rounds, 
    __global const uchar* prev, __global const uchar *data, 
    __global uchar *result, const uint blocks ) {

    const size_t id = get_global_id( 0 );
    if( id > blocks ) return;

    const size_t startPos = BlockSize * id;

    // Create Block
    uchar block[BlockSize];
    for( uint i = 0; i < BlockSize; i  ) block[i] = data[startPos i];

    // Calculate Result
    AddRoundKey( rkey, block, rounds );

    for( uint j = 1; j < rounds;   j ){
        const uint round = rounds - j;
        InverseShiftRows( block );
        InverseSubBytes( block );
        AddRoundKey( rkey, block, round );
        InverseMixColumns( block );
    }

    InverseSubBytes( block );
    InverseShiftRows( block );
    AddRoundKey( rkey, block, 0 );

    // Store Result
    for( uint i = 0; i < BlockSize; i   ) {
        result[startPos i] = block[i] ^ prev[startPos i];
    }
}

有了这个内核,我可以在一个数据包中击败一个包含125块数据的8核CPU.为了加速多个数据包,我试图将所有数据元素组合在一起.这涉及将输入数据组合成单个向量,然后复杂性来自每个内核需要知道在密钥内访问的位置,从而导致包含轮数和轮次偏移的两个额外数组.事实证明,这比每个数据包单独执行内核要慢.

解决方法:

将您的内核视为执行CBC工作的函数.正如您所发现的,它的链接性质意味着CBC任务本身是基本序列化的.此外,GPU更喜欢运行具有相同工作负载的16个线程.这基本上是多处理器核心中单个任务的大小,你往往有几十个;但是管理系统总体上只能为它们提供一些这样的任务,而内存系统很少能跟上它们.此外,循环是内核最糟糕的用途之一,因为GPU不是为了控制流量而设计的.

因此,看看AES,它运行在16字节块上,但仅在字节操作中运行.这将是您的第一个维度 – 每个块应该由16个线程处理(可能是opencl用语中的本地工作大小).确保将块传输到本地内存,其中所有线程可以锁步执行,以非常低的延迟进行随机访问.在AES块操作中展开所有内容,使用get_local_id(0)来了解每个线程操作的字节.如果工作组在可能用完锁步的处理器上运行,则与屏障(CLK_LOCAL_MEM_FENCE)同步.密钥可能会进入常量内存,因为这可以缓存.块链接可能是具有循环的适当级别,如果仅为了避免从全局存储器重新加载先前的块密文.使用async_work_group_copy()异步存储已完成的密文可能会有所帮助.你可以通过使用向量使线程做更多工作,但由于像shiftRows这样的步骤,这可能无济于事.

基本上,如果一组16个线程中的任何线程(可能随架构而变化)获得任何不同的控制流,那么GPU就会停滞不前.如果没有足够的这样的组来填充管道和多处理器,那么你的GPU就处于空闲状态.在您非常仔细地优化内存访问之前,它不会接近CPU速度,甚至在此之后,您需要同时处理几十个数据包以避免给GPU太小的工作组.问题是虽然GPU可以运行数千个线程,但它的控制结构只能随时处理几个工作组.

还有一件事需要注意;当您在工作组中使用障碍时,工作组中的每个线程都必须执行相同的屏障调用.这意味着即使你有额外的线程空闲运行(例如,那些解密组合工作组中的较短数据包),即使它们没有内存访问,它们也必须继续循环.

来源：https://www./content-4-436001.html