/* Crack callback */
static void ocl_androidfde_crack_callback(char *line, int self)
{
int a,b,c,e;
char plain[MAX];
cl_uint16 addline;
cl_uint16 salt;
cl_uint16 salt2;
size_t nws1;
size_t nws;
/* setup addline */
addline.s0=addline.s1=addline.s2=addline.s3=addline.s4=addline.s5=addline.s6=addline.s7=addline.sF=0;
addline.sF=strlen(line);
addline.s0=line[0]|(line[1]<<8)|(line[2]<<16)|(line[3]<<24);
addline.s1=line[4]|(line[5]<<8)|(line[6]<<16)|(line[7]<<24);
addline.s2=line[8]|(line[9]<<8)|(line[10]<<16)|(line[11]<<24);
addline.s3=line[12]|(line[13]<<8)|(line[14]<<16)|(line[15]<<24);
/* setup salt */
salt=ocl_get_salt();
salt2=ocl_get_salt2();
if (attack_over!=0) pthread_exit(NULL);
if (rule_counts[self][0]==-1) return;
nws = (rule_counts[self][0] / wthreads[self].vectorsize);
while ((nws%64)!=0) nws++;
nws1 = nws*wthreads[self].vectorsize;
if (nws1==0) nws1=64;
if (nws==0) nws=64;
_clSetKernelArg(rule_kernelend[self], 0, sizeof(cl_mem), (void*) &rule_buffer[self]);
_clSetKernelArg(rule_kernelend[self], 1, sizeof(cl_mem), (void*) &rule_images2_buf[self]);
_clSetKernelArg(rule_kernelend[self], 2, sizeof(cl_mem), (void*) &rule_images4_buf[self]);
_clSetKernelArg(rule_kernelend[self], 3, sizeof(cl_uint16), (void*) &addline);
_clSetKernelArg(rule_kernelend[self], 4, sizeof(cl_uint16), (void*) &salt);
_clSetKernelArg(rule_kernelend[self], 6, sizeof(cl_uint16), (void*) &salt2);
_clSetKernelArg(rule_kernelend[self], 7, sizeof(cl_mem), (void*) &rule_found_buf[self]);
_clSetKernelArg(rule_kernelend[self], 8, sizeof(cl_mem), (void*) &rule_found_ind_buf[self]);
_clSetKernelArg(rule_kernelend2[self], 0, sizeof(cl_mem), (void*) &rule_buffer[self]);
_clSetKernelArg(rule_kernelend2[self], 1, sizeof(cl_mem), (void*) &rule_images2_buf[self]);
_clSetKernelArg(rule_kernelend2[self], 2, sizeof(cl_mem), (void*) &rule_images4_buf[self]);
_clSetKernelArg(rule_kernelend2[self], 3, sizeof(cl_uint16), (void*) &addline);
_clSetKernelArg(rule_kernelend2[self], 4, sizeof(cl_uint16), (void*) &salt);
_clSetKernelArg(rule_kernelend2[self], 6, sizeof(cl_uint16), (void*) &salt2);
_clSetKernelArg(rule_kernelend2[self], 7, sizeof(cl_mem), (void*) &rule_found_buf[self]);
_clSetKernelArg(rule_kernelend2[self], 8, sizeof(cl_mem), (void*) &rule_found_ind_buf[self]);
_clSetKernelArg(rule_kernelmod[self], 0, sizeof(cl_mem), (void*) &rule_images2_buf[self]);
_clSetKernelArg(rule_kernelmod[self], 1, sizeof(cl_mem), (void*) &rule_images_buf[self]);
_clSetKernelArg(rule_kernelmod[self], 2, sizeof(cl_mem), (void*) &rule_images3_buf[self]);
_clSetKernelArg(rule_kernelmod[self], 3, sizeof(cl_uint16), (void*) &addline);
_clSetKernelArg(rule_kernelmod[self], 4, sizeof(cl_uint16), (void*) &salt);
_clSetKernelArg(rule_kernelmod[self], 6, sizeof(cl_uint16), (void*) &salt2);
_clSetKernelArg(rule_kernelmod[self], 7, sizeof(cl_mem), (void*) &rule_sizes_buf[self]);
_clSetKernelArg(rule_kernelpre1[self], 0, sizeof(cl_mem), (void*) &rule_images4_buf[self]);
_clSetKernelArg(rule_kernelpre1[self], 1, sizeof(cl_mem), (void*) &rule_images2_buf[self]);
_clSetKernelArg(rule_kernelpre1[self], 2, sizeof(cl_mem), (void*) &rule_images3_buf[self]);
_clSetKernelArg(rule_kernelpre1[self], 3, sizeof(cl_uint16), (void*) &addline);
_clSetKernelArg(rule_kernelpre1[self], 4, sizeof(cl_uint16), (void*) &salt);
_clSetKernelArg(rule_kernelpre1[self], 6, sizeof(cl_uint16), (void*) &salt2);
_clSetKernelArg(rule_kernelbl1[self], 0, sizeof(cl_mem), (void*) &rule_images4_buf[self]);
_clSetKernelArg(rule_kernelbl1[self], 1, sizeof(cl_mem), (void*) &rule_images2_buf[self]);
_clSetKernelArg(rule_kernelbl1[self], 2, sizeof(cl_mem), (void*) &rule_images3_buf[self]);
_clSetKernelArg(rule_kernelbl1[self], 3, sizeof(cl_uint16), (void*) &addline);
_clSetKernelArg(rule_kernelbl1[self], 4, sizeof(cl_uint16), (void*) &salt);
_clSetKernelArg(rule_kernelbl1[self], 6, sizeof(cl_uint16), (void*) &salt2);
_clSetKernelArg(rule_kernelpre2[self], 0, sizeof(cl_mem), (void*) &rule_images4_buf[self]);
_clSetKernelArg(rule_kernelpre2[self], 1, sizeof(cl_mem), (void*) &rule_images2_buf[self]);
_clSetKernelArg(rule_kernelpre2[self], 2, sizeof(cl_mem), (void*) &rule_images3_buf[self]);
_clSetKernelArg(rule_kernelpre2[self], 3, sizeof(cl_uint16), (void*) &addline);
_clSetKernelArg(rule_kernelpre2[self], 4, sizeof(cl_uint16), (void*) &salt);
_clSetKernelArg(rule_kernelpre2[self], 6, sizeof(cl_uint16), (void*) &salt2);
_clSetKernelArg(rule_kernelbl2[self], 0, sizeof(cl_mem), (void*) &rule_images4_buf[self]);
_clSetKernelArg(rule_kernelbl2[self], 1, sizeof(cl_mem), (void*) &rule_images2_buf[self]);
_clSetKernelArg(rule_kernelbl2[self], 2, sizeof(cl_mem), (void*) &rule_images3_buf[self]);
_clSetKernelArg(rule_kernelbl2[self], 3, sizeof(cl_uint16), (void*) &addline);
_clSetKernelArg(rule_kernelbl2[self], 4, sizeof(cl_uint16), (void*) &salt);
_clSetKernelArg(rule_kernelbl2[self], 6, sizeof(cl_uint16), (void*) &salt2);
_clEnqueueNDRangeKernel(rule_oclqueue[self], rule_kernelmod[self], 1, NULL, &nws1, rule_local_work_size, 0, NULL, NULL);
_clFinish(rule_oclqueue[self]);
_clEnqueueNDRangeKernel(rule_oclqueue[self], rule_kernelpre1[self], 1, NULL, &nws, rule_local_work_size, 0, NULL, NULL);
_clFinish(rule_oclqueue[self]);
b=(2*2000)/1000;
for (a=0;a<2000;a+=1000)
{
if (attack_over!=0) pthread_exit(NULL);
addline.sA=a;
addline.sB=a+1000;
if (a==0) addline.sA=1;
_clSetKernelArg(rule_kernelbl1[self], 3, sizeof(cl_uint16), (void*) &addline);
_clEnqueueNDRangeKernel(rule_oclqueue[self], rule_kernelbl1[self], 1, NULL, &nws, rule_local_work_size, 0, NULL, NULL);
_clFinish(rule_oclqueue[self]);
wthreads[self].tries+=(nws1)/b;
pthread_mutex_lock(&wthreads[self].tempmutex);
pthread_mutex_unlock(&wthreads[self].tempmutex);
}
_clEnqueueNDRangeKernel(rule_oclqueue[self], rule_kernelpre2[self], 1, NULL, &nws, rule_local_work_size, 0, NULL, NULL);
_clFinish(rule_oclqueue[self]);
for (a=0;a<2000;a+=1000)
{
if (attack_over!=0) pthread_exit(NULL);
addline.sA=a;
addline.sB=a+1000;
if (addline.sB>2000) addline.sB=2000;
if (a==0) addline.sA=1;
_clSetKernelArg(rule_kernelbl2[self], 3, sizeof(cl_uint16), (void*) &addline);
_clEnqueueNDRangeKernel(rule_oclqueue[self], rule_kernelbl2[self], 1, NULL, &nws, rule_local_work_size, 0, NULL, NULL);
_clFinish(rule_oclqueue[self]);
wthreads[self].tries+=(nws1)/b;
pthread_mutex_lock(&wthreads[self].tempmutex);
pthread_mutex_unlock(&wthreads[self].tempmutex);
}
_clEnqueueNDRangeKernel(rule_oclqueue[self], rule_kernelend[self], 1, NULL, &nws, rule_local_work_size, 0, NULL, NULL);
_clEnqueueReadBuffer(rule_oclqueue[self], rule_buffer[self], CL_TRUE, 0, hash_ret_len1*wthreads[self].vectorsize*nws, rule_ptr[self], 0, NULL, NULL);
for (a=0;a<nws;a++)
for (c=0;c<wthreads[self].vectorsize;c++)
{
e=(a)*wthreads[self].vectorsize+c;
if (check_androidfde(rule_ptr[self]+e*32)==hash_ok)
{
strcpy(plain,&rule_images[self][0]+(e*MAX));
strcat(plain,line);
add_cracked_list(hash_list->username, hash_list->hash, hash_list->salt, plain);
}
}
}
/* Crack callback */
static void ocl_wpa_crack_callback(char *line, int self)
{
int a,b,c,e;
int *found;
int err;
char plain[MAX];
cl_uint16 addline;
cl_uint16 salt;
cl_uint16 salt2;
/* setup addline */
addline.s0=addline.s1=addline.s2=addline.s3=addline.s4=addline.s5=addline.s6=addline.s7=addline.sF=0;
addline.sF=strlen(line);
addline.s0=line[0]|(line[1]<<8)|(line[2]<<16)|(line[3]<<24);
addline.s1=line[4]|(line[5]<<8)|(line[6]<<16)|(line[7]<<24);
addline.s2=line[8]|(line[9]<<8)|(line[10]<<16)|(line[11]<<24);
addline.s3=line[12]|(line[13]<<8)|(line[14]<<16)|(line[15]<<24);
/* setup salt */
salt=ocl_get_salt();
salt2=ocl_get_salt2();
if (attack_over!=0) pthread_exit(NULL);
pthread_mutex_lock(&wthreads[self].tempmutex);
pthread_mutex_unlock(&wthreads[self].tempmutex);
size_t nws1=ocl_rule_workset[self]*wthreads[self].vectorsize;
size_t nws=ocl_rule_workset[self];
_clSetKernelArg(rule_kernelend[self], 0, sizeof(cl_mem), (void*) &rule_buffer[self]);
_clSetKernelArg(rule_kernelend[self], 1, sizeof(cl_mem), (void*) &rule_images2_buf[self]);
_clSetKernelArg(rule_kernelend[self], 2, sizeof(cl_mem), (void*) &rule_images4_buf[self]);
_clSetKernelArg(rule_kernelend[self], 3, sizeof(cl_uint16), (void*) &addline);
_clSetKernelArg(rule_kernelend[self], 4, sizeof(cl_uint16), (void*) &salt);
_clSetKernelArg(rule_kernelend[self], 6, sizeof(cl_uint16), (void*) &salt2);
_clSetKernelArg(rule_kernelend[self], 7, sizeof(cl_mem), (void*) &block_buf[self]);
_clSetKernelArg(rule_kernelend[self], 8, sizeof(cl_mem), (void*) &eapol_buf[self]);
_clSetKernelArg(rule_kernelend[self], 9, sizeof(cl_mem), (void*) &rule_found_buf[self]);
_clSetKernelArg(rule_kernelend[self], 10, sizeof(cl_mem), (void*) &rule_found_ind_buf[self]);
_clSetKernelArg(rule_kernelmod[self], 0, sizeof(cl_mem), (void*) &rule_images2_buf[self]);
_clSetKernelArg(rule_kernelmod[self], 1, sizeof(cl_mem), (void*) &rule_images_buf[self]);
_clSetKernelArg(rule_kernelmod[self], 2, sizeof(cl_mem), (void*) &rule_images3_buf[self]);
_clSetKernelArg(rule_kernelmod[self], 3, sizeof(cl_uint16), (void*) &addline);
_clSetKernelArg(rule_kernelmod[self], 4, sizeof(cl_uint16), (void*) &salt);
_clSetKernelArg(rule_kernelmod[self], 6, sizeof(cl_uint16), (void*) &salt2);
_clSetKernelArg(rule_kernelmod[self], 7, sizeof(cl_mem), (void*) &block_buf[self]);
_clSetKernelArg(rule_kernelmod[self], 8, sizeof(cl_mem), (void*) &eapol_buf[self]);
_clSetKernelArg(rule_kernelmod[self], 9, sizeof(cl_mem), (void*) &rule_sizes_buf[self]);
_clSetKernelArg(rule_kernelpre1[self], 0, sizeof(cl_mem), (void*) &rule_images4_buf[self]);
_clSetKernelArg(rule_kernelpre1[self], 1, sizeof(cl_mem), (void*) &rule_images2_buf[self]);
_clSetKernelArg(rule_kernelpre1[self], 2, sizeof(cl_mem), (void*) &rule_images3_buf[self]);
_clSetKernelArg(rule_kernelpre1[self], 3, sizeof(cl_uint16), (void*) &addline);
_clSetKernelArg(rule_kernelpre1[self], 4, sizeof(cl_uint16), (void*) &salt);
_clSetKernelArg(rule_kernelpre1[self], 6, sizeof(cl_uint16), (void*) &salt2);
_clSetKernelArg(rule_kernelpre1[self], 7, sizeof(cl_mem), (void*) &block_buf[self]);
_clSetKernelArg(rule_kernelpre1[self], 8, sizeof(cl_mem), (void*) &eapol_buf[self]);
_clSetKernelArg(rule_kernelbl1[self], 0, sizeof(cl_mem), (void*) &rule_images4_buf[self]);
_clSetKernelArg(rule_kernelbl1[self], 1, sizeof(cl_mem), (void*) &rule_images2_buf[self]);
_clSetKernelArg(rule_kernelbl1[self], 2, sizeof(cl_mem), (void*) &rule_images3_buf[self]);
_clSetKernelArg(rule_kernelbl1[self], 3, sizeof(cl_uint16), (void*) &addline);
_clSetKernelArg(rule_kernelbl1[self], 4, sizeof(cl_uint16), (void*) &salt);
_clSetKernelArg(rule_kernelbl1[self], 6, sizeof(cl_uint16), (void*) &salt2);
_clSetKernelArg(rule_kernelbl1[self], 7, sizeof(cl_mem), (void*) &block_buf[self]);
_clSetKernelArg(rule_kernelbl1[self], 8, sizeof(cl_mem), (void*) &eapol_buf[self]);
_clSetKernelArg(rule_kernelpre2[self], 0, sizeof(cl_mem), (void*) &rule_images4_buf[self]);
_clSetKernelArg(rule_kernelpre2[self], 1, sizeof(cl_mem), (void*) &rule_images2_buf[self]);
_clSetKernelArg(rule_kernelpre2[self], 2, sizeof(cl_mem), (void*) &rule_images3_buf[self]);
_clSetKernelArg(rule_kernelpre2[self], 3, sizeof(cl_uint16), (void*) &addline);
_clSetKernelArg(rule_kernelpre2[self], 4, sizeof(cl_uint16), (void*) &salt);
_clSetKernelArg(rule_kernelpre2[self], 6, sizeof(cl_uint16), (void*) &salt2);
_clSetKernelArg(rule_kernelpre2[self], 7, sizeof(cl_mem), (void*) &block_buf[self]);
_clSetKernelArg(rule_kernelpre2[self], 8, sizeof(cl_mem), (void*) &eapol_buf[self]);
_clSetKernelArg(rule_kernelbl2[self], 0, sizeof(cl_mem), (void*) &rule_images4_buf[self]);
_clSetKernelArg(rule_kernelbl2[self], 1, sizeof(cl_mem), (void*) &rule_images2_buf[self]);
_clSetKernelArg(rule_kernelbl2[self], 2, sizeof(cl_mem), (void*) &rule_images3_buf[self]);
_clSetKernelArg(rule_kernelbl2[self], 3, sizeof(cl_uint16), (void*) &addline);
_clSetKernelArg(rule_kernelbl2[self], 4, sizeof(cl_uint16), (void*) &salt);
_clSetKernelArg(rule_kernelbl2[self], 6, sizeof(cl_uint16), (void*) &salt2);
_clSetKernelArg(rule_kernelbl2[self], 7, sizeof(cl_mem), (void*) &block_buf[self]);
_clSetKernelArg(rule_kernelbl2[self], 8, sizeof(cl_mem), (void*) &eapol_buf[self]);
_clEnqueueNDRangeKernel(rule_oclqueue[self], rule_kernelmod[self], 1, NULL, &nws1, rule_local_work_size, 0, NULL, NULL);
_clFinish(rule_oclqueue[self]);
_clEnqueueNDRangeKernel(rule_oclqueue[self], rule_kernelpre1[self], 1, NULL, &nws, rule_local_work_size, 0, NULL, NULL);
_clFinish(rule_oclqueue[self]);
for (a=0;a<7;a++)
{
if (attack_over==1) pthread_exit(NULL);
addline.sA=a*1170;
_clSetKernelArg(rule_kernelbl1[self], 3, sizeof(cl_uint16), (void*) &addline);
_clEnqueueNDRangeKernel(rule_oclqueue[self], rule_kernelbl1[self], 1, NULL, &nws, rule_local_work_size, 0, NULL, NULL);
_clFinish(rule_oclqueue[self]);
wthreads[self].tries+=(ocl_rule_workset[self]*wthreads[self].vectorsize)/14;
}
_clEnqueueNDRangeKernel(rule_oclqueue[self], rule_kernelpre2[self], 1, NULL, &nws, rule_local_work_size, 0, NULL, NULL);
_clFinish(rule_oclqueue[self]);
for (a=0;a<7;a++)
{
if (attack_over==1) pthread_exit(NULL);
addline.sA=a*1170;
_clSetKernelArg(rule_kernelbl2[self], 3, sizeof(cl_uint16), (void*) &addline);
_clEnqueueNDRangeKernel(rule_oclqueue[self], rule_kernelbl2[self], 1, NULL, &nws, rule_local_work_size, 0, NULL, NULL);
_clFinish(rule_oclqueue[self]);
wthreads[self].tries+=(ocl_rule_workset[self]*wthreads[self].vectorsize)/14;
}
_clEnqueueNDRangeKernel(rule_oclqueue[self], rule_kernelend[self], 1, NULL, &nws, rule_local_work_size, 0, NULL, NULL);
found = _clEnqueueMapBuffer(rule_oclqueue[self], rule_found_buf[self], CL_TRUE,CL_MAP_READ, 0, 4, 0, 0, NULL, &err);
if (err!=CL_SUCCESS) return;
if (*found>0)
{
_clEnqueueReadBuffer(rule_oclqueue[self], rule_found_ind_buf[self], CL_TRUE, 0, ocl_rule_workset[self]*sizeof(cl_uint), rule_found_ind[self], 0, NULL, NULL);
for (a=0;a<ocl_rule_workset[self];a++)
if (rule_found_ind[self][a]==1)
{
b=a*wthreads[self].vectorsize;
_clEnqueueReadBuffer(rule_oclqueue[self], rule_buffer[self], CL_TRUE, b*hash_ret_len1, hash_ret_len1*wthreads[self].vectorsize, rule_ptr[self]+b*hash_ret_len1, 0, NULL, NULL);
for (c=0;c<wthreads[self].vectorsize;c++)
{
e=(a)*wthreads[self].vectorsize+c;
if (memcmp(hccap.keymic, (char *)rule_ptr[self]+(e)*hash_ret_len1, hash_ret_len1-1) == 0)
{
strcpy(plain,&rule_images[self][0]+(e*MAX));
strcat(plain,line);
add_cracked_list(hash_list->username, hash_list->hash, hash_list->salt, plain);
}
}
}
bzero(rule_found_ind[self],ocl_rule_workset[self]*sizeof(cl_uint));
_clEnqueueWriteBuffer(rule_oclqueue[self], rule_found_ind_buf[self], CL_FALSE, 0, ocl_rule_workset[self]*sizeof(cl_uint), rule_found_ind[self], 0, NULL, NULL);
*found = 0;
_clEnqueueWriteBuffer(rule_oclqueue[self], rule_found_buf[self], CL_FALSE, 0, 4, found, 0, NULL, NULL);
}
_clEnqueueUnmapMemObject(rule_oclqueue[self],rule_found_buf[self],(void *)found,0,NULL,NULL);
}
