OpenCL SHA1流量优化
希望在OpenCL使用方面更有经验的人能在这里帮助我！我正在做一个项目(帮助我学习更多的密码，并尝试我的GPGPU编程)，在那里我试图实现我自己的SHA-1算法。
最终，我的问题是最大限度地提高吞吐量。目前，我看到的东西大概是56.1 MH/秒，这与我看过的开源程序相比非常糟糕，比如John the Ripper和OCLHashcat，它们分别提供了1,000 MH和1,500 MH/秒(见鬼，我对其中的三分之一非常满意！)
所以，我在做什么
我在OpenCL内核和C++主机应用程序中编写了一个SHA-1实现来将数据加载到GPU (使用CL1.2 C++包装器)。我正在以线程方式在CPU上生成候选数据块，并使用对C++的CL enqueueWriteBuffer调用将这些数据加载到全局GPU内存中(使用uchars表示要散列的字节)：
errorCode = dispatchQueue->enqueueWriteBuffer(
inputBuffer,
CL_FALSE,//CL_TRUE,
0,
sizeof(cl_uchar) * inputBufferSize,
passwordBuffer,
NULL,
&dispatchDelegate);我以以下方式使用enqueueNDRangeKernel对数据进行排队(其中全局worksize是用户定义的变量，目前我已将其设置为GPU的最大平均全局工作大小为每次运行16777万次)：
errorCode = dispatchQueue->enqueueNDRangeKernel(
*kernel,
NullRange,
NDRange(globalWorkgroupSize, 1),
NullRange,
NULL,
NULL);这意味着(每次分派)我使用get_global_offset(0)加载了一个1D数组中的16777万项，并将内核中的索引加载到这个数组中。
我的核心签名：
__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，在一个一维输入向量上执行1670万次内核执行？我应该使用二维空间并细分成局部工作组范围吗？我试过玩这个，但看上去没那么快。
或者，也许我的算法本身就是缓慢的来源？我花了大量的时间来优化它，并使用预处理器指令手动地展开所有的循环阶段。
我读过关于内存集中在硬件上的文章。这是我的问题吗？
任何建议都要感谢！如果我错过了什么重要的事情，请告诉我，我会更新的。
(预先谢谢！)
更新： 16,777,216是设备最大报告的工作组大小；256**3。布尔值的全局数组是一个布尔值。在内核队列开始时，它被设置为false，然后分支语句将只发现冲突时将其设置为true --这会强制收敛吗？passwordLen是当前输入值的长度，目标哈希是要检查的int4编码哈希。
Stack Overflow用户
发布于 2014-04-22 12:06:56
您的“最大扁平的全局工作大小”应该乘以passwordLen。它是可以运行的内核数，而不是输入数组的最大长度。您很可能会向GPU发送更多的数据。
其他潜在问题：“在CPU上以线程方式生成候选数据块”，在内核迭代之前尝试这样做，以查看延迟是在生成数据块还是在处理内核；您的sha1算法是另一个明显的潜在问题。我不确定通过‘展开’循环来真正优化它的程度，通常更大的优化问题是'if‘语句(如果一个工作组中的单个内核实例测试为真，那么所有锁定的工作组实例都必须并行地跟随该分支)。
DarkZeros是正确的，您应该手动处理本地工作组大小，使其成为全局大小的最高公共倍数和可以同时在卡上运行的内核数。最简单的方法是将全局工作组大小集合到卡容量的下一个倍数，并在内核中使用外部if{}语句，该语句只为global_id运行内核，小于要运行的内核的实际数量。
戴夫。
https://stackoverflow.com/questions/22976617
复制相似问题
腾讯云开发者
