c++ - OpenCL SHA1 吞吐量优化

Question

希望在 OpenCL 使用方面更有经验的人可以在这里帮助我!我正在做一个项目(帮助我学习更多的加密知识并尝试 GPGPU 编程)，我正在尝试实现我自己的 SHA-1 算法。

最终我的问题是关于最大化我的吞吐率。目前我看到的是 56.1 MH/sec，与我看过的开源程序相比非常糟糕，例如 John the Ripper 和 OCLHashcat，后者分别给出 1,000 和 1,500 MH/sec(见鬼，我会很高兴得到其中的三分之一!)。

那么，我在做什么

我已经在 OpenCL 内核和 C++ 主机应用程序中编写了一个 SHA-1 实现，以将数据加载到 GPU(使用 CL 1.2 C++ 包装器)。我正在生成候选数据 block 以在 CPU 上以线程方式散列，并使用 CL C++ 调用 enqueueWriteBuffer(使用 uchars 表示要散列的字节)将此数据加载到全局 GPU 内存中:

errorCode = dispatchQueue->enqueueWriteBuffer(
        inputBuffer,
        CL_FALSE,//CL_TRUE,
        0,
        sizeof(cl_uchar) * inputBufferSize,
        passwordBuffer,
        NULL,
        &dispatchDelegate);

我按以下方式使用 enqueueNDRangeKernel 对数据进行排队(其中全局 worksize 是用户定义的变量，目前我已将其设置为我的GPU 的最大扁平化全局工作规模为每次运行 1677.7 万):

errorCode = dispatchQueue->enqueueNDRangeKernel(
        *kernel,
        NullRange,
        NDRange(globalWorkgroupSize, 1), 
        NullRange, 
        NULL,
        NULL);

这意味着(每次分派(dispatch))我将 1677 万个项目加载到一维数组中，并使用 get_global_offset(0) 从我的内核中索引到这个数组中。

我的内核签名:

    __kernel void sha1Crack(__global uchar* out, __global uchar* in, 
                            __constant int* passLen, __constant int* targetHash, 
                            __global bool* collisionFound)
    {
        //Kernel Instance Global GPU Mem IO Mapping:
        __private int id = get_global_id(0);
        __private int inputIndexStart = id * passwordLen;

        //Select Password input key space:
        #pragma unroll
        for (i = 0; i < passwordLen; i++)
        {
            inputMem[i] = in[inputIndexStart + i];
        }

        //SHA1 Code omitted for brevity...
    }

那么，考虑到这一切:我在加载数据的方式上是否做错了什么？ 即 1 次调用 enqueueNDrange 以在一维输入 vector 上执行 1670 万次内核？我应该使用二维空间并 segmentation 为 localworkgroup 范围吗？我试过玩这个，但似乎并不快。

或者，也许我的算法本身就是缓慢的根源？我花了很多时间对其进行优化，并使用预处理器指令手动展开所有循环阶段。

我读过有关硬件内存合并的内容。那会是我的问题吗？ :S

任何建议都非常感谢!如果我错过了任何重要的事情，请告诉我，我会更新。

提前致谢! ;)

---

更新:16,777,216 是设备报告的最大工作组大小； 256**3。 bool 值的全局数组是一个 bool 值。它在内核入队开始时设置为 false，然后如果仅发现冲突，则分支语句将其设置为 true - 这会强制收敛吗？ passwordLen 是当前输入值的长度，target hash 是要检查的 int[4] 编码哈希。

Answer 1

您的“最大扁平化全局工作大小”应乘以 passwordLen。它是您可以运行的内核数，而不是输入数组的最大长度。您很可能可以向 GPU 发送比这更多的数据。

其他潜在问题:“生成候选数据 block 以在 CPU 上以线程方式散列”，尝试在内核迭代之前执行此操作，以查看延迟是在数据 block 的生成中还是在内核的处理；您的 sha1 算法是另一个明显的潜在问题。我不确定你通过“展开”循环真正优化了多少，通常更大的优化问题是“if”语句(如果工作组中的单个内核实例测试为真，那么所有锁步工作组实例必须并行跟随该分支)。

DarkZeros 是正确的，您应该手动调整本地工作组大小，使其成为全局大小和卡上可以同时运行的内核数量的最大公倍数。最简单的方法是将全局工作组大小四舍五入到卡容量的下一个倍数，并在内核中使用外部 if{} 语句，只运行内核 global_id 小于您想要的实际内核数跑。

戴夫。