diff --git a/include/ext_OpenCL.h b/include/ext_OpenCL.h index 681751a5b..5b730ad3d 100644 --- a/include/ext_OpenCL.h +++ b/include/ext_OpenCL.h @@ -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_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); -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); 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); diff --git a/include/types.h b/include/types.h index fd7449fb8..2fc417965 100644 --- a/include/types.h +++ b/include/types.h @@ -834,7 +834,7 @@ struct __hc_device_param uint device_id; uint platform_devices_id; // for mapping with hms devices - uint skipped; + bool skipped; uint sm_major; uint sm_minor; diff --git a/src/ext_OpenCL.c b/src/ext_OpenCL.c index 4e6d058df..f5fa4316f 100644 --- a/src/ext_OpenCL.c +++ b/src/ext_OpenCL.c @@ -376,33 +376,54 @@ cl_program hc_clCreateProgramWithBinary (OCL_PTR *ocl, cl_context context, cl_ui 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); 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; - size_t len = 0; + char *options_update = (char *) mymalloc (len + 1); - 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) diff --git a/src/oclHashcat.c b/src/oclHashcat.c index 17cfdd4c7..fd6c6e11d 100644 --- a/src/oclHashcat.c +++ b/src/oclHashcat.c @@ -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); - 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; @@ -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); - 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 @@ -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); } - 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); @@ -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); - 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; @@ -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); - 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); @@ -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); - 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; @@ -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); - 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);