/* 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); }