arno/hashcat
1
0
Fork 0
mirror of https://github.com/hashcat/hashcat.git synced 2024-12-22 22:58:30 +00:00
Code Issues Releases Wiki Activity

skip device if buildProgram() fail

Browse Source
This commit is contained in:
Gabriele 'matrix' Gristina 2016-02-07 21:20:10 +01:00
parent 6d37c123a5
commit 74be8e2e25
4 changed files with 69 additions and 21 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
include/ext_OpenCL.h
View File

@ -109,7 +109,7 @@ cl_context hc_clCreateContext (OCL_PTR *ocl, cl_context_properties *properties,
cl_kernel hc_clCreateKernel (OCL_PTR *ocl, cl_program program, const char *kernel_name); cl_kernel hc_clCreateKernel (OCL_PTR *ocl, cl_program program, const char *kernel_name);
cl_program hc_clCreateProgramWithSource (OCL_PTR *ocl, cl_context context, cl_uint count, const char **strings, const size_t *lengths); cl_program hc_clCreateProgramWithSource (OCL_PTR *ocl, cl_context context, cl_uint count, const char **strings, const size_t *lengths);
cl_program hc_clCreateProgramWithBinary (OCL_PTR *ocl, cl_context context, cl_uint num_devices, const cl_device_id *device_list, const size_t *lengths, const unsigned char **binaries, cl_int *binary_status); cl_program hc_clCreateProgramWithBinary (OCL_PTR *ocl, cl_context context, cl_uint num_devices, const cl_device_id *device_list, const size_t *lengths, const unsigned char **binaries, cl_int *binary_status);
void hc_clBuildProgram (OCL_PTR *ocl, cl_program program, cl_uint num_devices, const cl_device_id *device_list, const char *options, void (CL_CALLBACK *pfn_notify) (cl_program program, void *user_data), void *user_data); cl_int hc_clBuildProgram (OCL_PTR *ocl, cl_program program, cl_uint num_devices, const cl_device_id *device_list, const char *options, void (CL_CALLBACK *pfn_notify) (cl_program program, void *user_data), void *user_data, bool exitOnFail);
cl_int hc_clEnqueueNDRangeKernel (OCL_PTR *ocl, cl_command_queue command_queue, cl_kernel kernel, cl_uint work_dim, const size_t *global_work_offset, const size_t *global_work_size, const size_t *local_work_size, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event, bool exitOnFail); cl_int hc_clEnqueueNDRangeKernel (OCL_PTR *ocl, cl_command_queue command_queue, cl_kernel kernel, cl_uint work_dim, const size_t *global_work_offset, const size_t *global_work_size, const size_t *local_work_size, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event, bool exitOnFail);
void hc_clEnqueueReadBuffer (OCL_PTR *ocl, cl_command_queue command_queue, cl_mem buffer, cl_bool blocking_read, size_t offset, size_t cb, void *ptr, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event); void hc_clEnqueueReadBuffer (OCL_PTR *ocl, cl_command_queue command_queue, cl_mem buffer, cl_bool blocking_read, size_t offset, size_t cb, void *ptr, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event);
void hc_clEnqueueWriteBuffer (OCL_PTR *ocl, cl_command_queue command_queue, cl_mem buffer, cl_bool blocking_write, size_t offset, size_t cb, const void *ptr, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event); void hc_clEnqueueWriteBuffer (OCL_PTR *ocl, cl_command_queue command_queue, cl_mem buffer, cl_bool blocking_write, size_t offset, size_t cb, const void *ptr, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event);

2
include/types.h
View File

@ -834,7 +834,7 @@ struct __hc_device_param
uint device_id; uint device_id;
uint platform_devices_id; // for mapping with hms devices uint platform_devices_id; // for mapping with hms devices
uint skipped; bool skipped;
uint sm_major; uint sm_major;
uint sm_minor; uint sm_minor;

45
src/ext_OpenCL.c
View File

@ -376,33 +376,54 @@ cl_program hc_clCreateProgramWithBinary (OCL_PTR *ocl, cl_context context, cl_ui
return (program); return (program);
} }
void hc_clBuildProgram (OCL_PTR *ocl, cl_program program, cl_uint num_devices, const cl_device_id *device_list, const char *options, void (CL_CALLBACK *pfn_notify) (cl_program program, void *user_data), void *user_data) cl_int hc_clBuildProgram (OCL_PTR *ocl, cl_program program, cl_uint num_devices, const cl_device_id *device_list, const char *options, void (CL_CALLBACK *pfn_notify) (cl_program program, void *user_data), void *user_data, bool exitOnFail)
{ {
cl_int CL_err = ocl->clBuildProgram (program, num_devices, device_list, options, pfn_notify, user_data); cl_int CL_err = ocl->clBuildProgram (program, num_devices, device_list, options, pfn_notify, user_data);
if (CL_err != CL_SUCCESS) if (CL_err != CL_SUCCESS)
{ {
log_error ("ERROR: %s : %d : %s\n", "clBuildProgram()", CL_err, val2cstr_cl (CL_err)); size_t len = strlen (options) + 1 + 15;
char *buf = NULL; char *options_update = (char *) mymalloc (len + 1);
size_t len = 0;
cl_int err = hc_clGetProgramBuildInfo (ocl, program, *device_list, CL_PROGRAM_BUILD_LOG, 0, NULL, &len); options_update = strncat (options_update, options, len - 1 - 15);
options_update = strncat (options_update, " -cl-opt-disable", 1 + 15);
if (err == CL_SUCCESS && len > 0) if (data.quiet == 0) log_error ("\n=== Build failed, retry with optimization disabled ===\n");
CL_err = ocl->clBuildProgram (program, num_devices, device_list, options_update, pfn_notify, user_data);
myfree (options_update);
if (CL_err != CL_SUCCESS)
{ {
buf = (char *) mymalloc (len + 1); log_error ("ERROR: %s : %d : %s\n", "clBuildProgram()", CL_err, val2cstr_cl (CL_err));
if (hc_clGetProgramBuildInfo (ocl, program, *device_list, CL_PROGRAM_BUILD_LOG, len, buf, NULL) == CL_SUCCESS) log_error ("\n=== Build Options : %s ===\n", options);
size_t len = 0;
cl_int err = hc_clGetProgramBuildInfo (ocl, program, *device_list, CL_PROGRAM_BUILD_LOG, 0, NULL, &len);
if (err == CL_SUCCESS && len > 0)
{ {
fprintf (stderr, "\n=== Build Log (start) ===\n%s\n=== Build Log (end) ===\n", buf); char *buf = (char *) mymalloc (len + 1);
if (hc_clGetProgramBuildInfo (ocl, program, *device_list, CL_PROGRAM_BUILD_LOG, len, buf, NULL) == CL_SUCCESS)
{
fprintf (stderr, "\n=== Build Log (start) ===\n%s\n=== Build Log (end) ===\n", buf);
}
myfree (buf);
} }
myfree (buf); if (exitOnFail) exit (-1);
}
exit (-1); return (-1);
}
} }
return 0;
} }
cl_kernel hc_clCreateKernel (OCL_PTR *ocl, cl_program program, const char *kernel_name) cl_kernel hc_clCreateKernel (OCL_PTR *ocl, cl_program program, const char *kernel_name)

