arno/hashcat
1
0
Fork 0
mirror of https://github.com/hashcat/hashcat.git synced 2024-11-26 09:58:16 +00:00
Code Issues Releases Wiki Activity

Merge branch 'master' of https://github.com/hashcat/hashcat

Browse Source
This commit is contained in:
Jens Steube 2018-02-13 10:27:55 +01:00
commit 7d09ae93f2
1 changed files with 18 additions and 12 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

30
src/opencl.c
View File

@ -1416,8 +1416,6 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param,
wpa_t *wpa = &wpas[digests_offset + loops_pos]; wpa_t *wpa = &wpas[digests_offset + loops_pos];
printf ("%u\n", wpa->keyver);
if (wpa->keyver == 1) if (wpa->keyver == 1)
{ {
CL_rc = run_kernel (hashcat_ctx, device_param, KERN_RUN_AUX1, pws_cnt, false, 0); CL_rc = run_kernel (hashcat_ctx, device_param, KERN_RUN_AUX1, pws_cnt, false, 0);
@ -1742,20 +1740,28 @@ int run_kernel_mp (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param,
case KERN_RUN_MP_L: device_param->kernel_params_mp_l_buf64[9] = num; break; case KERN_RUN_MP_L: device_param->kernel_params_mp_l_buf64[9] = num; break;
} }
u64 kernel_threads = 0;
cl_kernel kernel = NULL; cl_kernel kernel = NULL;
switch (kern_run) switch (kern_run)
{ {
case KERN_RUN_MP: kernel = device_param->kernel_mp; break; case KERN_RUN_MP:
case KERN_RUN_MP_R: kernel = device_param->kernel_mp_r; break; kernel = device_param->kernel_mp;
case KERN_RUN_MP_L: kernel = device_param->kernel_mp_l; break; kernel_threads = device_param->kernel_wgs_mp;
break;
case KERN_RUN_MP_R:
kernel = device_param->kernel_mp_r;
kernel_threads = device_param->kernel_wgs_mp_r;
break;
case KERN_RUN_MP_L:
kernel = device_param->kernel_mp_l;
kernel_threads = device_param->kernel_wgs_mp_l;
break;
default: default:
event_log_error (hashcat_ctx, "Invalid kernel specified."); event_log_error (hashcat_ctx, "Invalid kernel specified.");
return -1; return -1;
} }
const u64 kernel_threads = device_param->device_maxworkgroup_size;
while (num_elements % kernel_threads) num_elements++; while (num_elements % kernel_threads) num_elements++;
int CL_rc; int CL_rc;
@ -1808,7 +1814,7 @@ int run_kernel_tm (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param)
{ {
const u64 num_elements = 1024; // fixed const u64 num_elements = 1024; // fixed
const u64 kernel_threads = device_param->device_maxworkgroup_size; const u64 kernel_threads = device_param->kernel_wgs_tm;
cl_kernel kernel = device_param->kernel_tm; cl_kernel kernel = device_param->kernel_tm;
@ -1838,7 +1844,7 @@ int run_kernel_amp (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param,
device_param->kernel_params_amp_buf64[6] = num_elements; device_param->kernel_params_amp_buf64[6] = num_elements;
const u64 kernel_threads = device_param->device_maxworkgroup_size; const u64 kernel_threads = device_param->kernel_wgs_amp;
while (num_elements % kernel_threads) num_elements++; while (num_elements % kernel_threads) num_elements++;
@ -1874,7 +1880,7 @@ int run_kernel_atinit (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_par
device_param->kernel_params_atinit_buf64[1] = num_elements; device_param->kernel_params_atinit_buf64[1] = num_elements;
const u64 kernel_threads = device_param->device_maxworkgroup_size; const u64 kernel_threads = device_param->kernel_wgs_atinit;
while (num_elements % kernel_threads) num_elements++; while (num_elements % kernel_threads) num_elements++;
@ -1918,7 +1924,7 @@ int run_kernel_memset (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_par
device_param->kernel_params_memset_buf32[1] = value; device_param->kernel_params_memset_buf32[1] = value;
device_param->kernel_params_memset_buf64[2] = num16d; device_param->kernel_params_memset_buf64[2] = num16d;
const u64 kernel_threads = device_param->device_maxworkgroup_size; const u64 kernel_threads = device_param->kernel_wgs_memset;
u64 num_elements = num16d; u64 num_elements = num16d;
@ -1973,7 +1979,7 @@ int run_kernel_decompress (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device
device_param->kernel_params_decompress_buf64[3] = num_elements; device_param->kernel_params_decompress_buf64[3] = num_elements;
const u64 kernel_threads = device_param->device_maxworkgroup_size; const u64 kernel_threads = device_param->kernel_wgs_decompress;
while (num_elements % kernel_threads) num_elements++; while (num_elements % kernel_threads) num_elements++;