arno/hashcat
1
0
Fork 0
mirror of https://github.com/hashcat/hashcat.git synced 2025-02-17 01:52:06 +00:00
Code Issues Releases Wiki Activity

More CUDA rewrites

Browse Source
This commit is contained in:
Jens Steube 2019-05-05 11:57:54 +02:00
parent ec9925f3b1
commit 08dc1acc02
4 changed files with 558 additions and 215 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

18
src/autotune.c
View File

@ -236,8 +236,6 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param
double exec_msec_pre_final = try_run (hashcat_ctx, device_param, kernel_accel, kernel_loops); double exec_msec_pre_final = try_run (hashcat_ctx, device_param, kernel_accel, kernel_loops);
printf ("%f\n", exec_msec_pre_final);
const u32 exec_left = (const u32) (target_msec / exec_msec_pre_final); const u32 exec_left = (const u32) (target_msec / exec_msec_pre_final);
const u32 accel_left = kernel_accel_max / kernel_accel; const u32 accel_left = kernel_accel_max / kernel_accel;
@ -255,23 +253,23 @@ printf ("%f\n", exec_msec_pre_final);
{ {
// reset them fake words // reset them fake words
CL_rc = run_cuda_kernel_memset (hashcat_ctx, device_param, device_param->cuda_d_pws_buf, 0, device_param->size_pws); CU_rc = run_cuda_kernel_memset (hashcat_ctx, device_param, device_param->cuda_d_pws_buf, 0, device_param->size_pws);
if (CL_rc == -1) return -1; if (CU_rc == -1) return -1;
// reset other buffers in case autotune cracked something // reset other buffers in case autotune cracked something
CL_rc = run_cuda_kernel_memset (hashcat_ctx, device_param, device_param->cuda_d_plain_bufs, 0, device_param->size_plains); CU_rc = run_cuda_kernel_memset (hashcat_ctx, device_param, device_param->cuda_d_plain_bufs, 0, device_param->size_plains);
if (CL_rc == -1) return -1; if (CU_rc == -1) return -1;
CL_rc = run_cuda_kernel_memset (hashcat_ctx, device_param, device_param->cuda_d_digests_shown, 0, device_param->size_shown); CU_rc = run_cuda_kernel_memset (hashcat_ctx, device_param, device_param->cuda_d_digests_shown, 0, device_param->size_shown);
if (CL_rc == -1) return -1; if (CU_rc == -1) return -1;
CL_rc = run_cuda_kernel_memset (hashcat_ctx, device_param, device_param->cuda_d_result, 0, device_param->size_results); CU_rc = run_cuda_kernel_memset (hashcat_ctx, device_param, device_param->cuda_d_result, 0, device_param->size_results);
if (CL_rc == -1) return -1; if (CU_rc == -1) return -1;
} }
if (device_param->is_opencl == true) if (device_param->is_opencl == true)

474
src/backend.c
View File

