OpenCL AES 并行化

Question

我正在尝试编写一些代码来为 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。为了加速多个数据包，我尝试将所有数据元素组合在一起。这涉及将输入数据组合成一个向量，然后复杂性来自每个内核需要知道在密钥中访问的位置，这导致两个额外的数组包含轮数和轮数偏移量。结果证明这比为每个数据包单独执行内核还要慢。

Answer 1

将您的内核视为执行 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 可以运行数千个线程，但它的控制结构在任何时候都只能处理几个工作组。

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

Answer 2

从您的描述中并不完全清楚，但我认为存在一些概念上的混淆。

不要遍历每个数据包并启动一个新内核。您不需要告诉 OpenCL 启动一堆内核。相反，将尽可能多的数据包上传到 GPU，然后只运行一次内核。当您指定工作组大小时，这就是 GPU 尝试同时运行的内核数。

您需要对内核进行编程，使其在您上传的数据中的不同位置查找它们的数据包。例如，如果您要将两个数组添加到第三个数组中，您的内核将如下所示：

__kernel void vectorAdd(__global const int* a,
                        __global const int* b,
                        __global int* c) {
  int idx = get_global_id(0);
  c[idx] = a[idx] + b[idx];
}

重要的部分是每个内核通过其全局 id 知道数组的索引。你会想做类似的事情。