OpenCL AES Parallelization

OpenCL AES Parallelization

本文关键字：Parallelization AES OpenCL      更新时间：2023-10-16

我正在尝试编写一些为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)来了解每个线程操作的字节。与barrier(CLK_LOCAL_MEM_FENCE)同步，以防工作组在处理器上运行时可能无法同步。键可能会进入常量内存，因为它可以被缓存。如果只是为了避免从全局内存中重新加载前一个块密文，块链可能是具有循环的适当级别。此外，使用async_work_group_copy()异步存储完整的密文可能会有所帮助。你可以通过使用向量让线程做更多的工作，但这可能没有帮助，因为有像shiftRows这样的步骤。

基本上，如果16个线程中的任何一个线程(可能因架构而异)得到任何不同的控制流，你的GPU就会停滞。如果没有足够的这样的组来填充管道和多处理器，那么您的GPU将处于闲置状态。除非你非常仔细地优化了内存访问，否则它不会接近CPU的速度，即使在那之后，你也需要一次处理几十个数据包，以避免给GPU提供太小的工作组。问题是，尽管GPU可以运行数千个线程，但它的控制结构在任何时候都只能处理几个工作组。

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

从你的描述中不完全清楚，但我认为有一些概念上的混淆。

不要遍历每个包并启动一个新内核。你不需要告诉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知道数组的索引。您将需要做类似的事情。

相关文章：