41
src/oclHashcat.c
View File

@ -13775,7 +13775,14 @@ int main (int argc, char **argv)
device_param->program = hc_clCreateProgramWithSource (data.ocl, device_param->context, 1, (const char **) kernel_sources, NULL); device_param->program = hc_clCreateProgramWithSource (data.ocl, device_param->context, 1, (const char **) kernel_sources, NULL);
hc_clBuildProgram (data.ocl, device_param->program, 1, &device_param->device, build_opts, NULL, NULL); int rc = hc_clBuildProgram (data.ocl, device_param->program, 1, &device_param->device, build_opts, NULL, NULL, false);
if (rc != 0)
{
device_param->skipped = true;
log_info ("Device #%u: Kernel %s build failure. Proceed without this device.", device_id + 1, source_file);
continue;
}
size_t binary_size; size_t binary_size;
@ -13797,7 +13804,7 @@ int main (int argc, char **argv)
device_param->program = hc_clCreateProgramWithBinary (data.ocl, device_param->context, 1, &device_param->device, kernel_lengths, (const u8 **) kernel_sources, NULL); device_param->program = hc_clCreateProgramWithBinary (data.ocl, device_param->context, 1, &device_param->device, kernel_lengths, (const u8 **) kernel_sources, NULL);
hc_clBuildProgram (data.ocl, device_param->program, 1, &device_param->device, build_opts, NULL, NULL); hc_clBuildProgram (data.ocl, device_param->program, 1, &device_param->device, build_opts, NULL, NULL, true);
} }
} }
else else
@ -13819,7 +13826,13 @@ int main (int argc, char **argv)
snprintf (build_opts_update, sizeof (build_opts_update) - 1, "%s -DSCRYPT_N=%d -DSCRYPT_R=%d -DSCRYPT_P=%d -DSCRYPT_TMTO=%d", build_opts, data.salts_buf[0].scrypt_N, data.salts_buf[0].scrypt_r, data.salts_buf[0].scrypt_p, 1 << data.salts_buf[0].scrypt_tmto); snprintf (build_opts_update, sizeof (build_opts_update) - 1, "%s -DSCRYPT_N=%d -DSCRYPT_R=%d -DSCRYPT_P=%d -DSCRYPT_TMTO=%d", build_opts, data.salts_buf[0].scrypt_N, data.salts_buf[0].scrypt_r, data.salts_buf[0].scrypt_p, 1 << data.salts_buf[0].scrypt_tmto);
} }
hc_clBuildProgram (data.ocl, device_param->program, 1, &device_param->device, build_opts_update, NULL, NULL); int rc = hc_clBuildProgram (data.ocl, device_param->program, 1, &device_param->device, build_opts_update, NULL, NULL, false);
if (rc != 0)
{
device_param->skipped = true;
log_info ("Device #%u: Kernel %s build failure. Proceed without this device.", device_id + 1, source_file);
}
} }
local_free (kernel_lengths); local_free (kernel_lengths);
@ -13883,7 +13896,14 @@ int main (int argc, char **argv)
device_param->program_mp = hc_clCreateProgramWithSource (data.ocl, device_param->context, 1, (const char **) kernel_sources, NULL); device_param->program_mp = hc_clCreateProgramWithSource (data.ocl, device_param->context, 1, (const char **) kernel_sources, NULL);
hc_clBuildProgram (data.ocl, device_param->program_mp, 1, &device_param->device, build_opts, NULL, NULL); int rc = hc_clBuildProgram (data.ocl, device_param->program_mp, 1, &device_param->device, build_opts, NULL, NULL, false);
if (rc != 0)
{
device_param->skipped = true;
log_info ("Device #%u: Kernel %s build failure. Proceed without this device.", device_id + 1, source_file);
continue;
}
size_t binary_size; size_t binary_size;
@ -13905,7 +13925,7 @@ int main (int argc, char **argv)
device_param->program_mp = hc_clCreateProgramWithBinary (data.ocl, device_param->context, 1, &device_param->device, kernel_lengths, (const u8 **) kernel_sources, NULL); device_param->program_mp = hc_clCreateProgramWithBinary (data.ocl, device_param->context, 1, &device_param->device, kernel_lengths, (const u8 **) kernel_sources, NULL);
hc_clBuildProgram (data.ocl, device_param->program_mp, 1, &device_param->device, build_opts, NULL, NULL); hc_clBuildProgram (data.ocl, device_param->program_mp, 1, &device_param->device, build_opts, NULL, NULL, true);
} }
local_free (kernel_lengths); local_free (kernel_lengths);
@ -13973,7 +13993,14 @@ int main (int argc, char **argv)
device_param->program_amp = hc_clCreateProgramWithSource (data.ocl, device_param->context, 1, (const char **) kernel_sources, NULL); device_param->program_amp = hc_clCreateProgramWithSource (data.ocl, device_param->context, 1, (const char **) kernel_sources, NULL);
hc_clBuildProgram (data.ocl, device_param->program_amp, 1, &device_param->device, build_opts, NULL, NULL); int rc = hc_clBuildProgram (data.ocl, device_param->program_amp, 1, &device_param->device, build_opts, NULL, NULL, false);
if (rc != 0)
{
device_param->skipped = true;
log_info ("Device #%u: Kernel %s build failure. Proceed without this device.", device_id + 1, source_file);
continue;
}
size_t binary_size; size_t binary_size;
@ -13995,7 +14022,7 @@ int main (int argc, char **argv)
device_param->program_amp = hc_clCreateProgramWithBinary (data.ocl, device_param->context, 1, &device_param->device, kernel_lengths, (const u8 **) kernel_sources, NULL); device_param->program_amp = hc_clCreateProgramWithBinary (data.ocl, device_param->context, 1, &device_param->device, kernel_lengths, (const u8 **) kernel_sources, NULL);
hc_clBuildProgram (data.ocl, device_param->program_amp, 1, &device_param->device, build_opts, NULL, NULL); hc_clBuildProgram (data.ocl, device_param->program_amp, 1, &device_param->device, build_opts, NULL, NULL, true);
} }
local_free (kernel_lengths); local_free (kernel_lengths);