@ -2438,22 +2438,47 @@ int gidd_to_pw_t (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, c
{ {
pw_idx_t pw_idx; pw_idx_t pw_idx;
int CL_rc; pw_idx.off = 0;
pw_idx.cnt = 0;
pw_idx.len = 0;
CL_rc = hc_clEnqueueReadBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_idx, CL_TRUE, gidd * sizeof (pw_idx_t), sizeof (pw_idx_t), &pw_idx, 0, NULL, NULL); if (device_param->is_cuda == true)
{
const int CU_rc = hc_cuMemcpyDtoH (hashcat_ctx, &pw_idx, device_param->cuda_d_pws_idx + (gidd * sizeof (pw_idx_t)), sizeof (pw_idx_t));
if (CU_rc == -1) return -1;
}
if (device_param->is_opencl == true)
{
const int CL_rc = hc_clEnqueueReadBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_idx, CL_TRUE, gidd * sizeof (pw_idx_t), sizeof (pw_idx_t), &pw_idx, 0, NULL, NULL);
if (CL_rc == -1) return -1; if (CL_rc == -1) return -1;
}
const u32 off = pw_idx.off; const u32 off = pw_idx.off;
const u32 cnt = pw_idx.cnt; const u32 cnt = pw_idx.cnt;
const u32 len = pw_idx.len; const u32 len = pw_idx.len;
if (device_param->is_cuda == true)
{
if (cnt > 0) if (cnt > 0)
{ {
CL_rc = hc_clEnqueueReadBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_comp_buf, CL_TRUE, off * sizeof (u32), cnt * sizeof (u32), pw->i, 0, NULL, NULL); const int CU_rc = hc_cuMemcpyDtoH (hashcat_ctx,pw->i, device_param->cuda_d_pws_comp_buf + (off * sizeof (u32)), cnt * sizeof (u32));
if (CU_rc == -1) return -1;
}
}
if (device_param->is_opencl == true)
{
if (cnt > 0)
{
const int CL_rc = hc_clEnqueueReadBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_comp_buf, CL_TRUE, off * sizeof (u32), cnt * sizeof (u32), pw->i, 0, NULL, NULL);
if (CL_rc == -1) return -1; if (CL_rc == -1) return -1;
} }
}
for (u32 i = cnt; i < 64; i++) for (u32 i = cnt; i < 64; i++)
{ {
@ -2478,6 +2503,7 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param,
return process_stdout (hashcat_ctx, device_param, pws_cnt); return process_stdout (hashcat_ctx, device_param, pws_cnt);
} }
int CU_rc;
int CL_rc; int CL_rc;
if (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL) if (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL)
@ -2493,20 +2519,40 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param,
{ {
const u32 size_tm = 32 * sizeof (bs_word_t); const u32 size_tm = 32 * sizeof (bs_word_t);
if (device_param->is_cuda == true)
{
CU_rc = run_cuda_kernel_bzero (hashcat_ctx, device_param, device_param->cuda_d_tm_c, size_tm);
if (CU_rc == -1) return -1;
}
if (device_param->is_opencl == true)
{
CL_rc = run_opencl_kernel_bzero (hashcat_ctx, device_param, device_param->opencl_d_tm_c, size_tm); CL_rc = run_opencl_kernel_bzero (hashcat_ctx, device_param, device_param->opencl_d_tm_c, size_tm);
if (CL_rc == -1) return -1; if (CL_rc == -1) return -1;
}
CL_rc = run_kernel_tm (hashcat_ctx, device_param); const int rc_kernel_tm = run_kernel_tm (hashcat_ctx, device_param);
if (CL_rc == -1) return -1; if (rc_kernel_tm == -1) return -1;
if (device_param->is_cuda == true)
{
const int CU_rc = hc_cuMemcpyDtoD (hashcat_ctx, device_param->cuda_d_bfs_c, device_param->cuda_d_tm_c, size_tm);
if (CU_rc == -1) return -1;
}
if (device_param->is_opencl == true)
{
CL_rc = hc_clEnqueueCopyBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_tm_c, device_param->opencl_d_bfs_c, 0, 0, size_tm, 0, NULL, NULL); CL_rc = hc_clEnqueueCopyBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_tm_c, device_param->opencl_d_bfs_c, 0, 0, size_tm, 0, NULL, NULL);
if (CL_rc == -1) return -1; if (CL_rc == -1) return -1;
} }
} }
} }
}
if (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) if (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL)
{ {
@ -2544,6 +2590,7 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param,
if (run_init == true) if (run_init == true)
{ {
//tbd
CL_rc = hc_clEnqueueCopyBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_amp_buf, device_param->opencl_d_pws_buf, 0, 0, pws_cnt * sizeof (pw_t), 0, NULL, NULL); CL_rc = hc_clEnqueueCopyBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_amp_buf, device_param->opencl_d_pws_buf, 0, 0, pws_cnt * sizeof (pw_t), 0, NULL, NULL);
if (CL_rc == -1) return -1; if (CL_rc == -1) return -1;
@ -2568,12 +2615,14 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param,
if (CL_rc == -1) return -1; if (CL_rc == -1) return -1;
//tbd
CL_rc = hc_clEnqueueReadBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_hooks, CL_TRUE, 0, device_param->size_hooks, device_param->hooks_buf, 0, NULL, NULL); CL_rc = hc_clEnqueueReadBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_hooks, CL_TRUE, 0, device_param->size_hooks, device_param->hooks_buf, 0, NULL, NULL);
if (CL_rc == -1) return -1; if (CL_rc == -1) return -1;
module_ctx->module_hook12 (device_param, hashes->hook_salts_buf, salt_pos, pws_cnt); module_ctx->module_hook12 (device_param, hashes->hook_salts_buf, salt_pos, pws_cnt);
//tbd
CL_rc = hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_hooks, CL_TRUE, 0, device_param->size_hooks, device_param->hooks_buf, 0, NULL, NULL); CL_rc = hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_hooks, CL_TRUE, 0, device_param->size_hooks, device_param->hooks_buf, 0, NULL, NULL);
if (CL_rc == -1) return -1; if (CL_rc == -1) return -1;
@ -2640,12 +2689,14 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param,
if (CL_rc == -1) return -1; if (CL_rc == -1) return -1;
//tbd
CL_rc = hc_clEnqueueReadBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_hooks, CL_TRUE, 0, device_param->size_hooks, device_param->hooks_buf, 0, NULL, NULL); CL_rc = hc_clEnqueueReadBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_hooks, CL_TRUE, 0, device_param->size_hooks, device_param->hooks_buf, 0, NULL, NULL);
if (CL_rc == -1) return -1; if (CL_rc == -1) return -1;
module_ctx->module_hook23 (device_param, hashes->hook_salts_buf, salt_pos, pws_cnt); module_ctx->module_hook23 (device_param, hashes->hook_salts_buf, salt_pos, pws_cnt);
//tbd
CL_rc = hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_hooks, CL_TRUE, 0, device_param->size_hooks, device_param->hooks_buf, 0, NULL, NULL); CL_rc = hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_hooks, CL_TRUE, 0, device_param->size_hooks, device_param->hooks_buf, 0, NULL, NULL);
if (CL_rc == -1) return -1; if (CL_rc == -1) return -1;
@ -3311,7 +3362,15 @@ int run_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, con
int run_kernel_mp (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const u32 kern_run, const u64 num) int run_kernel_mp (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const u32 kern_run, const u64 num)
{ {
puts ("run_kernel_mp"); u64 kernel_threads = 0;
switch (kern_run)
{
case KERN_RUN_MP: kernel_threads = device_param->kernel_wgs_mp; break;
case KERN_RUN_MP_R: kernel_threads = device_param->kernel_wgs_mp_r; break;
case KERN_RUN_MP_L: kernel_threads = device_param->kernel_wgs_mp_l; break;
}
u64 num_elements = num; u64 num_elements = num;
switch (kern_run) switch (kern_run)
@ -3321,62 +3380,81 @@ puts ("run_kernel_mp");
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; if (device_param->is_cuda == true)
cl_kernel kernel = NULL; {
CUfunction cuda_function = NULL;
void **cuda_args = NULL;
switch (kern_run) switch (kern_run)
{ {
case KERN_RUN_MP: case KERN_RUN_MP: cuda_function = device_param->cuda_function_mp;
kernel = device_param->opencl_kernel_mp; cuda_args = device_param->kernel_params_mp;
kernel_threads = device_param->kernel_wgs_mp;
break; break;
case KERN_RUN_MP_R: case KERN_RUN_MP_R: cuda_function = device_param->cuda_function_mp_r;
kernel = device_param->opencl_kernel_mp_r; cuda_args = device_param->kernel_params_mp_r;
kernel_threads = device_param->kernel_wgs_mp_r;
break; break;
case KERN_RUN_MP_L: case KERN_RUN_MP_L: cuda_function = device_param->cuda_function_mp_l;
kernel = device_param->opencl_kernel_mp_l; cuda_args = device_param->kernel_params_mp_l;
kernel_threads = device_param->kernel_wgs_mp_l; break;
}
num_elements = CEILDIV (num_elements, kernel_threads);
const int rc_cuLaunchKernel = hc_cuLaunchKernel (hashcat_ctx, cuda_function, num_elements, 1, 1, kernel_threads, 1, 1, 0, NULL, cuda_args, NULL);
if (rc_cuLaunchKernel == -1) return -1;
const int rc_cuCtxSynchronize = hc_cuCtxSynchronize (hashcat_ctx);
if (rc_cuCtxSynchronize == -1) return -1;
}
if (device_param->is_opencl == true)
{
int CL_rc;
cl_kernel opencl_kernel = NULL;
switch (kern_run)
{
case KERN_RUN_MP: opencl_kernel = device_param->opencl_kernel_mp; break;
case KERN_RUN_MP_R: opencl_kernel = device_param->opencl_kernel_mp_r; break;
case KERN_RUN_MP_L: opencl_kernel = device_param->opencl_kernel_mp_l; break;
}
switch (kern_run)
{
case KERN_RUN_MP: CL_rc = hc_clSetKernelArg (hashcat_ctx, opencl_kernel, 3, sizeof (cl_ulong), device_param->kernel_params_mp[3]); if (CL_rc == -1) return -1;
CL_rc = hc_clSetKernelArg (hashcat_ctx, opencl_kernel, 4, sizeof (cl_uint), device_param->kernel_params_mp[4]); if (CL_rc == -1) return -1;
CL_rc = hc_clSetKernelArg (hashcat_ctx, opencl_kernel, 5, sizeof (cl_uint), device_param->kernel_params_mp[5]); if (CL_rc == -1) return -1;
CL_rc = hc_clSetKernelArg (hashcat_ctx, opencl_kernel, 6, sizeof (cl_uint), device_param->kernel_params_mp[6]); if (CL_rc == -1) return -1;
CL_rc = hc_clSetKernelArg (hashcat_ctx, opencl_kernel, 7, sizeof (cl_uint), device_param->kernel_params_mp[7]); if (CL_rc == -1) return -1;
CL_rc = hc_clSetKernelArg (hashcat_ctx, opencl_kernel, 8, sizeof (cl_ulong), device_param->kernel_params_mp[8]); if (CL_rc == -1) return -1;
break;
case KERN_RUN_MP_R: CL_rc = hc_clSetKernelArg (hashcat_ctx, opencl_kernel, 3, sizeof (cl_ulong), device_param->kernel_params_mp_r[3]); if (CL_rc == -1) return -1;
CL_rc = hc_clSetKernelArg (hashcat_ctx, opencl_kernel, 4, sizeof (cl_uint), device_param->kernel_params_mp_r[4]); if (CL_rc == -1) return -1;
CL_rc = hc_clSetKernelArg (hashcat_ctx, opencl_kernel, 5, sizeof (cl_uint), device_param->kernel_params_mp_r[5]); if (CL_rc == -1) return -1;
CL_rc = hc_clSetKernelArg (hashcat_ctx, opencl_kernel, 6, sizeof (cl_uint), device_param->kernel_params_mp_r[6]); if (CL_rc == -1) return -1;
CL_rc = hc_clSetKernelArg (hashcat_ctx, opencl_kernel, 7, sizeof (cl_uint), device_param->kernel_params_mp_r[7]); if (CL_rc == -1) return -1;
CL_rc = hc_clSetKernelArg (hashcat_ctx, opencl_kernel, 8, sizeof (cl_ulong), device_param->kernel_params_mp_r[8]); if (CL_rc == -1) return -1;
break;
case KERN_RUN_MP_L: CL_rc = hc_clSetKernelArg (hashcat_ctx, opencl_kernel, 3, sizeof (cl_ulong), device_param->kernel_params_mp_l[3]); if (CL_rc == -1) return -1;
CL_rc = hc_clSetKernelArg (hashcat_ctx, opencl_kernel, 4, sizeof (cl_uint), device_param->kernel_params_mp_l[4]); if (CL_rc == -1) return -1;
CL_rc = hc_clSetKernelArg (hashcat_ctx, opencl_kernel, 5, sizeof (cl_uint), device_param->kernel_params_mp_l[5]); if (CL_rc == -1) return -1;
CL_rc = hc_clSetKernelArg (hashcat_ctx, opencl_kernel, 6, sizeof (cl_uint), device_param->kernel_params_mp_l[6]); if (CL_rc == -1) return -1;
CL_rc = hc_clSetKernelArg (hashcat_ctx, opencl_kernel, 7, sizeof (cl_uint), device_param->kernel_params_mp_l[7]); if (CL_rc == -1) return -1;
CL_rc = hc_clSetKernelArg (hashcat_ctx, opencl_kernel, 8, sizeof (cl_uint), device_param->kernel_params_mp_l[8]); if (CL_rc == -1) return -1;
CL_rc = hc_clSetKernelArg (hashcat_ctx, opencl_kernel, 9, sizeof (cl_ulong), device_param->kernel_params_mp_l[9]); if (CL_rc == -1) return -1;
break; break;
default:
event_log_error (hashcat_ctx, "Invalid kernel specified.");
return -1;
} }
num_elements = round_up_multiple_64 (num_elements, kernel_threads); num_elements = round_up_multiple_64 (num_elements, kernel_threads);
int CL_rc;
switch (kern_run)
{
case KERN_RUN_MP: CL_rc = hc_clSetKernelArg (hashcat_ctx, kernel, 3, sizeof (cl_ulong), device_param->kernel_params_mp[3]); if (CL_rc == -1) return -1;
CL_rc = hc_clSetKernelArg (hashcat_ctx, kernel, 4, sizeof (cl_uint), device_param->kernel_params_mp[4]); if (CL_rc == -1) return -1;
CL_rc = hc_clSetKernelArg (hashcat_ctx, kernel, 5, sizeof (cl_uint), device_param->kernel_params_mp[5]); if (CL_rc == -1) return -1;
CL_rc = hc_clSetKernelArg (hashcat_ctx, kernel, 6, sizeof (cl_uint), device_param->kernel_params_mp[6]); if (CL_rc == -1) return -1;
CL_rc = hc_clSetKernelArg (hashcat_ctx, kernel, 7, sizeof (cl_uint), device_param->kernel_params_mp[7]); if (CL_rc == -1) return -1;
CL_rc = hc_clSetKernelArg (hashcat_ctx, kernel, 8, sizeof (cl_ulong), device_param->kernel_params_mp[8]); if (CL_rc == -1) return -1;
break;
case KERN_RUN_MP_R: CL_rc = hc_clSetKernelArg (hashcat_ctx, kernel, 3, sizeof (cl_ulong), device_param->kernel_params_mp_r[3]); if (CL_rc == -1) return -1;
CL_rc = hc_clSetKernelArg (hashcat_ctx, kernel, 4, sizeof (cl_uint), device_param->kernel_params_mp_r[4]); if (CL_rc == -1) return -1;
CL_rc = hc_clSetKernelArg (hashcat_ctx, kernel, 5, sizeof (cl_uint), device_param->kernel_params_mp_r[5]); if (CL_rc == -1) return -1;
CL_rc = hc_clSetKernelArg (hashcat_ctx, kernel, 6, sizeof (cl_uint), device_param->kernel_params_mp_r[6]); if (CL_rc == -1) return -1;
CL_rc = hc_clSetKernelArg (hashcat_ctx, kernel, 7, sizeof (cl_uint), device_param->kernel_params_mp_r[7]); if (CL_rc == -1) return -1;
CL_rc = hc_clSetKernelArg (hashcat_ctx, kernel, 8, sizeof (cl_ulong), device_param->kernel_params_mp_r[8]); if (CL_rc == -1) return -1;
break;
case KERN_RUN_MP_L: CL_rc = hc_clSetKernelArg (hashcat_ctx, kernel, 3, sizeof (cl_ulong), device_param->kernel_params_mp_l[3]); if (CL_rc == -1) return -1;
CL_rc = hc_clSetKernelArg (hashcat_ctx, kernel, 4, sizeof (cl_uint), device_param->kernel_params_mp_l[4]); if (CL_rc == -1) return -1;
CL_rc = hc_clSetKernelArg (hashcat_ctx, kernel, 5, sizeof (cl_uint), device_param->kernel_params_mp_l[5]); if (CL_rc == -1) return -1;
CL_rc = hc_clSetKernelArg (hashcat_ctx, kernel, 6, sizeof (cl_uint), device_param->kernel_params_mp_l[6]); if (CL_rc == -1) return -1;
CL_rc = hc_clSetKernelArg (hashcat_ctx, kernel, 7, sizeof (cl_uint), device_param->kernel_params_mp_l[7]); if (CL_rc == -1) return -1;
CL_rc = hc_clSetKernelArg (hashcat_ctx, kernel, 8, sizeof (cl_uint), device_param->kernel_params_mp_l[8]); if (CL_rc == -1) return -1;
CL_rc = hc_clSetKernelArg (hashcat_ctx, kernel, 9, sizeof (cl_ulong), device_param->kernel_params_mp_l[9]); if (CL_rc == -1) return -1;
break;
}
const size_t global_work_size[3] = { num_elements, 1, 1 }; const size_t global_work_size[3] = { num_elements, 1, 1 };
const size_t local_work_size[3] = { kernel_threads, 1, 1 }; const size_t local_work_size[3] = { kernel_threads, 1, 1 };
CL_rc = hc_clEnqueueNDRangeKernel (hashcat_ctx, device_param->opencl_command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL); CL_rc = hc_clEnqueueNDRangeKernel (hashcat_ctx, device_param->opencl_command_queue, opencl_kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);
if (CL_rc == -1) return -1; if (CL_rc == -1) return -1;
@ -3387,10 +3465,12 @@ puts ("run_kernel_mp");
CL_rc = hc_clFinish (hashcat_ctx, device_param->opencl_command_queue); CL_rc = hc_clFinish (hashcat_ctx, device_param->opencl_command_queue);
if (CL_rc == -1) return -1; if (CL_rc == -1) return -1;
}
return 0; return 0;
} }
//tbd
int run_kernel_tm (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param) int run_kernel_tm (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param)
{ {
puts ("run_kernel_tm"); puts ("run_kernel_tm");
@ -3420,6 +3500,7 @@ puts ("run_kernel_tm");
return 0; return 0;
} }
//tbd
int run_kernel_amp (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const u64 num) int run_kernel_amp (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const u64 num)
{ {
puts ("run_kernel_amp"); puts ("run_kernel_amp");
@ -3459,27 +3540,43 @@ puts ("run_kernel_amp");
int run_kernel_decompress (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const u64 num) int run_kernel_decompress (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const u64 num)
{ {
puts ("run_kernel_decompress"); device_param->kernel_params_decompress_buf64[3] = num;
u64 num_elements = num;
device_param->kernel_params_decompress_buf64[3] = num_elements; u64 num_elements = num;
const u64 kernel_threads = device_param->kernel_wgs_decompress; const u64 kernel_threads = device_param->kernel_wgs_decompress;
if (device_param->is_cuda == true)
{
num_elements = CEILDIV (num_elements, kernel_threads);
CUfunction cuda_function = device_param->cuda_function_decompress;
const int rc_cuLaunchKernel = hc_cuLaunchKernel (hashcat_ctx, cuda_function, num_elements, 1, 1, kernel_threads, 1, 1, 0, NULL, device_param->kernel_params_decompress, NULL);
if (rc_cuLaunchKernel == -1) return -1;
const int rc_cuCtxSynchronize = hc_cuCtxSynchronize (hashcat_ctx);
if (rc_cuCtxSynchronize == -1) return -1;
}
if (device_param->is_opencl == true)
{
num_elements = round_up_multiple_64 (num_elements, kernel_threads); num_elements = round_up_multiple_64 (num_elements, kernel_threads);
cl_kernel kernel = device_param->opencl_kernel_decompress; cl_kernel opencl_kernel = device_param->opencl_kernel_decompress;
const size_t global_work_size[3] = { num_elements, 1, 1 }; const size_t global_work_size[3] = { num_elements, 1, 1 };
const size_t local_work_size[3] = { kernel_threads, 1, 1 }; const size_t local_work_size[3] = { kernel_threads, 1, 1 };
int CL_rc; int CL_rc;
CL_rc = hc_clSetKernelArg (hashcat_ctx, kernel, 3, sizeof (cl_ulong), device_param->kernel_params_decompress[3]); CL_rc = hc_clSetKernelArg (hashcat_ctx, opencl_kernel, 3, sizeof (cl_ulong), device_param->kernel_params_decompress[3]);
if (CL_rc == -1) return -1; if (CL_rc == -1) return -1;
CL_rc = hc_clEnqueueNDRangeKernel (hashcat_ctx, device_param->opencl_command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL); CL_rc = hc_clEnqueueNDRangeKernel (hashcat_ctx, device_param->opencl_command_queue, opencl_kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);
if (CL_rc == -1) return -1; if (CL_rc == -1) return -1;
@ -3490,13 +3587,13 @@ puts ("run_kernel_decompress");
CL_rc = hc_clFinish (hashcat_ctx, device_param->opencl_command_queue); CL_rc = hc_clFinish (hashcat_ctx, device_param->opencl_command_queue);
if (CL_rc == -1) return -1; if (CL_rc == -1) return -1;
}
return 0; return 0;
} }
int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const u64 pws_cnt) int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const u64 pws_cnt)
{ {
puts ("run_copy");
combinator_ctx_t *combinator_ctx = hashcat_ctx->combinator_ctx; combinator_ctx_t *combinator_ctx = hashcat_ctx->combinator_ctx;
hashconfig_t *hashconfig = hashcat_ctx->hashconfig; hashconfig_t *hashconfig = hashcat_ctx->hashconfig;
user_options_t *user_options = hashcat_ctx->user_options; user_options_t *user_options = hashcat_ctx->user_options;
@ -3517,6 +3614,28 @@ puts ("run_copy");
#endif #endif
if (user_options->slow_candidates == true) if (user_options->slow_candidates == true)
{
if (device_param->is_cuda == true)
{
int CU_rc;
CU_rc = hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t));
if (CU_rc == -1) return -1;
const pw_idx_t *pw_idx = device_param->pws_idx + pws_cnt;
const u32 off = pw_idx->off;
if (off)
{
CU_rc = hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32));
if (CU_rc == -1) return -1;
}
}
if (device_param->is_opencl == true)
{ {
int CL_rc; int CL_rc;
@ -3534,14 +3653,37 @@ puts ("run_copy");
if (CL_rc == -1) return -1; if (CL_rc == -1) return -1;
} }
}
CL_rc = run_kernel_decompress (hashcat_ctx, device_param, pws_cnt); const int rc_kernel_decompress = run_kernel_decompress (hashcat_ctx, device_param, pws_cnt);
if (CL_rc == -1) return -1; if (rc_kernel_decompress == -1) return -1;
} }
else else
{ {
if (user_options_extra->attack_kern == ATTACK_KERN_STRAIGHT) if (user_options_extra->attack_kern == ATTACK_KERN_STRAIGHT)
{
if (device_param->is_cuda == true)
{
int CU_rc;
CU_rc = hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t));
if (CU_rc == -1) return -1;
const pw_idx_t *pw_idx = device_param->pws_idx + pws_cnt;
const u32 off = pw_idx->off;
if (off)
{
CU_rc = hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32));
if (CU_rc == -1) return -1;
}
}
if (device_param->is_opencl == true)
{ {
int CL_rc; int CL_rc;
@ -3559,10 +3701,11 @@ puts ("run_copy");
if (CL_rc == -1) return -1; if (CL_rc == -1) return -1;
} }
}
CL_rc = run_kernel_decompress (hashcat_ctx, device_param, pws_cnt); const int rc_kernel_decompress = run_kernel_decompress (hashcat_ctx, device_param, pws_cnt);
if (CL_rc == -1) return -1; if (rc_kernel_decompress == -1) return -1;
} }
else if (user_options_extra->attack_kern == ATTACK_KERN_COMBI) else if (user_options_extra->attack_kern == ATTACK_KERN_COMBI)
{ {
@ -3602,6 +3745,28 @@ puts ("run_copy");
} }
} }
if (device_param->is_cuda == true)
{
int CU_rc;
CU_rc = hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t));
if (CU_rc == -1) return -1;
const pw_idx_t *pw_idx = device_param->pws_idx + pws_cnt;
const u32 off = pw_idx->off;
if (off)
{
CU_rc = hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32));
if (CU_rc == -1) return -1;
}
}
if (device_param->is_opencl == true)
{
int CL_rc; int CL_rc;
CL_rc = hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_idx, CL_TRUE, 0, pws_cnt * sizeof (pw_idx_t), device_param->pws_idx, 0, NULL, NULL); CL_rc = hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_idx, CL_TRUE, 0, pws_cnt * sizeof (pw_idx_t), device_param->pws_idx, 0, NULL, NULL);
@ -3618,14 +3783,37 @@ puts ("run_copy");
if (CL_rc == -1) return -1; if (CL_rc == -1) return -1;
} }
}
CL_rc = run_kernel_decompress (hashcat_ctx, device_param, pws_cnt); const int rc_kernel_decompress = run_kernel_decompress (hashcat_ctx, device_param, pws_cnt);
if (CL_rc == -1) return -1; if (rc_kernel_decompress == -1) return -1;
} }
else else
{ {
if (user_options->attack_mode == ATTACK_MODE_COMBI) if (user_options->attack_mode == ATTACK_MODE_COMBI)
{
if (device_param->is_cuda == true)
{
int CU_rc;
CU_rc = hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t));
if (CU_rc == -1) return -1;
const pw_idx_t *pw_idx = device_param->pws_idx + pws_cnt;
const u32 off = pw_idx->off;
if (off)
{
CU_rc = hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32));
if (CU_rc == -1) return -1;
}
}
if (device_param->is_opencl == true)
{ {
int CL_rc; int CL_rc;
@ -3643,12 +3831,35 @@ puts ("run_copy");
if (CL_rc == -1) return -1; if (CL_rc == -1) return -1;
} }
}
CL_rc = run_kernel_decompress (hashcat_ctx, device_param, pws_cnt); const int rc_kernel_decompress = run_kernel_decompress (hashcat_ctx, device_param, pws_cnt);
if (CL_rc == -1) return -1; if (rc_kernel_decompress == -1) return -1;
} }
else if (user_options->attack_mode == ATTACK_MODE_HYBRID1) else if (user_options->attack_mode == ATTACK_MODE_HYBRID1)
{
if (device_param->is_cuda == true)
{
int CU_rc;
CU_rc = hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t));
if (CU_rc == -1) return -1;
const pw_idx_t *pw_idx = device_param->pws_idx + pws_cnt;
const u32 off = pw_idx->off;
if (off)
{
CU_rc = hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32));
if (CU_rc == -1) return -1;
}
}
if (device_param->is_opencl == true)
{ {
int CL_rc; int CL_rc;
@ -3666,10 +3877,11 @@ puts ("run_copy");
if (CL_rc == -1) return -1; if (CL_rc == -1) return -1;
} }
}
CL_rc = run_kernel_decompress (hashcat_ctx, device_param, pws_cnt); const int rc_kernel_decompress = run_kernel_decompress (hashcat_ctx, device_param, pws_cnt);
if (CL_rc == -1) return -1; if (rc_kernel_decompress == -1) return -1;
} }
else if (user_options->attack_mode == ATTACK_MODE_HYBRID2) else if (user_options->attack_mode == ATTACK_MODE_HYBRID2)
{ {
@ -3700,7 +3912,6 @@ puts ("run_copy");
int run_cracker (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const u64 pws_cnt) int run_cracker (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const u64 pws_cnt)
{ {
puts ("run_cracker");
combinator_ctx_t *combinator_ctx = hashcat_ctx->combinator_ctx; combinator_ctx_t *combinator_ctx = hashcat_ctx->combinator_ctx;
hashconfig_t *hashconfig = hashcat_ctx->hashconfig; hashconfig_t *hashconfig = hashcat_ctx->hashconfig;
hashes_t *hashes = hashcat_ctx->hashes; hashes_t *hashes = hashcat_ctx->hashes;
@ -3861,11 +4072,21 @@ puts ("run_cracker");
else else
{ {
if (user_options_extra->attack_kern == ATTACK_KERN_STRAIGHT) if (user_options_extra->attack_kern == ATTACK_KERN_STRAIGHT)
{
if (device_param->is_cuda == true)
{
const int CU_rc = hc_cuMemcpyDtoD (hashcat_ctx, device_param->cuda_d_rules_c, device_param->cuda_d_rules + (innerloop_pos * sizeof (kernel_rule_t)), innerloop_left * sizeof (kernel_rule_t));
if (CU_rc == -1) return -1;
}
if (device_param->is_opencl == true)
{ {
const int CL_rc = hc_clEnqueueCopyBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_rules, device_param->opencl_d_rules_c, innerloop_pos * sizeof (kernel_rule_t), 0, innerloop_left * sizeof (kernel_rule_t), 0, NULL, NULL); const int CL_rc = hc_clEnqueueCopyBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_rules, device_param->opencl_d_rules_c, innerloop_pos * sizeof (kernel_rule_t), 0, innerloop_left * sizeof (kernel_rule_t), 0, NULL, NULL);
if (CL_rc == -1) return -1; if (CL_rc == -1) return -1;
} }
}
else if (user_options_extra->attack_kern == ATTACK_KERN_COMBI) else if (user_options_extra->attack_kern == ATTACK_KERN_COMBI)
{ {
if (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) if (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL)
@ -3968,10 +4189,20 @@ puts ("run_cracker");
innerloop_left = i; innerloop_left = i;
if (device_param->is_cuda == true)
{
const int CU_rc = hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_combs_c, device_param->combs_buf, innerloop_left * sizeof (pw_t));
if (CU_rc == -1) return -1;
}
if (device_param->is_opencl == true)
{
const int CL_rc = hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_combs_c, CL_TRUE, 0, innerloop_left * sizeof (pw_t), device_param->combs_buf, 0, NULL, NULL); const int CL_rc = hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_combs_c, CL_TRUE, 0, innerloop_left * sizeof (pw_t), device_param->combs_buf, 0, NULL, NULL);
if (CL_rc == -1) return -1; if (CL_rc == -1) return -1;
} }
}
else if (user_options->attack_mode == ATTACK_MODE_HYBRID1) else if (user_options->attack_mode == ATTACK_MODE_HYBRID1)
{ {
u64 off = innerloop_pos; u64 off = innerloop_pos;
@ -3984,27 +4215,45 @@ puts ("run_cracker");
if (CL_rc == -1) return -1; if (CL_rc == -1) return -1;
if (device_param->is_cuda == true)
{
const int CU_rc = hc_cuMemcpyDtoD (hashcat_ctx, device_param->cuda_d_combs_c, device_param->cuda_d_combs, innerloop_left * sizeof (pw_t));
if (CU_rc == -1) return -1;
}
if (device_param->is_opencl == true)
{
CL_rc = hc_clEnqueueCopyBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_combs, device_param->opencl_d_combs_c, 0, 0, innerloop_left * sizeof (pw_t), 0, NULL, NULL); CL_rc = hc_clEnqueueCopyBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_combs, device_param->opencl_d_combs_c, 0, 0, innerloop_left * sizeof (pw_t), 0, NULL, NULL);
if (CL_rc == -1) return -1; if (CL_rc == -1) return -1;
} }
}
else if (user_options->attack_mode == ATTACK_MODE_HYBRID2) else if (user_options->attack_mode == ATTACK_MODE_HYBRID2)
{ {
u64 off = innerloop_pos; u64 off = innerloop_pos;
device_param->kernel_params_mp_buf64[3] = off; device_param->kernel_params_mp_buf64[3] = off;
int CL_rc; const int rc_kernel_mp = run_kernel_mp (hashcat_ctx, device_param, KERN_RUN_MP, innerloop_left);
CL_rc = run_kernel_mp (hashcat_ctx, device_param, KERN_RUN_MP, innerloop_left); if (rc_kernel_mp == -1) return -1;
if (CL_rc == -1) return -1; if (device_param->is_cuda == true)
{
CL_rc = hc_clEnqueueCopyBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_combs, device_param->opencl_d_combs_c, 0, 0, innerloop_left * sizeof (pw_t), 0, NULL, NULL); const int CU_rc = hc_cuMemcpyDtoD (hashcat_ctx, device_param->cuda_d_combs_c, device_param->cuda_d_combs, innerloop_left * sizeof (pw_t));
if (CU_rc == -1) return -1;
}
if (device_param->is_opencl == true)
{
const int CL_rc = hc_clEnqueueCopyBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_combs, device_param->opencl_d_combs_c, 0, 0, innerloop_left * sizeof (pw_t), 0, NULL, NULL);
if (CL_rc == -1) return -1; if (CL_rc == -1) return -1;
} }
} }
}
else else
{ {
if ((user_options->attack_mode == ATTACK_MODE_COMBI) || (user_options->attack_mode == ATTACK_MODE_HYBRID2)) if ((user_options->attack_mode == ATTACK_MODE_COMBI) || (user_options->attack_mode == ATTACK_MODE_HYBRID2))
@ -4107,45 +4356,71 @@ puts ("run_cracker");
innerloop_left = i; innerloop_left = i;
if (device_param->is_cuda == true)
{
const int CU_rc = hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_combs_c, device_param->combs_buf, innerloop_left * sizeof (pw_t));
if (CU_rc == -1) return -1;
}
if (device_param->is_opencl == true)
{
const int CL_rc = hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_combs_c, CL_TRUE, 0, innerloop_left * sizeof (pw_t), device_param->combs_buf, 0, NULL, NULL); const int CL_rc = hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_combs_c, CL_TRUE, 0, innerloop_left * sizeof (pw_t), device_param->combs_buf, 0, NULL, NULL);
if (CL_rc == -1) return -1; if (CL_rc == -1) return -1;
} }
}
else if (user_options->attack_mode == ATTACK_MODE_HYBRID1) else if (user_options->attack_mode == ATTACK_MODE_HYBRID1)
{ {
u64 off = innerloop_pos; u64 off = innerloop_pos;
device_param->kernel_params_mp_buf64[3] = off; device_param->kernel_params_mp_buf64[3] = off;
int CL_rc; const int rc_kernel_mp = run_kernel_mp (hashcat_ctx, device_param, KERN_RUN_MP, innerloop_left);
CL_rc = run_kernel_mp (hashcat_ctx, device_param, KERN_RUN_MP, innerloop_left); if (rc_kernel_mp == -1) return -1;
if (CL_rc == -1) return -1; if (device_param->is_cuda == true)
{
CL_rc = hc_clEnqueueCopyBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_combs, device_param->opencl_d_combs_c, 0, 0, innerloop_left * sizeof (pw_t), 0, NULL, NULL); const int CU_rc = hc_cuMemcpyDtoD (hashcat_ctx, device_param->cuda_d_combs_c, device_param->cuda_d_combs, innerloop_left * sizeof (pw_t));
if (CU_rc == -1) return -1;
}
if (device_param->is_opencl == true)
{
const int CL_rc = hc_clEnqueueCopyBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_combs, device_param->opencl_d_combs_c, 0, 0, innerloop_left * sizeof (pw_t), 0, NULL, NULL);
if (CL_rc == -1) return -1; if (CL_rc == -1) return -1;
} }
} }
} }
}
else if (user_options_extra->attack_kern == ATTACK_KERN_BF) else if (user_options_extra->attack_kern == ATTACK_KERN_BF)
{ {
u64 off = innerloop_pos; u64 off = innerloop_pos;
device_param->kernel_params_mp_r_buf64[3] = off; device_param->kernel_params_mp_r_buf64[3] = off;
int CL_rc; const int rc_kernel_mp = run_kernel_mp (hashcat_ctx, device_param, KERN_RUN_MP_R, innerloop_left);
CL_rc = run_kernel_mp (hashcat_ctx, device_param, KERN_RUN_MP_R, innerloop_left); if (rc_kernel_mp == -1) return -1;
if (CL_rc == -1) return -1; if (device_param->is_cuda == true)
{
CL_rc = hc_clEnqueueCopyBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_bfs, device_param->opencl_d_bfs_c, 0, 0, innerloop_left * sizeof (bf_t), 0, NULL, NULL); const int CU_rc = hc_cuMemcpyDtoD (hashcat_ctx, device_param->cuda_d_bfs_c, device_param->cuda_d_bfs, innerloop_left * sizeof (bf_t));
if (CU_rc == -1) return -1;
}
if (device_param->is_opencl == true)
{
const int CL_rc = hc_clEnqueueCopyBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_bfs, device_param->opencl_d_bfs_c, 0, 0, innerloop_left * sizeof (bf_t), 0, NULL, NULL);
if (CL_rc == -1) return -1; if (CL_rc == -1) return -1;
} }
} }
}
const int rc = choose_kernel (hashcat_ctx, device_param, highest_pw_len, pws_cnt, fast_iteration, salt_pos); const int rc = choose_kernel (hashcat_ctx, device_param, highest_pw_len, pws_cnt, fast_iteration, salt_pos);
@ -10562,9 +10837,20 @@ int backend_session_update_mp (hashcat_ctx_t *hashcat_ctx)
device_param->kernel_params_mp_buf64[3] = 0; device_param->kernel_params_mp_buf64[3] = 0;
device_param->kernel_params_mp_buf32[4] = mask_ctx->css_cnt; device_param->kernel_params_mp_buf32[4] = mask_ctx->css_cnt;
if (device_param->is_cuda == true)
{
int CU_rc;
//for (u32 i = 3; i < 4; i++) { CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_mp, i, sizeof (cl_ulong), device_param->kernel_params_mp[i]); if (CL_rc == -1) return -1; }
//for (u32 i = 4; i < 8; i++) { CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_mp, i, sizeof (cl_uint), device_param->kernel_params_mp[i]); if (CL_rc == -1) return -1; }
CU_rc = hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_root_css_buf, mask_ctx->root_css_buf, device_param->size_root_css); if (CU_rc == -1) return -1;
CU_rc = hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_markov_css_buf, mask_ctx->markov_css_buf, device_param->size_markov_css); if (CU_rc == -1) return -1;
}
if (device_param->is_opencl == true) if (device_param->is_opencl == true)
{ {
int CL_rc = CL_SUCCESS; int CL_rc;
for (u32 i = 3; i < 4; i++) { CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_mp, i, sizeof (cl_ulong), device_param->kernel_params_mp[i]); if (CL_rc == -1) return -1; } for (u32 i = 3; i < 4; i++) { CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_mp, i, sizeof (cl_ulong), device_param->kernel_params_mp[i]); if (CL_rc == -1) return -1; }
for (u32 i = 4; i < 8; i++) { CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_mp, i, sizeof (cl_uint), device_param->kernel_params_mp[i]); if (CL_rc == -1) return -1; } for (u32 i = 4; i < 8; i++) { CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_mp, i, sizeof (cl_uint), device_param->kernel_params_mp[i]); if (CL_rc == -1) return -1; }
@ -10602,6 +10888,22 @@ int backend_session_update_mp_rl (hashcat_ctx_t *hashcat_ctx, const u32 css_cnt_
device_param->kernel_params_mp_r_buf64[3] = 0; device_param->kernel_params_mp_r_buf64[3] = 0;
device_param->kernel_params_mp_r_buf32[4] = css_cnt_r; device_param->kernel_params_mp_r_buf32[4] = css_cnt_r;
if (device_param->is_cuda == true)
{
int CU_rc;
//for (u32 i = 3; i < 4; i++) { CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_mp_l, i, sizeof (cl_ulong), device_param->kernel_params_mp_l[i]); if (CL_rc == -1) return -1; }
//for (u32 i = 4; i < 8; i++) { CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_mp_l, i, sizeof (cl_uint), device_param->kernel_params_mp_l[i]); if (CL_rc == -1) return -1; }
//for (u32 i = 9; i < 9; i++) { CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_mp_l, i, sizeof (cl_ulong), device_param->kernel_params_mp_l[i]); if (CL_rc == -1) return -1; }
//for (u32 i = 3; i < 4; i++) { CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_mp_r, i, sizeof (cl_ulong), device_param->kernel_params_mp_r[i]); if (CL_rc == -1) return -1; }
//for (u32 i = 4; i < 7; i++) { CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_mp_r, i, sizeof (cl_uint), device_param->kernel_params_mp_r[i]); if (CL_rc == -1) return -1; }
//for (u32 i = 8; i < 8; i++) { CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_mp_r, i, sizeof (cl_ulong), device_param->kernel_params_mp_r[i]); if (CL_rc == -1) return -1; }
CU_rc = hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_root_css_buf, mask_ctx->root_css_buf, device_param->size_root_css); if (CU_rc == -1) return -1;
CU_rc = hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_markov_css_buf, mask_ctx->markov_css_buf, device_param->size_markov_css); if (CU_rc == -1) return -1;
}
if (device_param->is_opencl == true) if (device_param->is_opencl == true)
{ {
int CL_rc = CL_SUCCESS; int CL_rc = CL_SUCCESS;

14
src/dispatch.c
View File

@ -349,6 +349,13 @@ HC_API_CALL void *thread_calc_stdin (void *p)
if (device_param->skipped_warning == true) return NULL; if (device_param->skipped_warning == true) return NULL;
if (device_param->is_cuda == true)
{
const int rc_cuCtxSetCurrent = hc_cuCtxSetCurrent (hashcat_ctx, device_param->cuda_context);
if (rc_cuCtxSetCurrent == -1) return NULL;
}
const int rc_calc = calc_stdin (hashcat_ctx, device_param); const int rc_calc = calc_stdin (hashcat_ctx, device_param);
if (rc_calc == -1) if (rc_calc == -1)
@ -1668,6 +1675,13 @@ HC_API_CALL void *thread_calc (void *p)
if (device_param->skipped_warning == true) return NULL; if (device_param->skipped_warning == true) return NULL;
if (device_param->is_cuda == true)
{
const int rc_cuCtxSetCurrent = hc_cuCtxSetCurrent (hashcat_ctx, device_param->cuda_context);
if (rc_cuCtxSetCurrent == -1) return NULL;
}
const int rc_calc = calc (hashcat_ctx, device_param); const int rc_calc = calc (hashcat_ctx, device_param);
if (rc_calc == -1) if (rc_calc == -1)

71
src/hashes.c
View File

@ -309,8 +309,16 @@ void check_hash (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, pl
{ {
tmps = hcmalloc (hashconfig->tmp_size); tmps = hcmalloc (hashconfig->tmp_size);
if (device_param->is_cuda == true)
{
hc_cuMemcpyDtoH (hashcat_ctx, tmps, device_param->cuda_d_tmps + (plain->gidvid * hashconfig->tmp_size), hashconfig->tmp_size);
}
if (device_param->is_opencl == true)
{
hc_clEnqueueReadBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_tmps, CL_TRUE, plain->gidvid * hashconfig->tmp_size, hashconfig->tmp_size, tmps, 0, NULL, NULL); hc_clEnqueueReadBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_tmps, CL_TRUE, plain->gidvid * hashconfig->tmp_size, hashconfig->tmp_size, tmps, 0, NULL, NULL);
} }
}
// hash // hash
@ -460,15 +468,21 @@ int check_cracked (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param,
u32 num_cracked; u32 num_cracked;
cl_int CL_err; int CU_rc;
int CL_rc;
CL_err = hc_clEnqueueReadBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_result, CL_TRUE, 0, sizeof (u32), &num_cracked, 0, NULL, NULL); if (device_param->is_cuda == true)
if (CL_err != CL_SUCCESS)
{ {
event_log_error (hashcat_ctx, "clEnqueueReadBuffer(): %s", val2cstr_cl (CL_err)); CU_rc = hc_cuMemcpyDtoH (hashcat_ctx, &num_cracked, device_param->cuda_d_result, sizeof (u32));
return -1; if (CU_rc == -1) return -1;
}
if (device_param->is_opencl == true)
{
CL_rc = hc_clEnqueueReadBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_result, CL_TRUE, 0, sizeof (u32), &num_cracked, 0, NULL, NULL);
if (CL_rc == -1) return -1;
} }
if (user_options->speed_only == true) if (user_options->speed_only == true)
@ -483,13 +497,18 @@ int check_cracked (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param,
{ {
plain_t *cracked = (plain_t *) hccalloc (num_cracked, sizeof (plain_t)); plain_t *cracked = (plain_t *) hccalloc (num_cracked, sizeof (plain_t));
CL_err = hc_clEnqueueReadBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_plain_bufs, CL_TRUE, 0, num_cracked * sizeof (plain_t), cracked, 0, NULL, NULL); if (device_param->is_cuda == true)
if (CL_err != CL_SUCCESS)
{ {
event_log_error (hashcat_ctx, "clEnqueueReadBuffer(): %s", val2cstr_cl (CL_err)); CU_rc = hc_cuMemcpyDtoH (hashcat_ctx, cracked, device_param->cuda_d_plain_bufs, num_cracked * sizeof (plain_t));
return -1; if (CU_rc == -1) return -1;
}
if (device_param->is_opencl == true)
{
CL_rc = hc_clEnqueueReadBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_plain_bufs, CL_TRUE, 0, num_cracked * sizeof (plain_t), cracked, 0, NULL, NULL);
if (CL_rc == -1) return -1;
} }
u32 cpt_cracked = 0; u32 cpt_cracked = 0;
@ -553,25 +572,35 @@ int check_cracked (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param,
memset (hashes->digests_shown_tmp, 0, salt_buf->digests_cnt * sizeof (u32)); memset (hashes->digests_shown_tmp, 0, salt_buf->digests_cnt * sizeof (u32));
CL_err = hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_digests_shown, CL_TRUE, salt_buf->digests_offset * sizeof (u32), salt_buf->digests_cnt * sizeof (u32), &hashes->digests_shown_tmp[salt_buf->digests_offset], 0, NULL, NULL); if (device_param->is_cuda == true)
if (CL_err != CL_SUCCESS)
{ {
event_log_error (hashcat_ctx, "clEnqueueWriteBuffer(): %s", val2cstr_cl (CL_err)); CU_rc = hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_digests_shown + (salt_buf->digests_offset * sizeof (u32)), &hashes->digests_shown_tmp[salt_buf->digests_offset], salt_buf->digests_cnt * sizeof (u32));
return -1; if (CU_rc == -1) return -1;
}
if (device_param->is_opencl == true)
{
CL_rc = hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_digests_shown, CL_TRUE, salt_buf->digests_offset * sizeof (u32), salt_buf->digests_cnt * sizeof (u32), &hashes->digests_shown_tmp[salt_buf->digests_offset], 0, NULL, NULL);
if (CL_rc == -1) return -1;
} }
} }
num_cracked = 0; num_cracked = 0;
CL_err = hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_result, CL_TRUE, 0, sizeof (u32), &num_cracked, 0, NULL, NULL); if (device_param->is_cuda == true)
if (CL_err != CL_SUCCESS)
{ {
event_log_error (hashcat_ctx, "clEnqueueWriteBuffer(): %s", val2cstr_cl (CL_err)); CU_rc = hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_result, &num_cracked, sizeof (u32));
return -1; if (CU_rc == -1) return -1;
}
if (device_param->is_opencl == true)
{
CL_rc = hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_result, CL_TRUE, 0, sizeof (u32), &num_cracked, 0, NULL, NULL);
if (CL_rc == -1) return -1;
} }
} }