skip device if buildProgram() fail

diff --git a/include/ext_OpenCL.h b/include/ext_OpenCL.h
index 681751a..5b730ad 100644 (file)
--- 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 fd7449f..2fc4179 100644 (file)
--- 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 4e6d058..f5fa431 100644 (file)
--- 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));
+
+      log_error ("\n=== Build Options : %s ===\n", options);
 
-      if (hc_clGetProgramBuildInfo (ocl, program, *device_list, CL_PROGRAM_BUILD_LOG, len, buf, NULL) == CL_SUCCESS)
+      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 17cfdd4..fd6c6e1 100644 (file)
--- 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);