如何解决OpenCL 强力 TEA 块 32 位,密钥 64 位
我决定自己学习 OpenCL 并为 TEA 算法编写一个暴力密码,我是否正确理解 OpenCL?你能在速度方面改进一些东西吗?我犯了什么错误?
我循环准备前5个字节,剩下的3个字节由内核整理,255个线程,每个65535
在主程序中:
for (int x5 = KEY[0]; x5 >= 0; x5--) {
KEY[0]=x5;
for (int x4 = KEY[1]; x4 >= 0; x4--) {
KEY[1]=x4;
for (int x3 = KEY[2]; x3 >= 0; x3--) {
KEY[2]=x3;
for (int x2 = KEY[3]; x2 >= 0; x2--) {
KEY[3]=x2;
for (int x = KEY[4]; x >= 0; x--) {
KEY[4]=x;
ret = clEnqueueWriteBuffer(command_queue,key_mem_obj,CL_TRUE,8 * sizeof(int),KEY,NULL,NULL);
ret = clEnqueueWriteBuffer(command_queue,cadr_mem_obj,1 * sizeof(int),CADR,NULL);
ret = clSetKernelArg(kernel,sizeof(cl_mem),(void *)&key_mem_obj);
ret = clSetKernelArg(kernel,2,(void *)&cadr_mem_obj);
NDRange = 0x0100;
ret = clEnqueueNDRangeKernel(command_queue,kernel,1,&NDRange,NULL);
if (ret != CL_SUCCESS) {
break;
}
ret = clEnqueueReadBuffer(command_queue,NULL);
if (CADR[0]>0) {
uint16_t k=CADR[0];
ret = clEnqueueReadBuffer(command_queue,retc_mem_obj,524280 * sizeof(int),RETC,NULL);
for ((i = 0); i < k; i++) {
Form1->Memo1->Lines->BeginUpdate();
Form1->Memo1->Lines->Add(IntToHex(RETC[i*8],2)+IntToHex(RETC[i*8+1],2)+
IntToHex(RETC[i*8+2],2)+IntToHex(RETC[i*8+3],2)+IntToHex(RETC[i*8+4],2)+
IntToHex(RETC[i*8+5],2)+IntToHex(RETC[i*8+6],2)+IntToHex(RETC[i*8+7],2));
Form1->Memo1->Lines->EndUpdate();
Form1->Label6->Caption=IntToStr(Form1->Memo1->Lines->Count-1);
}
CADR[0]=0;
}
KEY2[0]=KEY[0];
KEY2[1]=KEY[1];
KEY2[2]=KEY[2];
KEY2[3]=KEY[3];
KEY2[4]=KEY[4];
KEY2[5]=KEY[5];
KEY2[6]=KEY[6];
KEY2[7]=KEY[7];
if(Terminated){
break;
}
}
KEY[4]=0xFF;
}
KEY[3]=0xFF;
}
KEY[2]=0xFF;
}
KEY[1]=0xFF;
}
KEY[0]=0xFF;`
内核:
#pragma OPENCL EXTENSION cl_khr_byte_addressable_store : enable
#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable
#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable
#pragma OPENCL EXTENSION cl_khr_global_int32_extended_atomics : enable
#pragma OPENCL EXTENSION cl_khr_local_int32_extended_atomics : enable
__kernel void brute(__global const int *KEY,__global const int *DAT,__global int
*CADR,__global int *RETC)
{
int i = get_global_id(0);
ushort Data[2];
ushort Key[4];
Key[0]=(KEY[0]<<8)+KEY[1];
Key[1]=(KEY[2]<<8)+KEY[3];
// Key[2]=(KEY[4]<<8)+KEY[5];
Key[3]=(KEY[6]<<8)+KEY[7];
Key[2] = (KEY[4]<<8) + i;
for (int j=0xFFFF; j>=0; j--){
Key[3]=j;
Data[0]=(DAT[0]<<8)+DAT[1];
Data[1]=(DAT[2]<<8)+DAT[3];
ushort delta = 0x9e37;
ushort sum = (delta<<5);
for (uint n = 0;n < 32; ++n){
Data[1]-=(((Data[0])+Key[2])^(Data[0]+sum)^((Data[0]>>5)+Key[3]));
Data[0]-=(((Data[1]<<4)+Key[0])^(Data[1]+sum)^(Data[1]+Key[1]));
sum -= delta;
}
if ((Data[0]==0x0000) && (Data[1]==0x0000)){
int a=CADR[0];
atomic_inc(CADR);
RETC[a*8]=(Key[0] >> 8);
RETC[a*8+1]=(Key[0] & 0xFF);
RETC[a*8+2]=(Key[1] >> 8);
RETC[a*8+3]=(Key[1] & 0xFF);
RETC[a*8+4]=(Key[2] >> 8);
RETC[a*8+5]=(Key[2] & 0xFF);
RETC[a*8+6]=(Key[3] >> 8);
RETC[a*8+7]=(Key[3] & 0xFF);
}
}
}
解决方法
如果你只启动 256 个线程,每个线程做 65536 次迭代,你的 GPU 不会饱和,性能会很差。 GPU 有数千个“核心”,如果您启动 256 个线程,您将只使用其中的 256 个,而其余线程保持空闲。
GPU 并行化的想法是将工作拆分为尽可能多的独立问题。在您的情况下,这意味着:启动 256*65536 个线程,每个线程执行一次迭代。那么性能会好很多。
版权声明:本文内容由互联网用户自发贡献,该文观点与技术仅代表作者本人。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。如发现本站有涉嫌侵权/违法违规的内容, 请发送邮件至 dio@foxmail.com 举报,一经查实,本站将立刻删除。