X-Git-Url: https://www.flypig.org.uk/git/?a=blobdiff_plain;f=src%2Fhashcat.c;h=057663d28d4bb5653a7d5ee3fb9f7bd80964c39a;hb=9b4e6056d452fa9429112eafaf7cced19e63f09c;hp=f3a97a9b6473a473d3bb1e65fa3e79bdab0dad39;hpb=8fc1306b0b16950a97e5e9bde56577db180337ab;p=hashcat.git diff --git a/src/hashcat.c b/src/hashcat.c index f3a97a9..057663d 100644 --- a/src/hashcat.c +++ b/src/hashcat.c @@ -6,7 +6,11 @@ * License.....: MIT */ -#ifdef DARWIN +#ifdef __APPLE__ +#include +#endif + +#ifdef __FreeBSD__ #include #endif @@ -72,6 +76,7 @@ double TARGET_MS_PROFILE[4] = { 2, 12, 96, 480 }; #define INCREMENT 0 #define INCREMENT_MIN 1 #define INCREMENT_MAX PW_MAX +#define MANGLE 0 #define SEPARATOR ':' #define BITMAP_MIN 16 #define BITMAP_MAX 24 @@ -86,7 +91,7 @@ double TARGET_MS_PROFILE[4] = { 2, 12, 96, 480 }; #define KERNEL_COMBS 1024 #define KERNEL_BFS 1024 #define KERNEL_THREADS_MAX 256 -#define KERNEL_THREADS_MAX_CPU 16 +#define KERNEL_THREADS_MAX_CPU 1 #define POWERTUNE_ENABLE 0 #define LOGFILE_DISABLE 0 #define SCRYPT_TMTO 0 @@ -443,6 +448,7 @@ const char *USAGE_BIG[] = " -i, --increment | | Enable mask increment mode |", " --increment-min | Num | Start mask incrementing at X | --increment-min=4", " --increment-max | Num | Stop mask incrementing at X | --increment-max=8", + " --mangle | | Mangle password before hashing |", "", "- [ Hash modes ] -", "", @@ -468,14 +474,14 @@ const char *USAGE_BIG[] = " 40 | md5($salt.unicode($pass)) | Raw Hash, Salted and / or Iterated", " 3800 | md5($salt.$pass.$salt) | Raw Hash, Salted and / or Iterated", " 3710 | md5($salt.md5($pass)) | Raw Hash, Salted and / or Iterated", - " 2600 | md5(md5($pass) | Raw Hash, Salted and / or Iterated", + " 2600 | md5(md5($pass)) | Raw Hash, Salted and / or Iterated", " 4300 | md5(strtoupper(md5($pass))) | Raw Hash, Salted and / or Iterated", " 4400 | md5(sha1($pass)) | Raw Hash, Salted and / or Iterated", " 110 | sha1($pass.$salt) | Raw Hash, Salted and / or Iterated", " 120 | sha1($salt.$pass) | Raw Hash, Salted and / or Iterated", " 130 | sha1(unicode($pass).$salt) | Raw Hash, Salted and / or Iterated", " 140 | sha1($salt.unicode($pass)) | Raw Hash, Salted and / or Iterated", - " 4500 | sha1(sha1($pass) | Raw Hash, Salted and / or Iterated", + " 4500 | sha1(sha1($pass)) | Raw Hash, Salted and / or Iterated", " 4700 | sha1(md5($pass)) | Raw Hash, Salted and / or Iterated", " 4900 | sha1($salt.$pass.$salt) | Raw Hash, Salted and / or Iterated", " 1410 | sha256($pass.$salt) | Raw Hash, Salted and / or Iterated", @@ -925,6 +931,9 @@ void status_display () if (data.devices_status == STATUS_INIT) return; if (data.devices_status == STATUS_STARTING) return; + // in this case some required buffers are free'd, ascii_digest() would run into segfault + if (data.shutdown_inner == 1) return; + if (data.machine_readable == 1) { status_display_machine_readable (); @@ -1339,26 +1348,22 @@ void status_display () } else { - char display_etc[32] = { 0 }; + char display_etc[32] = { 0 }; + char display_runtime[32] = { 0 }; struct tm tm_etc; + struct tm tm_runtime; struct tm *tmp = NULL; #ifdef WIN - tmp = _gmtime64 (&sec_etc); - #else - tmp = gmtime (&sec_etc); - #endif if (tmp != NULL) { - memset (&tm_etc, 0, sizeof (tm_etc)); - memcpy (&tm_etc, tmp, sizeof (tm_etc)); format_timer_display (&tm_etc, display_etc, sizeof (display_etc)); @@ -1376,7 +1381,43 @@ void status_display () if (etc[etc_len - 1] == '\n') etc[etc_len - 1] = 0; if (etc[etc_len - 2] == '\r') etc[etc_len - 2] = 0; - log_info ("Time.Estimated.: %s (%s)", etc, display_etc); + if (data.runtime) + { + time_t runtime_cur; + + time (&runtime_cur); + + #ifdef WIN + + __time64_t runtime_left = data.proc_start + data.runtime + data.prepare_time - runtime_cur; + + tmp = _gmtime64 (&runtime_left); + + #else + + time_t runtime_left = data.proc_start + data.runtime + data.prepare_time - runtime_cur; + + tmp = gmtime (&runtime_left); + + #endif + + if ((tmp != NULL) && (runtime_left > 0) && (runtime_left < sec_etc)) + { + memcpy (&tm_runtime, tmp, sizeof (tm_runtime)); + + format_timer_display (&tm_runtime, display_runtime, sizeof (display_runtime)); + + log_info ("Time.Estimated.: %s (%s), but limited (%s)", etc, display_etc, display_runtime); + } + else + { + log_info ("Time.Estimated.: %s (%s), but limit exceeded", etc, display_etc); + } + } + else + { + log_info ("Time.Estimated.: %s (%s)", etc, display_etc); + } } } } @@ -1683,6 +1724,8 @@ static void status_benchmark () if (data.devices_status == STATUS_INIT) return; if (data.devices_status == STATUS_STARTING) return; + if (data.shutdown_inner == 1) return; + if (data.machine_readable == 1) { status_benchmark_automate (); @@ -1775,31 +1818,43 @@ static void status_benchmark () * hashcat -only- functions */ -static void generate_source_kernel_filename (const uint attack_exec, const uint attack_kern, const uint kern_type, char *shared_dir, char *source_file) +static void generate_source_kernel_filename (const uint attack_exec, const uint attack_kern, const uint kern_type, uint mangle, char *shared_dir, char *source_file) { + char const * postfix = ""; + if (mangle == 1) + { + postfix = "_m"; + } + if (attack_exec == ATTACK_EXEC_INSIDE_KERNEL) { if (attack_kern == ATTACK_KERN_STRAIGHT) - snprintf (source_file, 255, "%s/OpenCL/m%05d_a0.cl", shared_dir, (int) kern_type); + snprintf (source_file, 255, "%s/OpenCL/m%05d_a0%s.cl", shared_dir, (int) kern_type, postfix); else if (attack_kern == ATTACK_KERN_COMBI) - snprintf (source_file, 255, "%s/OpenCL/m%05d_a1.cl", shared_dir, (int) kern_type); + snprintf (source_file, 255, "%s/OpenCL/m%05d_a1%s.cl", shared_dir, (int) kern_type, postfix); else if (attack_kern == ATTACK_KERN_BF) - snprintf (source_file, 255, "%s/OpenCL/m%05d_a3.cl", shared_dir, (int) kern_type); + snprintf (source_file, 255, "%s/OpenCL/m%05d_a3%s.cl", shared_dir, (int) kern_type, postfix); } else snprintf (source_file, 255, "%s/OpenCL/m%05d.cl", shared_dir, (int) kern_type); } -static void generate_cached_kernel_filename (const uint attack_exec, const uint attack_kern, const uint kern_type, char *profile_dir, const char *device_name_chksum, char *cached_file) +static void generate_cached_kernel_filename (const uint attack_exec, const uint attack_kern, const uint kern_type, uint mangle, char *profile_dir, const char *device_name_chksum, char *cached_file) { + char const * postfix = ""; + if (mangle == 1) + { + postfix = "_m"; + } + if (attack_exec == ATTACK_EXEC_INSIDE_KERNEL) { if (attack_kern == ATTACK_KERN_STRAIGHT) - snprintf (cached_file, 255, "%s/kernels/m%05d_a0.%s.kernel", profile_dir, (int) kern_type, device_name_chksum); + snprintf (cached_file, 255, "%s/kernels/m%05d_a0%s.%s.kernel", profile_dir, (int) kern_type, postfix, device_name_chksum); else if (attack_kern == ATTACK_KERN_COMBI) - snprintf (cached_file, 255, "%s/kernels/m%05d_a1.%s.kernel", profile_dir, (int) kern_type, device_name_chksum); + snprintf (cached_file, 255, "%s/kernels/m%05d_a1%s.%s.kernel", profile_dir, (int) kern_type, postfix, device_name_chksum); else if (attack_kern == ATTACK_KERN_BF) - snprintf (cached_file, 255, "%s/kernels/m%05d_a3.%s.kernel", profile_dir, (int) kern_type, device_name_chksum); + snprintf (cached_file, 255, "%s/kernels/m%05d_a3%s.%s.kernel", profile_dir, (int) kern_type, postfix, device_name_chksum); } else { @@ -1918,9 +1973,18 @@ static void clear_prompt () fflush (stdout); } -static void gidd_to_pw_t (hc_device_param_t *device_param, const u64 gidd, pw_t *pw) +static int gidd_to_pw_t (hc_device_param_t *device_param, const u64 gidd, pw_t *pw) { - hc_clEnqueueReadBuffer (data.ocl, device_param->command_queue, device_param->d_pws_buf, CL_TRUE, gidd * sizeof (pw_t), sizeof (pw_t), pw, 0, NULL, NULL); + cl_int CL_err = hc_clEnqueueReadBuffer (data.ocl, device_param->command_queue, device_param->d_pws_buf, CL_TRUE, gidd * sizeof (pw_t), sizeof (pw_t), pw, 0, NULL, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clEnqueueReadBuffer(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } + + return 0; } static void check_hash (hc_device_param_t *device_param, plain_t *plain) @@ -2255,13 +2319,22 @@ static void check_hash (hc_device_param_t *device_param, plain_t *plain) } } -static void check_cracked (hc_device_param_t *device_param, const uint salt_pos) +static int check_cracked (hc_device_param_t *device_param, const uint salt_pos) { salt_t *salt_buf = &data.salts_buf[salt_pos]; u32 num_cracked; - hc_clEnqueueReadBuffer (data.ocl, device_param->command_queue, device_param->d_result, CL_TRUE, 0, sizeof (u32), &num_cracked, 0, NULL, NULL); + cl_int CL_err; + + CL_err = hc_clEnqueueReadBuffer (data.ocl, device_param->command_queue, device_param->d_result, CL_TRUE, 0, sizeof (u32), &num_cracked, 0, NULL, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clEnqueueReadBuffer(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } if (num_cracked) { @@ -2271,7 +2344,14 @@ static void check_cracked (hc_device_param_t *device_param, const uint salt_pos) plain_t *cracked = (plain_t *) mycalloc (num_cracked, sizeof (plain_t)); - hc_clEnqueueReadBuffer (data.ocl, device_param->command_queue, device_param->d_plain_bufs, CL_TRUE, 0, num_cracked * sizeof (plain_t), cracked, 0, NULL, NULL); + CL_err = hc_clEnqueueReadBuffer (data.ocl, device_param->command_queue, device_param->d_plain_bufs, CL_TRUE, 0, num_cracked * sizeof (plain_t), cracked, 0, NULL, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clEnqueueReadBuffer(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } uint cpt_cracked = 0; @@ -2334,13 +2414,29 @@ static void check_cracked (hc_device_param_t *device_param, const uint salt_pos) memset (data.digests_shown_tmp, 0, salt_buf->digests_cnt * sizeof (uint)); - hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_digests_shown, CL_TRUE, salt_buf->digests_offset * sizeof (uint), salt_buf->digests_cnt * sizeof (uint), &data.digests_shown_tmp[salt_buf->digests_offset], 0, NULL, NULL); + CL_err = hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_digests_shown, CL_TRUE, salt_buf->digests_offset * sizeof (uint), salt_buf->digests_cnt * sizeof (uint), &data.digests_shown_tmp[salt_buf->digests_offset], 0, NULL, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clEnqueueWriteBuffer(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } } num_cracked = 0; - hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_result, CL_TRUE, 0, sizeof (u32), &num_cracked, 0, NULL, NULL); + CL_err = hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_result, CL_TRUE, 0, sizeof (u32), &num_cracked, 0, NULL, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clEnqueueWriteBuffer(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } } + + return 0; } // stolen from princeprocessor ;) @@ -2645,8 +2741,10 @@ static void save_hash () unlink (old_hashfile); } -static void run_kernel (const uint kern_run, hc_device_param_t *device_param, const uint num, const uint event_update, const uint iteration) +static int run_kernel (const uint kern_run, hc_device_param_t *device_param, const uint num, const uint event_update, const uint iteration) { + cl_int CL_err = CL_SUCCESS; + uint num_elements = num; device_param->kernel_params_buf32[33] = data.combs_mode; @@ -2667,17 +2765,24 @@ static void run_kernel (const uint kern_run, hc_device_param_t *device_param, co case KERN_RUN_3: kernel = device_param->kernel3; break; } - hc_clSetKernelArg (data.ocl, kernel, 24, sizeof (cl_uint), device_param->kernel_params[24]); - hc_clSetKernelArg (data.ocl, kernel, 25, sizeof (cl_uint), device_param->kernel_params[25]); - hc_clSetKernelArg (data.ocl, kernel, 26, sizeof (cl_uint), device_param->kernel_params[26]); - hc_clSetKernelArg (data.ocl, kernel, 27, sizeof (cl_uint), device_param->kernel_params[27]); - hc_clSetKernelArg (data.ocl, kernel, 28, sizeof (cl_uint), device_param->kernel_params[28]); - hc_clSetKernelArg (data.ocl, kernel, 29, sizeof (cl_uint), device_param->kernel_params[29]); - hc_clSetKernelArg (data.ocl, kernel, 30, sizeof (cl_uint), device_param->kernel_params[30]); - hc_clSetKernelArg (data.ocl, kernel, 31, sizeof (cl_uint), device_param->kernel_params[31]); - hc_clSetKernelArg (data.ocl, kernel, 32, sizeof (cl_uint), device_param->kernel_params[32]); - hc_clSetKernelArg (data.ocl, kernel, 33, sizeof (cl_uint), device_param->kernel_params[33]); - hc_clSetKernelArg (data.ocl, kernel, 34, sizeof (cl_uint), device_param->kernel_params[34]); + CL_err |= hc_clSetKernelArg (data.ocl, kernel, 24, sizeof (cl_uint), device_param->kernel_params[24]); + CL_err |= hc_clSetKernelArg (data.ocl, kernel, 25, sizeof (cl_uint), device_param->kernel_params[25]); + CL_err |= hc_clSetKernelArg (data.ocl, kernel, 26, sizeof (cl_uint), device_param->kernel_params[26]); + CL_err |= hc_clSetKernelArg (data.ocl, kernel, 27, sizeof (cl_uint), device_param->kernel_params[27]); + CL_err |= hc_clSetKernelArg (data.ocl, kernel, 28, sizeof (cl_uint), device_param->kernel_params[28]); + CL_err |= hc_clSetKernelArg (data.ocl, kernel, 29, sizeof (cl_uint), device_param->kernel_params[29]); + CL_err |= hc_clSetKernelArg (data.ocl, kernel, 30, sizeof (cl_uint), device_param->kernel_params[30]); + CL_err |= hc_clSetKernelArg (data.ocl, kernel, 31, sizeof (cl_uint), device_param->kernel_params[31]); + CL_err |= hc_clSetKernelArg (data.ocl, kernel, 32, sizeof (cl_uint), device_param->kernel_params[32]); + CL_err |= hc_clSetKernelArg (data.ocl, kernel, 33, sizeof (cl_uint), device_param->kernel_params[33]); + CL_err |= hc_clSetKernelArg (data.ocl, kernel, 34, sizeof (cl_uint), device_param->kernel_params[34]); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clSetKernelArg(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } cl_event event; @@ -2686,7 +2791,14 @@ static void run_kernel (const uint kern_run, hc_device_param_t *device_param, co const size_t global_work_size[3] = { num_elements, 32, 1 }; const size_t local_work_size[3] = { kernel_threads / 32, 32, 1 }; - hc_clEnqueueNDRangeKernel (data.ocl, device_param->command_queue, kernel, 2, NULL, global_work_size, local_work_size, 0, NULL, &event); + CL_err = hc_clEnqueueNDRangeKernel (data.ocl, device_param->command_queue, kernel, 2, NULL, global_work_size, local_work_size, 0, NULL, &event); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clEnqueueNDRangeKernel(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } } else { @@ -2703,10 +2815,24 @@ static void run_kernel (const uint kern_run, hc_device_param_t *device_param, co const size_t global_work_size[3] = { num_elements, 1, 1 }; const size_t local_work_size[3] = { kernel_threads, 1, 1 }; - hc_clEnqueueNDRangeKernel (data.ocl, device_param->command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, &event); + CL_err = hc_clEnqueueNDRangeKernel (data.ocl, device_param->command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, &event); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clEnqueueNDRangeKernel(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } } - hc_clFlush (data.ocl, device_param->command_queue); + CL_err = hc_clFlush (data.ocl, device_param->command_queue); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clFlush(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } if (device_param->nvidia_spin_damp) { @@ -2724,13 +2850,27 @@ static void run_kernel (const uint kern_run, hc_device_param_t *device_param, co } } - hc_clWaitForEvents (data.ocl, 1, &event); + CL_err = hc_clWaitForEvents (data.ocl, 1, &event); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clWaitForEvents(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } cl_ulong time_start; cl_ulong time_end; - hc_clGetEventProfilingInfo (data.ocl, event, CL_PROFILING_COMMAND_START, sizeof (time_start), &time_start, NULL); - hc_clGetEventProfilingInfo (data.ocl, event, CL_PROFILING_COMMAND_END, sizeof (time_end), &time_end, NULL); + CL_err |= hc_clGetEventProfilingInfo (data.ocl, event, CL_PROFILING_COMMAND_START, sizeof (time_start), &time_start, NULL); + CL_err |= hc_clGetEventProfilingInfo (data.ocl, event, CL_PROFILING_COMMAND_END, sizeof (time_end), &time_end, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetEventProfilingInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } const double exec_us = (double) (time_end - time_start) / 1000; @@ -2763,13 +2903,31 @@ static void run_kernel (const uint kern_run, hc_device_param_t *device_param, co device_param->exec_pos = exec_pos; } - hc_clReleaseEvent (data.ocl, event); + CL_err = hc_clReleaseEvent (data.ocl, event); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clReleaseEvent(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } + + CL_err = hc_clFinish (data.ocl, device_param->command_queue); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clFinish(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } - hc_clFinish (data.ocl, device_param->command_queue); + return 0; } -static void run_kernel_mp (const uint kern_run, hc_device_param_t *device_param, const uint num) +static int run_kernel_mp (const uint kern_run, hc_device_param_t *device_param, const uint num) { + cl_int CL_err = CL_SUCCESS; + uint num_elements = num; switch (kern_run) @@ -2797,42 +2955,74 @@ static void run_kernel_mp (const uint kern_run, hc_device_param_t *device_param, switch (kern_run) { - case KERN_RUN_MP: hc_clSetKernelArg (data.ocl, kernel, 3, sizeof (cl_ulong), device_param->kernel_params_mp[3]); - hc_clSetKernelArg (data.ocl, kernel, 4, sizeof (cl_uint), device_param->kernel_params_mp[4]); - hc_clSetKernelArg (data.ocl, kernel, 5, sizeof (cl_uint), device_param->kernel_params_mp[5]); - hc_clSetKernelArg (data.ocl, kernel, 6, sizeof (cl_uint), device_param->kernel_params_mp[6]); - hc_clSetKernelArg (data.ocl, kernel, 7, sizeof (cl_uint), device_param->kernel_params_mp[7]); - hc_clSetKernelArg (data.ocl, kernel, 8, sizeof (cl_uint), device_param->kernel_params_mp[8]); + case KERN_RUN_MP: CL_err |= hc_clSetKernelArg (data.ocl, kernel, 3, sizeof (cl_ulong), device_param->kernel_params_mp[3]); + CL_err |= hc_clSetKernelArg (data.ocl, kernel, 4, sizeof (cl_uint), device_param->kernel_params_mp[4]); + CL_err |= hc_clSetKernelArg (data.ocl, kernel, 5, sizeof (cl_uint), device_param->kernel_params_mp[5]); + CL_err |= hc_clSetKernelArg (data.ocl, kernel, 6, sizeof (cl_uint), device_param->kernel_params_mp[6]); + CL_err |= hc_clSetKernelArg (data.ocl, kernel, 7, sizeof (cl_uint), device_param->kernel_params_mp[7]); + CL_err |= hc_clSetKernelArg (data.ocl, kernel, 8, sizeof (cl_uint), device_param->kernel_params_mp[8]); break; - case KERN_RUN_MP_R: hc_clSetKernelArg (data.ocl, kernel, 3, sizeof (cl_ulong), device_param->kernel_params_mp_r[3]); - hc_clSetKernelArg (data.ocl, kernel, 4, sizeof (cl_uint), device_param->kernel_params_mp_r[4]); - hc_clSetKernelArg (data.ocl, kernel, 5, sizeof (cl_uint), device_param->kernel_params_mp_r[5]); - hc_clSetKernelArg (data.ocl, kernel, 6, sizeof (cl_uint), device_param->kernel_params_mp_r[6]); - hc_clSetKernelArg (data.ocl, kernel, 7, sizeof (cl_uint), device_param->kernel_params_mp_r[7]); - hc_clSetKernelArg (data.ocl, kernel, 8, sizeof (cl_uint), device_param->kernel_params_mp_r[8]); + case KERN_RUN_MP_R: CL_err |= hc_clSetKernelArg (data.ocl, kernel, 3, sizeof (cl_ulong), device_param->kernel_params_mp_r[3]); + CL_err |= hc_clSetKernelArg (data.ocl, kernel, 4, sizeof (cl_uint), device_param->kernel_params_mp_r[4]); + CL_err |= hc_clSetKernelArg (data.ocl, kernel, 5, sizeof (cl_uint), device_param->kernel_params_mp_r[5]); + CL_err |= hc_clSetKernelArg (data.ocl, kernel, 6, sizeof (cl_uint), device_param->kernel_params_mp_r[6]); + CL_err |= hc_clSetKernelArg (data.ocl, kernel, 7, sizeof (cl_uint), device_param->kernel_params_mp_r[7]); + CL_err |= hc_clSetKernelArg (data.ocl, kernel, 8, sizeof (cl_uint), device_param->kernel_params_mp_r[8]); break; - case KERN_RUN_MP_L: hc_clSetKernelArg (data.ocl, kernel, 3, sizeof (cl_ulong), device_param->kernel_params_mp_l[3]); - hc_clSetKernelArg (data.ocl, kernel, 4, sizeof (cl_uint), device_param->kernel_params_mp_l[4]); - hc_clSetKernelArg (data.ocl, kernel, 5, sizeof (cl_uint), device_param->kernel_params_mp_l[5]); - hc_clSetKernelArg (data.ocl, kernel, 6, sizeof (cl_uint), device_param->kernel_params_mp_l[6]); - hc_clSetKernelArg (data.ocl, kernel, 7, sizeof (cl_uint), device_param->kernel_params_mp_l[7]); - hc_clSetKernelArg (data.ocl, kernel, 8, sizeof (cl_uint), device_param->kernel_params_mp_l[8]); - hc_clSetKernelArg (data.ocl, kernel, 9, sizeof (cl_uint), device_param->kernel_params_mp_l[9]); + case KERN_RUN_MP_L: CL_err |= hc_clSetKernelArg (data.ocl, kernel, 3, sizeof (cl_ulong), device_param->kernel_params_mp_l[3]); + CL_err |= hc_clSetKernelArg (data.ocl, kernel, 4, sizeof (cl_uint), device_param->kernel_params_mp_l[4]); + CL_err |= hc_clSetKernelArg (data.ocl, kernel, 5, sizeof (cl_uint), device_param->kernel_params_mp_l[5]); + CL_err |= hc_clSetKernelArg (data.ocl, kernel, 6, sizeof (cl_uint), device_param->kernel_params_mp_l[6]); + CL_err |= hc_clSetKernelArg (data.ocl, kernel, 7, sizeof (cl_uint), device_param->kernel_params_mp_l[7]); + CL_err |= hc_clSetKernelArg (data.ocl, kernel, 8, sizeof (cl_uint), device_param->kernel_params_mp_l[8]); + CL_err |= hc_clSetKernelArg (data.ocl, kernel, 9, sizeof (cl_uint), device_param->kernel_params_mp_l[9]); break; } + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clSetKernelArg(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } + const size_t global_work_size[3] = { num_elements, 1, 1 }; const size_t local_work_size[3] = { kernel_threads, 1, 1 }; - hc_clEnqueueNDRangeKernel (data.ocl, device_param->command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL); + CL_err = hc_clEnqueueNDRangeKernel (data.ocl, device_param->command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clEnqueueNDRangeKernel(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } + + CL_err = hc_clFlush (data.ocl, device_param->command_queue); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clFlush(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } + + CL_err = hc_clFinish (data.ocl, device_param->command_queue); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clFinish(): %s\n", val2cstr_cl (CL_err)); - hc_clFlush (data.ocl, device_param->command_queue); + return -1; + } - hc_clFinish (data.ocl, device_param->command_queue); + return 0; } -static void run_kernel_tm (hc_device_param_t *device_param) +static int run_kernel_tm (hc_device_param_t *device_param) { + cl_int CL_err = CL_SUCCESS; + const uint num_elements = 1024; // fixed uint kernel_threads = 32; @@ -2842,15 +3032,40 @@ static void run_kernel_tm (hc_device_param_t *device_param) const size_t global_work_size[3] = { num_elements, 1, 1 }; const size_t local_work_size[3] = { kernel_threads, 1, 1 }; - hc_clEnqueueNDRangeKernel (data.ocl, device_param->command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL); + CL_err = hc_clEnqueueNDRangeKernel (data.ocl, device_param->command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clEnqueueNDRangeKernel(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } + + CL_err = hc_clFlush (data.ocl, device_param->command_queue); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clFlush(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } + + CL_err = hc_clFinish (data.ocl, device_param->command_queue); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clFinish(): %s\n", val2cstr_cl (CL_err)); - hc_clFlush (data.ocl, device_param->command_queue); + return -1; + } - hc_clFinish (data.ocl, device_param->command_queue); + return 0; } -static void run_kernel_amp (hc_device_param_t *device_param, const uint num) +static int run_kernel_amp (hc_device_param_t *device_param, const uint num) { + cl_int CL_err = CL_SUCCESS; + uint num_elements = num; device_param->kernel_params_amp_buf32[5] = data.combs_mode; @@ -2865,21 +3080,53 @@ static void run_kernel_amp (hc_device_param_t *device_param, const uint num) cl_kernel kernel = device_param->kernel_amp; - hc_clSetKernelArg (data.ocl, kernel, 5, sizeof (cl_uint), device_param->kernel_params_amp[5]); - hc_clSetKernelArg (data.ocl, kernel, 6, sizeof (cl_uint), device_param->kernel_params_amp[6]); + CL_err |= hc_clSetKernelArg (data.ocl, kernel, 5, sizeof (cl_uint), device_param->kernel_params_amp[5]); + CL_err |= hc_clSetKernelArg (data.ocl, kernel, 6, sizeof (cl_uint), device_param->kernel_params_amp[6]); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clSetKernelArg(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } const size_t global_work_size[3] = { num_elements, 1, 1 }; const size_t local_work_size[3] = { kernel_threads, 1, 1 }; - hc_clEnqueueNDRangeKernel (data.ocl, device_param->command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL); + CL_err = hc_clEnqueueNDRangeKernel (data.ocl, device_param->command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clEnqueueNDRangeKernel(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } + + CL_err = hc_clFlush (data.ocl, device_param->command_queue); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clFlush(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } + + CL_err = hc_clFinish (data.ocl, device_param->command_queue); - hc_clFlush (data.ocl, device_param->command_queue); + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clFinish(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } - hc_clFinish (data.ocl, device_param->command_queue); + return 0; } -static void run_kernel_memset (hc_device_param_t *device_param, cl_mem buf, const uint value, const uint num) +static int run_kernel_memset (hc_device_param_t *device_param, cl_mem buf, const uint value, const uint num) { + cl_int CL_err = CL_SUCCESS; + const u32 num16d = num / 16; const u32 num16m = num % 16; @@ -2896,81 +3143,84 @@ static void run_kernel_memset (hc_device_param_t *device_param, cl_mem buf, cons cl_kernel kernel = device_param->kernel_memset; - hc_clSetKernelArg (data.ocl, kernel, 0, sizeof (cl_mem), (void *) &buf); - hc_clSetKernelArg (data.ocl, kernel, 1, sizeof (cl_uint), device_param->kernel_params_memset[1]); - hc_clSetKernelArg (data.ocl, kernel, 2, sizeof (cl_uint), device_param->kernel_params_memset[2]); + CL_err |= hc_clSetKernelArg (data.ocl, kernel, 0, sizeof (cl_mem), (void *) &buf); + CL_err |= hc_clSetKernelArg (data.ocl, kernel, 1, sizeof (cl_uint), device_param->kernel_params_memset[1]); + CL_err |= hc_clSetKernelArg (data.ocl, kernel, 2, sizeof (cl_uint), device_param->kernel_params_memset[2]); - const size_t global_work_size[3] = { num_elements, 1, 1 }; - const size_t local_work_size[3] = { kernel_threads, 1, 1 }; + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clSetKernelArg(): %s\n", val2cstr_cl (CL_err)); - hc_clEnqueueNDRangeKernel (data.ocl, device_param->command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL); + return -1; + } - hc_clFlush (data.ocl, device_param->command_queue); + const size_t global_work_size[3] = { num_elements, 1, 1 }; + const size_t local_work_size[3] = { kernel_threads, 1, 1 }; - hc_clFinish (data.ocl, device_param->command_queue); - } + CL_err = hc_clEnqueueNDRangeKernel (data.ocl, device_param->command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL); - if (num16m) - { - u32 tmp[4]; + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clEnqueueNDRangeKernel(): %s\n", val2cstr_cl (CL_err)); - tmp[0] = value; - tmp[1] = value; - tmp[2] = value; - tmp[3] = value; + return -1; + } - hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, buf, CL_TRUE, num16d * 16, num16m, tmp, 0, NULL, NULL); - } -} + CL_err = hc_clFlush (data.ocl, device_param->command_queue); -static void run_kernel_bzero (hc_device_param_t *device_param, cl_mem buf, const size_t size) -{ - run_kernel_memset (device_param, buf, 0, size); + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clFlush(): %s\n", val2cstr_cl (CL_err)); - /* - int rc = -1; + return -1; + } - if (device_param->opencl_v12 && device_param->platform_vendor_id == VENDOR_ID_AMD) - { - // So far tested, amd is the only supporting this OpenCL 1.2 function without segfaulting + CL_err = hc_clFinish (data.ocl, device_param->command_queue); - const cl_uchar zero = 0; + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clFinish(): %s\n", val2cstr_cl (CL_err)); - rc = hc_clEnqueueFillBuffer (data.ocl, device_param->command_queue, buf, &zero, sizeof (cl_uchar), 0, size, 0, NULL, NULL); + return -1; + } } - if (rc != 0) + if (num16m) { - // NOTE: clEnqueueFillBuffer () always fails with -59 - // IOW, it's not supported by Nvidia drivers <= 352.21, also pocl segfaults, also on apple - // How's that possible, OpenCL 1.2 support is advertised?? - // We need to workaround... + u32 tmp[4]; - #define FILLSZ 0x100000 + tmp[0] = value; + tmp[1] = value; + tmp[2] = value; + tmp[3] = value; - char *tmp = (char *) mymalloc (FILLSZ); + CL_err = hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, buf, CL_TRUE, num16d * 16, num16m, tmp, 0, NULL, NULL); - for (size_t i = 0; i < size; i += FILLSZ) + if (CL_err != CL_SUCCESS) { - const size_t left = size - i; + log_error ("ERROR: clEnqueueWriteBuffer(): %s\n", val2cstr_cl (CL_err)); - const size_t fillsz = MIN (FILLSZ, left); - - hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, buf, CL_TRUE, i, fillsz, tmp, 0, NULL, NULL); + return -1; } - - myfree (tmp); } - */ + + return 0; +} + +static int run_kernel_bzero (hc_device_param_t *device_param, cl_mem buf, const size_t size) +{ + return run_kernel_memset (device_param, buf, 0, size); } -static void choose_kernel (hc_device_param_t *device_param, const uint attack_exec, const uint attack_mode, const uint opts_type, const salt_t *salt_buf, const uint highest_pw_len, const uint pws_cnt, const uint fast_iteration) +static int choose_kernel (hc_device_param_t *device_param, const uint attack_exec, const uint attack_mode, const uint opts_type, const salt_t *salt_buf, const uint highest_pw_len, const uint pws_cnt, const uint fast_iteration) { + cl_int CL_err = CL_SUCCESS; + if (data.hash_mode == 2000) { process_stdout (device_param, pws_cnt); - return; + return 0; } if (attack_exec == ATTACK_EXEC_INSIDE_KERNEL) @@ -2985,7 +3235,14 @@ static void choose_kernel (hc_device_param_t *device_param, const uint attack_ex run_kernel_tm (device_param); - hc_clEnqueueCopyBuffer (data.ocl, device_param->command_queue, device_param->d_tm_c, device_param->d_bfs_c, 0, 0, size_tm, 0, NULL, NULL); + CL_err = hc_clEnqueueCopyBuffer (data.ocl, device_param->command_queue, device_param->d_tm_c, device_param->d_bfs_c, 0, 0, size_tm, 0, NULL, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clEnqueueCopyBuffer(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } } } @@ -3012,11 +3269,25 @@ static void choose_kernel (hc_device_param_t *device_param, const uint attack_ex { run_kernel (KERN_RUN_12, device_param, pws_cnt, false, 0); - hc_clEnqueueReadBuffer (data.ocl, device_param->command_queue, device_param->d_hooks, CL_TRUE, 0, device_param->size_hooks, device_param->hooks_buf, 0, NULL, NULL); + CL_err = hc_clEnqueueReadBuffer (data.ocl, device_param->command_queue, device_param->d_hooks, CL_TRUE, 0, device_param->size_hooks, device_param->hooks_buf, 0, NULL, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clEnqueueReadBuffer(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } // do something with data - hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_hooks, CL_TRUE, 0, device_param->size_hooks, device_param->hooks_buf, 0, NULL, NULL); + CL_err = hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_hooks, CL_TRUE, 0, device_param->size_hooks, device_param->hooks_buf, 0, NULL, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clEnqueueWriteBuffer(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } } uint iter = salt_buf->salt_iter; @@ -3067,15 +3338,31 @@ static void choose_kernel (hc_device_param_t *device_param, const uint attack_ex { run_kernel (KERN_RUN_23, device_param, pws_cnt, false, 0); - hc_clEnqueueReadBuffer (data.ocl, device_param->command_queue, device_param->d_hooks, CL_TRUE, 0, device_param->size_hooks, device_param->hooks_buf, 0, NULL, NULL); + CL_err = hc_clEnqueueReadBuffer (data.ocl, device_param->command_queue, device_param->d_hooks, CL_TRUE, 0, device_param->size_hooks, device_param->hooks_buf, 0, NULL, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clEnqueueReadBuffer(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } // do something with data - hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_hooks, CL_TRUE, 0, device_param->size_hooks, device_param->hooks_buf, 0, NULL, NULL); + CL_err = hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_hooks, CL_TRUE, 0, device_param->size_hooks, device_param->hooks_buf, 0, NULL, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clEnqueueWriteBuffer(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } } run_kernel (KERN_RUN_3, device_param, pws_cnt, false, 0); } + + return 0; } static int run_rule_engine (const int rule_len, const char *rule_buf) @@ -3092,11 +3379,20 @@ static int run_rule_engine (const int rule_len, const char *rule_buf) return 1; } -static void run_copy (hc_device_param_t *device_param, const uint pws_cnt) +static int run_copy (hc_device_param_t *device_param, const uint pws_cnt) { + cl_int CL_err = CL_SUCCESS; + if (data.attack_kern == ATTACK_KERN_STRAIGHT) { - hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_pws_buf, CL_TRUE, 0, pws_cnt * sizeof (pw_t), device_param->pws_buf, 0, NULL, NULL); + CL_err = hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_pws_buf, CL_TRUE, 0, pws_cnt * sizeof (pw_t), device_param->pws_buf, 0, NULL, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clEnqueueWriteBuffer(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } } else if (data.attack_kern == ATTACK_KERN_COMBI) { @@ -3154,7 +3450,14 @@ static void run_copy (hc_device_param_t *device_param, const uint pws_cnt) } } - hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_pws_buf, CL_TRUE, 0, pws_cnt * sizeof (pw_t), device_param->pws_buf, 0, NULL, NULL); + CL_err = hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_pws_buf, CL_TRUE, 0, pws_cnt * sizeof (pw_t), device_param->pws_buf, 0, NULL, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clEnqueueWriteBuffer(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } } else if (data.attack_kern == ATTACK_KERN_BF) { @@ -3164,6 +3467,8 @@ static void run_copy (hc_device_param_t *device_param, const uint pws_cnt) run_kernel_mp (KERN_RUN_MP_L, device_param, pws_cnt); } + + return 0; } static double try_run (hc_device_param_t *device_param, const u32 kernel_accel, const u32 kernel_loops) @@ -3188,7 +3493,7 @@ static double try_run (hc_device_param_t *device_param, const u32 kernel_accel, return exec_ms_prev; } -static void autotune (hc_device_param_t *device_param) +static int autotune (hc_device_param_t *device_param) { const double target_ms = TARGET_MS_PROFILE[data.workload_profile - 1]; @@ -3222,7 +3527,7 @@ static void autotune (hc_device_param_t *device_param) device_param->kernel_power = kernel_power; - return; + return 0; } // from here it's clear we are allowed to autotune @@ -3243,14 +3548,28 @@ static void autotune (hc_device_param_t *device_param) device_param->pws_buf[i].pw_len = 7 + (i & 7); } - hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_pws_buf, CL_TRUE, 0, kernel_power_max * sizeof (pw_t), device_param->pws_buf, 0, NULL, NULL); + cl_int CL_err = hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_pws_buf, CL_TRUE, 0, kernel_power_max * sizeof (pw_t), device_param->pws_buf, 0, NULL, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clEnqueueWriteBuffer(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } } if (data.attack_exec == ATTACK_EXEC_INSIDE_KERNEL) { if (data.kernel_rules_cnt > 1) { - hc_clEnqueueCopyBuffer (data.ocl, device_param->command_queue, device_param->d_rules, device_param->d_rules_c, 0, 0, MIN (kernel_loops_max, KERNEL_RULES) * sizeof (kernel_rule_t), 0, NULL, NULL); + cl_int CL_err = hc_clEnqueueCopyBuffer (data.ocl, device_param->command_queue, device_param->d_rules, device_param->d_rules_c, 0, 0, MIN (kernel_loops_max, KERNEL_RULES) * sizeof (kernel_rule_t), 0, NULL, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clEnqueueCopyBuffer(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } } } else @@ -3287,7 +3606,7 @@ static void autotune (hc_device_param_t *device_param) { for (int i = 0; i < STEPS_CNT; i++) { - const u32 kernel_accel_try = 1 << i; + const u32 kernel_accel_try = 1u << i; if (kernel_accel_try < kernel_accel_min) continue; if (kernel_accel_try > kernel_accel_max) break; @@ -3429,9 +3748,11 @@ static void autotune (hc_device_param_t *device_param) } #endif + + return 0; } -static void run_cracker (hc_device_param_t *device_param, const uint pws_cnt) +static int run_cracker (hc_device_param_t *device_param, const uint pws_cnt) { char *line_buf = (char *) mymalloc (HCBUFSIZ); @@ -3663,23 +3984,58 @@ static void run_cracker (hc_device_param_t *device_param, const uint pws_cnt) if (data.attack_mode == ATTACK_MODE_STRAIGHT) { - hc_clEnqueueCopyBuffer (data.ocl, device_param->command_queue, device_param->d_rules, device_param->d_rules_c, innerloop_pos * sizeof (kernel_rule_t), 0, innerloop_left * sizeof (kernel_rule_t), 0, NULL, NULL); + cl_int CL_err = hc_clEnqueueCopyBuffer (data.ocl, device_param->command_queue, device_param->d_rules, device_param->d_rules_c, innerloop_pos * sizeof (kernel_rule_t), 0, innerloop_left * sizeof (kernel_rule_t), 0, NULL, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clEnqueueCopyBuffer(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } } else if (data.attack_mode == ATTACK_MODE_COMBI) { - hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_combs_c, CL_TRUE, 0, innerloop_left * sizeof (comb_t), device_param->combs_buf, 0, NULL, NULL); + cl_int CL_err = hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_combs_c, CL_TRUE, 0, innerloop_left * sizeof (comb_t), device_param->combs_buf, 0, NULL, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clEnqueueWriteBuffer(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } } else if (data.attack_mode == ATTACK_MODE_BF) { - hc_clEnqueueCopyBuffer (data.ocl, device_param->command_queue, device_param->d_bfs, device_param->d_bfs_c, 0, 0, innerloop_left * sizeof (bf_t), 0, NULL, NULL); + cl_int CL_err = hc_clEnqueueCopyBuffer (data.ocl, device_param->command_queue, device_param->d_bfs, device_param->d_bfs_c, 0, 0, innerloop_left * sizeof (bf_t), 0, NULL, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clEnqueueCopyBuffer(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } } else if (data.attack_mode == ATTACK_MODE_HYBRID1) { - hc_clEnqueueCopyBuffer (data.ocl, device_param->command_queue, device_param->d_combs, device_param->d_combs_c, 0, 0, innerloop_left * sizeof (comb_t), 0, NULL, NULL); + cl_int CL_err = hc_clEnqueueCopyBuffer (data.ocl, device_param->command_queue, device_param->d_combs, device_param->d_combs_c, 0, 0, innerloop_left * sizeof (comb_t), 0, NULL, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clEnqueueCopyBuffer(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } } else if (data.attack_mode == ATTACK_MODE_HYBRID2) { - hc_clEnqueueCopyBuffer (data.ocl, device_param->command_queue, device_param->d_combs, device_param->d_combs_c, 0, 0, innerloop_left * sizeof (comb_t), 0, NULL, NULL); + cl_int CL_err = hc_clEnqueueCopyBuffer (data.ocl, device_param->command_queue, device_param->d_combs, device_param->d_combs_c, 0, 0, innerloop_left * sizeof (comb_t), 0, NULL, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clEnqueueCopyBuffer(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } } if (data.benchmark == 1) @@ -3687,7 +4043,9 @@ static void run_cracker (hc_device_param_t *device_param, const uint pws_cnt) hc_timer_set (&device_param->timer_speed); } - choose_kernel (device_param, data.attack_exec, data.attack_mode, data.opts_type, salt_buf, highest_pw_len, pws_cnt, fast_iteration); + int rc = choose_kernel (device_param, data.attack_exec, data.attack_mode, data.opts_type, salt_buf, highest_pw_len, pws_cnt, fast_iteration); + + if (rc == -1) return -1; if (data.devices_status == STATUS_STOP_AT_CHECKPOINT) check_checkpoint (); @@ -3755,6 +4113,8 @@ static void run_cracker (hc_device_param_t *device_param, const uint pws_cnt) device_param->speed_pos = speed_pos; myfree (line_buf); + + return 0; } static void load_segment (wl_data_t *wl_data, FILE *fd) @@ -4179,8 +4539,11 @@ static void *thread_monitor (void *p) { if (data.hm_nvapi) { - NV_GPU_PERF_POLICIES_INFO_PARAMS_V1 perfPolicies_info = { 0 }; - NV_GPU_PERF_POLICIES_STATUS_PARAMS_V1 perfPolicies_status = { 0 }; + NV_GPU_PERF_POLICIES_INFO_PARAMS_V1 perfPolicies_info; + NV_GPU_PERF_POLICIES_STATUS_PARAMS_V1 perfPolicies_status; + + memset (&perfPolicies_info, 0, sizeof (NV_GPU_PERF_POLICIES_INFO_PARAMS_V1)); + memset (&perfPolicies_status, 0, sizeof (NV_GPU_PERF_POLICIES_STATUS_PARAMS_V1)); perfPolicies_info.version = MAKE_NVAPI_VERSION (NV_GPU_PERF_POLICIES_INFO_PARAMS_V1, 1); perfPolicies_status.version = MAKE_NVAPI_VERSION (NV_GPU_PERF_POLICIES_STATUS_PARAMS_V1, 1); @@ -4306,7 +4669,7 @@ static void *thread_monitor (void *p) hm_set_fanspeed_with_device_id_nvapi (device_id, fan_speed_new, 1); #endif - #ifdef LINUX + #ifdef __linux__ hm_set_fanspeed_with_device_id_xnvctrl (device_id, fan_speed_new); #endif } @@ -4343,7 +4706,7 @@ static void *thread_monitor (void *p) time (&runtime_cur); - int runtime_left = data.proc_start + data.runtime - runtime_cur; + int runtime_left = data.proc_start + data.runtime + data.prepare_time - runtime_cur; if (runtime_left <= 0) { @@ -5798,6 +6161,22 @@ int main (int argc, char **argv) umask (077); + /** + * There's some buggy OpenCL runtime that do not support -I. + * A workaround is to chdir() to the OpenCL folder, + * then compile the kernels, + * then chdir() back to where we came from so we need to save it first + */ + + char cwd[1024]; + + if (getcwd (cwd, sizeof (cwd) - 1) == NULL) + { + log_error ("ERROR: getcwd(): %s", strerror (errno)); + + return -1; + } + /** * Real init */ @@ -5810,6 +6189,10 @@ int main (int argc, char **argv) data.proc_start = proc_start; + time_t prepare_start; + + time (&prepare_start); + int myargc = argc; char **myargv = argv; @@ -5874,6 +6257,7 @@ int main (int argc, char **argv) uint increment = INCREMENT; uint increment_min = INCREMENT_MIN; uint increment_max = INCREMENT_MAX; + uint mangle = MANGLE; char *cpu_affinity = NULL; OCL_PTR *ocl = NULL; char *opencl_devices = NULL; @@ -5938,6 +6322,7 @@ int main (int argc, char **argv) #define IDX_INCREMENT 'i' #define IDX_INCREMENT_MIN 0xff12 #define IDX_INCREMENT_MAX 0xff13 + #define IDX_MANGLE 0xff80 #define IDX_OUTFILE 'o' #define IDX_OUTFILE_FORMAT 0xff14 #define IDX_OUTFILE_AUTOHEX_DISABLE 0xff39 @@ -6065,6 +6450,7 @@ int main (int argc, char **argv) {"increment", no_argument, 0, IDX_INCREMENT}, {"increment-min", required_argument, 0, IDX_INCREMENT_MIN}, {"increment-max", required_argument, 0, IDX_INCREMENT_MAX}, + {"mangle", no_argument, 0, IDX_MANGLE}, {"custom-charset1", required_argument, 0, IDX_CUSTOM_CHARSET_1}, {"custom-charset2", required_argument, 0, IDX_CUSTOM_CHARSET_2}, {"custom-charset3", required_argument, 0, IDX_CUSTOM_CHARSET_3}, @@ -6093,7 +6479,7 @@ int main (int argc, char **argv) case IDX_SESSION: session = optarg; break; case IDX_SHOW: show = 1; break; case IDX_LEFT: left = 1; break; - case '?': return (-1); + case '?': return -1; } } @@ -6101,7 +6487,7 @@ int main (int argc, char **argv) { log_error ("ERROR: Invalid argument specified"); - return (-1); + return -1; } /** @@ -6112,14 +6498,14 @@ int main (int argc, char **argv) { log_info ("%s", VERSION_TAG); - return (0); + return 0; } if (usage) { usage_big_print (PROGNAME); - return (0); + return 0; } /** @@ -6134,11 +6520,26 @@ int main (int argc, char **argv) char *exec_path = get_exec_path (); - #ifdef LINUX + + #if defined(__linux__) || defined(__APPLE__) || defined(__FreeBSD__) char *resolved_install_folder = realpath (INSTALL_FOLDER, NULL); char *resolved_exec_path = realpath (exec_path, NULL); + if (resolved_install_folder == NULL) + { + log_error ("ERROR: %s: %s", resolved_install_folder, strerror (errno)); + + return -1; + } + + if (resolved_exec_path == NULL) + { + log_error ("ERROR: %s: %s", resolved_exec_path, strerror (errno)); + + return -1; + } + char *install_dir = get_install_dir (resolved_exec_path); char *profile_dir = NULL; char *session_dir = NULL; @@ -6183,6 +6584,66 @@ int main (int argc, char **argv) myfree (exec_path); + /** + * There's alot of problem related to bad support -I parameters when building the kernel. + * Each OpenCL runtime handles it slightly different. + * The most problematic is with new AMD drivers on Windows, which can not handle quote characters! + * The best workaround found so far is to modify the TMP variable (only inside hashcat process) before the runtime is load + */ + + char cpath[1024] = { 0 }; + + #if _WIN + + snprintf (cpath, sizeof (cpath) - 1, "%s\\OpenCL\\", shared_dir); + + char *cpath_real = mymalloc (MAX_PATH); + + if (GetFullPathName (cpath, MAX_PATH, cpath_real, NULL) == 0) + { + log_error ("ERROR: %s: %s", cpath, "GetFullPathName()"); + + return -1; + } + + #else + + snprintf (cpath, sizeof (cpath) - 1, "%s/OpenCL/", shared_dir); + + char *cpath_real = mymalloc (PATH_MAX); + + if (realpath (cpath, cpath_real) == NULL) + { + log_error ("ERROR: %s: %s", cpath, strerror (errno)); + + return -1; + } + + #endif + + //if (getenv ("TMP") == NULL) + if (1) + { + char tmp[1000]; + + snprintf (tmp, sizeof (tmp) - 1, "TMP=%s", cpath_real); + + putenv (tmp); + } + + #if _WIN + + naive_replace (cpath_real, '\\', '/'); + + // not escaping here, windows using quotes later + // naive_escape (cpath_real, PATH_MAX, ' ', '\\'); + + #else + + naive_escape (cpath_real, PATH_MAX, ' ', '\\'); + + #endif + /** * kernel cache, we need to make sure folder exist */ @@ -6219,7 +6680,7 @@ int main (int argc, char **argv) if (show == 1) log_error ("ERROR: Mixing --restore parameter and --show is not supported"); else log_error ("ERROR: Mixing --restore parameter and --left is not supported"); - return (-1); + return -1; } // this allows the user to use --show and --left while cracking (i.e. while another instance of hashcat is running) @@ -6248,7 +6709,7 @@ int main (int argc, char **argv) { log_error ("ERROR: Incompatible restore-file version"); - return (-1); + return -1; } myargc = rd->argc; @@ -6376,6 +6837,7 @@ int main (int argc, char **argv) increment_min_chgd = 1; break; case IDX_INCREMENT_MAX: increment_max = atoi (optarg); increment_max_chgd = 1; break; + case IDX_MANGLE: mangle = 1; break; case IDX_CUSTOM_CHARSET_1: custom_charset_1 = optarg; break; case IDX_CUSTOM_CHARSET_2: custom_charset_2 = optarg; break; case IDX_CUSTOM_CHARSET_3: custom_charset_3 = optarg; break; @@ -6383,7 +6845,7 @@ int main (int argc, char **argv) default: log_error ("ERROR: Invalid argument specified"); - return (-1); + return -1; } } @@ -6391,7 +6853,7 @@ int main (int argc, char **argv) { log_error ("ERROR: Invalid argument specified"); - return (-1); + return -1; } /** @@ -6442,21 +6904,21 @@ int main (int argc, char **argv) { log_error ("ERROR: Invalid attack-mode specified"); - return (-1); + return -1; } if (runtime_chgd && runtime == 0) // just added to remove compiler warnings for runtime_chgd { log_error ("ERROR: Invalid runtime specified"); - return (-1); + return -1; } if (hash_mode_chgd && hash_mode > 13800) // just added to remove compiler warnings for hash_mode_chgd { log_error ("ERROR: Invalid hash-type specified"); - return (-1); + return -1; } // renamed hash modes @@ -6475,7 +6937,7 @@ int main (int argc, char **argv) { log_error ("Old -m specified, use -m %d instead", n); - return (-1); + return -1; } } @@ -6485,7 +6947,7 @@ int main (int argc, char **argv) { log_error ("ERROR: Mixing support for user names and hashes of type %s is not supported", strhashtype (hash_mode)); - return (-1); + return -1; } } @@ -6493,7 +6955,7 @@ int main (int argc, char **argv) { log_error ("ERROR: Invalid outfile-format specified"); - return (-1); + return -1; } if (left == 1) @@ -6504,7 +6966,7 @@ int main (int argc, char **argv) { log_error ("ERROR: Mixing outfile-format > 1 with left parameter is not allowed"); - return (-1); + return -1; } } else @@ -6521,7 +6983,7 @@ int main (int argc, char **argv) { log_error ("ERROR: Mixing outfile-format > 7 with show parameter is not allowed"); - return (-1); + return -1; } } } @@ -6530,49 +6992,49 @@ int main (int argc, char **argv) { log_error ("ERROR: Invalid increment-min specified"); - return (-1); + return -1; } if (increment_max > INCREMENT_MAX) { log_error ("ERROR: Invalid increment-max specified"); - return (-1); + return -1; } if (increment_min > increment_max) { log_error ("ERROR: Invalid increment-min specified"); - return (-1); + return -1; } if ((increment == 1) && (attack_mode == ATTACK_MODE_STRAIGHT)) { log_error ("ERROR: Increment is not allowed in attack-mode 0"); - return (-1); + return -1; } if ((increment == 0) && (increment_min_chgd == 1)) { log_error ("ERROR: Increment-min is only supported combined with increment switch"); - return (-1); + return -1; } if ((increment == 0) && (increment_max_chgd == 1)) { log_error ("ERROR: Increment-max is only supported combined with increment switch"); - return (-1); + return -1; } if (rp_files_cnt && rp_gen) { log_error ("ERROR: Use of both rules-file and rules-generate is not supported"); - return (-1); + return -1; } if (rp_files_cnt || rp_gen) @@ -6581,7 +7043,7 @@ int main (int argc, char **argv) { log_error ("ERROR: Use of rules-file or rules-generate only allowed in attack-mode 0"); - return (-1); + return -1; } } @@ -6589,7 +7051,7 @@ int main (int argc, char **argv) { log_error ("ERROR: Invalid rp-gen-func-min specified"); - return (-1); + return -1; } if (kernel_accel_chgd == 1) @@ -6601,21 +7063,21 @@ int main (int argc, char **argv) log_info ("You can use --force to override this but do not post error reports if you do so"); log_info (""); - return (-1); + return -1; } if (kernel_accel < 1) { log_error ("ERROR: Invalid kernel-accel specified"); - return (-1); + return -1; } if (kernel_accel > 1024) { log_error ("ERROR: Invalid kernel-accel specified"); - return (-1); + return -1; } } @@ -6628,21 +7090,21 @@ int main (int argc, char **argv) log_info ("You can use --force to override this but do not post error reports if you do so"); log_info (""); - return (-1); + return -1; } if (kernel_loops < 1) { log_error ("ERROR: Invalid kernel-loops specified"); - return (-1); + return -1; } if (kernel_loops > 1024) { log_error ("ERROR: Invalid kernel-loops specified"); - return (-1); + return -1; } } @@ -6650,14 +7112,14 @@ int main (int argc, char **argv) { log_error ("ERROR: workload-profile %i not available", workload_profile); - return (-1); + return -1; } if (opencl_vector_width_chgd && (!is_power_of_2(opencl_vector_width) || opencl_vector_width > 16)) { log_error ("ERROR: opencl-vector-width %i not allowed", opencl_vector_width); - return (-1); + return -1; } if (show == 1 || left == 1) @@ -6668,14 +7130,14 @@ int main (int argc, char **argv) { log_error ("ERROR: Mixing remove parameter not allowed with show parameter or left parameter"); - return (-1); + return -1; } if (potfile_disable == 1) { log_error ("ERROR: Mixing potfile-disable parameter not allowed with show parameter or left parameter"); - return (-1); + return -1; } } @@ -6696,7 +7158,7 @@ int main (int argc, char **argv) { log_error ("ERROR: Invalid argument for benchmark mode specified"); - return (-1); + return -1; } if (attack_mode_chgd == 1) @@ -6705,7 +7167,7 @@ int main (int argc, char **argv) { log_error ("ERROR: Only attack-mode 3 allowed in benchmark mode"); - return (-1); + return -1; } } } @@ -6736,7 +7198,7 @@ int main (int argc, char **argv) { usage_mini_print (myargv[0]); - return (-1); + return -1; } } else if (attack_kern == ATTACK_KERN_STRAIGHT) @@ -6745,7 +7207,7 @@ int main (int argc, char **argv) { usage_mini_print (myargv[0]); - return (-1); + return -1; } } else if (attack_kern == ATTACK_KERN_COMBI) @@ -6754,7 +7216,7 @@ int main (int argc, char **argv) { usage_mini_print (myargv[0]); - return (-1); + return -1; } } else if (attack_kern == ATTACK_KERN_BF) @@ -6763,14 +7225,14 @@ int main (int argc, char **argv) { usage_mini_print (myargv[0]); - return (-1); + return -1; } } else { usage_mini_print (myargv[0]); - return (-1); + return -1; } } @@ -6785,13 +7247,13 @@ int main (int argc, char **argv) { log_error ("ERROR: Combining show parameter with keyspace parameter is not allowed"); - return (-1); + return -1; } else if (left == 1) { log_error ("ERROR: Combining left parameter with keyspace parameter is not allowed"); - return (-1); + return -1; } potfile_disable = 1; @@ -6831,14 +7293,14 @@ int main (int argc, char **argv) { log_error ("ERROR: Parameter remove-timer require parameter remove enabled"); - return (-1); + return -1; } if (remove_timer < 1) { log_error ("ERROR: Parameter remove-timer must have a value greater than or equal to 1"); - return (-1); + return -1; } } @@ -6850,14 +7312,14 @@ int main (int argc, char **argv) { log_error ("ERROR: Parameter loopback not allowed without rules-file or rules-generate"); - return (-1); + return -1; } } else { log_error ("ERROR: Parameter loopback allowed in attack-mode 0 only"); - return (-1); + return -1; } } @@ -6867,14 +7329,14 @@ int main (int argc, char **argv) { log_error ("ERROR: Parameter debug-mode option is only available with attack-mode 0"); - return (-1); + return -1; } if ((rp_files_cnt == 0) && (rp_gen == 0)) { log_error ("ERROR: Parameter debug-mode not allowed without rules-file or rules-generate"); - return (-1); + return -1; } } @@ -6882,7 +7344,7 @@ int main (int argc, char **argv) { log_error ("ERROR: Invalid debug-mode specified"); - return (-1); + return -1; } if (debug_file != NULL) @@ -6891,7 +7353,7 @@ int main (int argc, char **argv) { log_error ("ERROR: Parameter debug-file requires parameter debug-mode to be set"); - return (-1); + return -1; } } @@ -6901,7 +7363,7 @@ int main (int argc, char **argv) { log_error ("ERROR: Parameter induction-dir not allowed with brute-force attacks"); - return (-1); + return -1; } } @@ -6911,7 +7373,7 @@ int main (int argc, char **argv) { log_error ("ERROR: setting --weak-hash-threshold allowed only in straight-attack mode"); - return (-1); + return -1; } weak_hash_threshold = 0; @@ -6921,7 +7383,7 @@ int main (int argc, char **argv) { log_error ("ERROR: setting --nvidia-spin-damp must be between 0 and 100 (inclusive)"); - return (-1); + return -1; } @@ -6959,14 +7421,14 @@ int main (int argc, char **argv) { log_error ("ERROR: Rename directory %s to %s: %s", induction_directory, induction_directory_mv, strerror (errno)); - return (-1); + return -1; } } else { log_error ("ERROR: %s: %s", induction_directory, strerror (errno)); - return (-1); + return -1; } } @@ -6974,7 +7436,7 @@ int main (int argc, char **argv) { log_error ("ERROR: %s: %s", induction_directory, strerror (errno)); - return (-1); + return -1; } } } @@ -7035,7 +7497,7 @@ int main (int argc, char **argv) { log_error ("ERROR: Directory specified in outfile-check '%s' is not a valid directory", outfile_check_directory); - return (-1); + return -1; } } else if (outfile_check_dir == NULL) @@ -7044,7 +7506,7 @@ int main (int argc, char **argv) { log_error ("ERROR: %s: %s", outfile_check_directory, strerror (errno)); - return (-1); + return -1; } } } @@ -7199,6 +7661,7 @@ int main (int argc, char **argv) logfile_top_uint (increment); logfile_top_uint (increment_max); logfile_top_uint (increment_min); + logfile_top_uint (mangle); logfile_top_uint (keyspace); logfile_top_uint (left); logfile_top_uint (logfile_disable); @@ -7453,6 +7916,32 @@ int main (int argc, char **argv) dgst_pos1 = 3; dgst_pos2 = 2; dgst_pos3 = 1; +// mangle = 0; + break; + + case 1: hash_type = HASH_TYPE_MD5; + salt_type = SALT_TYPE_NONE; + attack_exec = ATTACK_EXEC_INSIDE_KERNEL; + opts_type = OPTS_TYPE_PT_GENERATE_LE + | OPTS_TYPE_PT_ADD80 + | OPTS_TYPE_PT_ADDBITS14; + kern_type = KERN_TYPE_MD5; + dgst_size = DGST_SIZE_4_4; + parse_func = md5_parse_hash; + sort_by_digest = sort_by_digest_4_4; + opti_type = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_PRECOMPUTE_INIT + | OPTI_TYPE_PRECOMPUTE_MERKLE + | OPTI_TYPE_MEET_IN_MIDDLE + | OPTI_TYPE_EARLY_SKIP + | OPTI_TYPE_NOT_ITERATED + | OPTI_TYPE_NOT_SALTED + | OPTI_TYPE_RAW_HASH; + dgst_pos0 = 0; + dgst_pos1 = 3; + dgst_pos2 = 2; + dgst_pos3 = 1; +// mangle = 1; break; case 10: hash_type = HASH_TYPE_MD5; @@ -11211,7 +11700,7 @@ int main (int argc, char **argv) dgst_pos3 = 6; break; - default: usage_mini_print (PROGNAME); return (-1); + default: usage_mini_print (PROGNAME); return -1; } /** @@ -11234,7 +11723,7 @@ int main (int argc, char **argv) { log_error ("ERROR: Parameter hex-salt not valid for hash-type %u", hash_mode); - return (-1); + return -1; } } @@ -11456,7 +11945,7 @@ int main (int argc, char **argv) { log_error ("ERROR: %s: %s", potfile, strerror (errno)); - return (-1); + return -1; } if (outfile != NULL) @@ -11467,7 +11956,7 @@ int main (int argc, char **argv) fclose (pot_fp); - return (-1); + return -1; } } else @@ -11485,7 +11974,7 @@ int main (int argc, char **argv) { log_error ("ERROR: %s: %s", potfile, strerror (errno)); - return (-1); + return -1; } data.pot_fp = pot_fp; @@ -11770,7 +12259,7 @@ int main (int argc, char **argv) { log_error ("ERROR: %s: %s", data.hashfile, strerror (errno)); - return (-1); + return -1; } hashes_avail = st.st_size / sizeof (hccap_t); @@ -11794,7 +12283,7 @@ int main (int argc, char **argv) { log_error ("ERROR: %s: %s", hashfile, strerror (errno)); - return (-1); + return -1; } if (data.quiet == 0) log_info_nn ("Counting lines in %s", hashfile); @@ -11809,7 +12298,7 @@ int main (int argc, char **argv) fclose (fp); - return (-1); + return -1; } hashlist_format = hlfmt_detect (fp, 100); // 100 = max numbers to "scan". could be hashes_avail, too @@ -11820,7 +12309,7 @@ int main (int argc, char **argv) fclose (fp); - return (-1); + return -1; } fclose (fp); @@ -11972,7 +12461,7 @@ int main (int argc, char **argv) { log_error ("ERROR: hccap file not specified"); - return (-1); + return -1; } hashlist_mode = HL_MODE_FILE; @@ -11985,7 +12474,7 @@ int main (int argc, char **argv) { log_error ("ERROR: %s: %s", hash_buf, strerror (errno)); - return (-1); + return -1; } if (hashes_avail < 1) @@ -11994,7 +12483,7 @@ int main (int argc, char **argv) fclose (fp); - return (-1); + return -1; } uint hccap_size = sizeof (hccap_t); @@ -12165,7 +12654,7 @@ int main (int argc, char **argv) { log_error ("ERROR: %s: %s", hashfile, strerror (errno)); - return (-1); + return -1; } uint line_num = 0; @@ -12731,7 +13220,7 @@ int main (int argc, char **argv) if (data.quiet == 0) log_info_nn (""); - return (0); + return 0; } if ((keyspace == 0) && (stdout_flag == 0)) @@ -12740,7 +13229,7 @@ int main (int argc, char **argv) { log_error ("ERROR: No hashes loaded"); - return (-1); + return -1; } } @@ -12821,7 +13310,7 @@ int main (int argc, char **argv) { log_error ("ERROR: Hashfile and Outfile are not allowed to point to the same file"); - return (-1); + return -1; } #endif @@ -12830,7 +13319,7 @@ int main (int argc, char **argv) { log_error ("ERROR: Hashfile and Outfile are not allowed to point to the same file"); - return (-1); + return -1; } #endif } @@ -12966,6 +13455,7 @@ int main (int argc, char **argv) // here we have in line_buf: ESSID:MAC1:MAC2 (without the plain) // manipulate salt_buf + memset (line_buf_cpy, 0, HCBUFSIZ); memcpy (line_buf_cpy, line_buf, i); char *mac2_pos = strrchr (line_buf_cpy, ':'); @@ -13400,14 +13890,14 @@ int main (int argc, char **argv) if (bitmap_max < bitmap_min) bitmap_max = bitmap_min; - uint *bitmap_s1_a = (uint *) mymalloc ((1 << bitmap_max) * sizeof (uint)); - uint *bitmap_s1_b = (uint *) mymalloc ((1 << bitmap_max) * sizeof (uint)); - uint *bitmap_s1_c = (uint *) mymalloc ((1 << bitmap_max) * sizeof (uint)); - uint *bitmap_s1_d = (uint *) mymalloc ((1 << bitmap_max) * sizeof (uint)); - uint *bitmap_s2_a = (uint *) mymalloc ((1 << bitmap_max) * sizeof (uint)); - uint *bitmap_s2_b = (uint *) mymalloc ((1 << bitmap_max) * sizeof (uint)); - uint *bitmap_s2_c = (uint *) mymalloc ((1 << bitmap_max) * sizeof (uint)); - uint *bitmap_s2_d = (uint *) mymalloc ((1 << bitmap_max) * sizeof (uint)); + uint *bitmap_s1_a = (uint *) mymalloc ((1u << bitmap_max) * sizeof (uint)); + uint *bitmap_s1_b = (uint *) mymalloc ((1u << bitmap_max) * sizeof (uint)); + uint *bitmap_s1_c = (uint *) mymalloc ((1u << bitmap_max) * sizeof (uint)); + uint *bitmap_s1_d = (uint *) mymalloc ((1u << bitmap_max) * sizeof (uint)); + uint *bitmap_s2_a = (uint *) mymalloc ((1u << bitmap_max) * sizeof (uint)); + uint *bitmap_s2_b = (uint *) mymalloc ((1u << bitmap_max) * sizeof (uint)); + uint *bitmap_s2_c = (uint *) mymalloc ((1u << bitmap_max) * sizeof (uint)); + uint *bitmap_s2_d = (uint *) mymalloc ((1u << bitmap_max) * sizeof (uint)); uint bitmap_bits; uint bitmap_nums; @@ -13418,7 +13908,7 @@ int main (int argc, char **argv) { if (data.quiet == 0) log_info_nn ("Generating bitmap tables with %u bits...", bitmap_bits); - bitmap_nums = 1 << bitmap_bits; + bitmap_nums = 1u << bitmap_bits; bitmap_mask = bitmap_nums - 1; @@ -13432,7 +13922,7 @@ int main (int argc, char **argv) break; } - bitmap_nums = 1 << bitmap_bits; + bitmap_nums = 1u << bitmap_bits; bitmap_mask = bitmap_nums - 1; @@ -13494,7 +13984,7 @@ int main (int argc, char **argv) { log_error ("ERROR: %s: %s", rp_file, strerror (errno)); - return (-1); + return -1; } while (!feof (fp)) @@ -13660,7 +14150,7 @@ int main (int argc, char **argv) { log_error ("ERROR: No valid rules left"); - return (-1); + return -1; } /** @@ -13675,7 +14165,14 @@ int main (int argc, char **argv) if (keyspace == 0) { - hc_clGetPlatformIDs (data.ocl, CL_PLATFORMS_MAX, platforms, &platforms_cnt); + cl_int CL_err = hc_clGetPlatformIDs (data.ocl, CL_PLATFORMS_MAX, platforms, &platforms_cnt); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetPlatformIDs(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } if (platforms_cnt == 0) { @@ -13688,7 +14185,7 @@ int main (int argc, char **argv) log_info (" NVidia users require NVidia drivers 346.59 or later (recommended 361.x or later)"); log_info (""); - return (-1); + return -1; } if (opencl_platforms_filter != (uint) -1) @@ -13699,7 +14196,7 @@ int main (int argc, char **argv) { log_error ("ERROR: The platform selected by the --opencl-platforms parameter is larger than the number of available platforms (%d)", platforms_cnt); - return (-1); + return -1; } } } @@ -13715,11 +14212,22 @@ int main (int argc, char **argv) for (uint platform_id = 0; platform_id < platforms_cnt; platform_id++) { - if ((opencl_platforms_filter & (1 << platform_id)) == 0) continue; + if ((opencl_platforms_filter & (1u << platform_id)) == 0) continue; cl_platform_id platform = platforms[platform_id]; - hc_clGetDeviceIDs (data.ocl, platform, CL_DEVICE_TYPE_ALL, DEVICES_MAX, platform_devices, &platform_devices_cnt); + cl_int CL_err = hc_clGetDeviceIDs (data.ocl, platform, CL_DEVICE_TYPE_ALL, DEVICES_MAX, platform_devices, &platform_devices_cnt); + + if (CL_err != CL_SUCCESS) + { + //log_error ("ERROR: clGetDeviceIDs(): %s\n", val2cstr_cl (CL_err)); + + //return -1; + + // Silently ignore at this point, it will be reused later and create a note for the user at that point + + continue; + } for (uint platform_devices_id = 0; platform_devices_id < platform_devices_cnt; platform_devices_id++) { @@ -13727,7 +14235,14 @@ int main (int argc, char **argv) cl_device_type device_type; - hc_clGetDeviceInfo (data.ocl, device, CL_DEVICE_TYPE, sizeof (device_type), &device_type, NULL); + cl_int CL_err = hc_clGetDeviceInfo (data.ocl, device, CL_DEVICE_TYPE, sizeof (device_type), &device_type, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetDeviceInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } device_types_all |= device_type; } @@ -13771,13 +14286,20 @@ int main (int argc, char **argv) for (uint platform_id = 0; platform_id < platforms_cnt; platform_id++) { - cl_platform_id platform = platforms[platform_id]; + cl_int CL_err = CL_SUCCESS; - hc_clGetDeviceIDs (data.ocl, platform, CL_DEVICE_TYPE_ALL, DEVICES_MAX, platform_devices, &platform_devices_cnt); + cl_platform_id platform = platforms[platform_id]; char platform_vendor[INFOSZ] = { 0 }; - hc_clGetPlatformInfo (data.ocl, platform, CL_PLATFORM_VENDOR, sizeof (platform_vendor), platform_vendor, NULL); + CL_err = hc_clGetPlatformInfo (data.ocl, platform, CL_PLATFORM_VENDOR, sizeof (platform_vendor), platform_vendor, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetPlatformInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } // find our own platform vendor because pocl and mesa are pushing original vendor_id through opencl // this causes trouble with vendor id based macros @@ -13822,7 +14344,18 @@ int main (int argc, char **argv) platform_vendor_id = VENDOR_ID_GENERIC; } - const uint platform_skipped = ((opencl_platforms_filter & (1 << platform_id)) == 0); + uint platform_skipped = ((opencl_platforms_filter & (1u << platform_id)) == 0); + + CL_err = hc_clGetDeviceIDs (data.ocl, platform, CL_DEVICE_TYPE_ALL, DEVICES_MAX, platform_devices, &platform_devices_cnt); + + if (CL_err != CL_SUCCESS) + { + //log_error ("ERROR: clGetDeviceIDs(): %s\n", val2cstr_cl (CL_err)); + + //return -1; + + platform_skipped = 2; + } if ((benchmark == 1 || quiet == 0) && (algorithm_pos == 0)) { @@ -13838,15 +14371,21 @@ int main (int argc, char **argv) log_info (line); } - else + else if (platform_skipped == 1) { log_info ("OpenCL Platform #%u: %s, skipped", platform_id + 1, platform_vendor); log_info (""); } + else if (platform_skipped == 2) + { + log_info ("OpenCL Platform #%u: %s, skipped! No OpenCL compatible devices found", platform_id + 1, platform_vendor); + log_info (""); + } } } if (platform_skipped == 1) continue; + if (platform_skipped == 2) continue; for (uint platform_devices_id = 0; platform_devices_id < platform_devices_cnt; platform_devices_id++) { @@ -13864,11 +14403,20 @@ int main (int argc, char **argv) device_param->platform_devices_id = platform_devices_id; + device_param->platform = platform; + // device_type cl_device_type device_type; - hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_TYPE, sizeof (device_type), &device_type, NULL); + CL_err = hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_TYPE, sizeof (device_type), &device_type, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetDeviceInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } device_type &= ~CL_DEVICE_TYPE_DEFAULT; @@ -13876,21 +14424,49 @@ int main (int argc, char **argv) // device_name - hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_NAME, 0, NULL, ¶m_value_size); + CL_err = hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_NAME, 0, NULL, ¶m_value_size); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetDeviceInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } char *device_name = (char *) mymalloc (param_value_size); - hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_NAME, param_value_size, device_name, NULL); + CL_err = hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_NAME, param_value_size, device_name, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetDeviceInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } device_param->device_name = device_name; // device_vendor - hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_VENDOR, 0, NULL, ¶m_value_size); + CL_err = hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_VENDOR, 0, NULL, ¶m_value_size); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetDeviceInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } char *device_vendor = (char *) mymalloc (param_value_size); - hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_VENDOR, param_value_size, device_vendor, NULL); + CL_err = hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_VENDOR, param_value_size, device_vendor, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetDeviceInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } device_param->device_vendor = device_vendor; @@ -13941,21 +14517,49 @@ int main (int argc, char **argv) // device_version - hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_VERSION, 0, NULL, ¶m_value_size); + CL_err = hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_VERSION, 0, NULL, ¶m_value_size); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetDeviceInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } char *device_version = (char *) mymalloc (param_value_size); - hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_VERSION, param_value_size, device_version, NULL); + CL_err = hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_VERSION, param_value_size, device_version, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetDeviceInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } device_param->device_version = device_version; // device_opencl_version - hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_OPENCL_C_VERSION, 0, NULL, ¶m_value_size); + CL_err = hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_OPENCL_C_VERSION, 0, NULL, ¶m_value_size); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetDeviceInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } char *device_opencl_version = (char *) mymalloc (param_value_size); - hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_OPENCL_C_VERSION, param_value_size, device_opencl_version, NULL); + CL_err = hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_OPENCL_C_VERSION, param_value_size, device_opencl_version, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetDeviceInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } device_param->opencl_v12 = device_opencl_version[9] > '1' || device_opencl_version[11] >= '2'; @@ -13971,11 +14575,25 @@ int main (int argc, char **argv) { if (opti_type & OPTI_TYPE_USES_BITS_64) { - hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_NATIVE_VECTOR_WIDTH_LONG, sizeof (vector_width), &vector_width, NULL); + CL_err = hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_NATIVE_VECTOR_WIDTH_LONG, sizeof (vector_width), &vector_width, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetDeviceInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } } else { - hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_NATIVE_VECTOR_WIDTH_INT, sizeof (vector_width), &vector_width, NULL); + CL_err = hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_NATIVE_VECTOR_WIDTH_INT, sizeof (vector_width), &vector_width, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetDeviceInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } } } else @@ -13996,7 +14614,14 @@ int main (int argc, char **argv) cl_uint device_processors; - hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_MAX_COMPUTE_UNITS, sizeof (device_processors), &device_processors, NULL); + CL_err = hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_MAX_COMPUTE_UNITS, sizeof (device_processors), &device_processors, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetDeviceInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } device_param->device_processors = device_processors; @@ -14005,7 +14630,14 @@ int main (int argc, char **argv) cl_ulong device_maxmem_alloc; - hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof (device_maxmem_alloc), &device_maxmem_alloc, NULL); + CL_err = hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof (device_maxmem_alloc), &device_maxmem_alloc, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetDeviceInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } device_param->device_maxmem_alloc = MIN (device_maxmem_alloc, 0x7fffffff); @@ -14013,7 +14645,14 @@ int main (int argc, char **argv) cl_ulong device_global_mem; - hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof (device_global_mem), &device_global_mem, NULL); + CL_err = hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof (device_global_mem), &device_global_mem, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetDeviceInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } device_param->device_global_mem = device_global_mem; @@ -14021,7 +14660,14 @@ int main (int argc, char **argv) size_t device_maxworkgroup_size; - hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof (device_maxworkgroup_size), &device_maxworkgroup_size, NULL); + CL_err = hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof (device_maxworkgroup_size), &device_maxworkgroup_size, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetDeviceInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } device_param->device_maxworkgroup_size = device_maxworkgroup_size; @@ -14029,7 +14675,14 @@ int main (int argc, char **argv) cl_uint device_maxclock_frequency; - hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_MAX_CLOCK_FREQUENCY, sizeof (device_maxclock_frequency), &device_maxclock_frequency, NULL); + CL_err = hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_MAX_CLOCK_FREQUENCY, sizeof (device_maxclock_frequency), &device_maxclock_frequency, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetDeviceInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } device_param->device_maxclock_frequency = device_maxclock_frequency; @@ -14037,7 +14690,14 @@ int main (int argc, char **argv) cl_bool device_endian_little; - hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_ENDIAN_LITTLE, sizeof (device_endian_little), &device_endian_little, NULL); + CL_err = hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_ENDIAN_LITTLE, sizeof (device_endian_little), &device_endian_little, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetDeviceInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } if (device_endian_little == CL_FALSE) { @@ -14050,7 +14710,14 @@ int main (int argc, char **argv) cl_bool device_available; - hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_AVAILABLE, sizeof (device_available), &device_available, NULL); + CL_err = hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_AVAILABLE, sizeof (device_available), &device_available, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetDeviceInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } if (device_available == CL_FALSE) { @@ -14063,20 +14730,34 @@ int main (int argc, char **argv) cl_bool device_compiler_available; - hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_COMPILER_AVAILABLE, sizeof (device_compiler_available), &device_compiler_available, NULL); + CL_err = hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_COMPILER_AVAILABLE, sizeof (device_compiler_available), &device_compiler_available, NULL); - if (device_compiler_available == CL_FALSE) + if (CL_err != CL_SUCCESS) { - log_info ("- Device #%u: WARNING: No compiler available for device", device_id + 1); + log_error ("ERROR: clGetDeviceInfo(): %s\n", val2cstr_cl (CL_err)); - device_param->skipped = 1; + return -1; + } + + if (device_compiler_available == CL_FALSE) + { + log_info ("- Device #%u: WARNING: No compiler available for device", device_id + 1); + + device_param->skipped = 1; } // device_execution_capabilities cl_device_exec_capabilities device_execution_capabilities; - hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_EXECUTION_CAPABILITIES, sizeof (device_execution_capabilities), &device_execution_capabilities, NULL); + CL_err = hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_EXECUTION_CAPABILITIES, sizeof (device_execution_capabilities), &device_execution_capabilities, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetDeviceInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } if ((device_execution_capabilities & CL_EXEC_KERNEL) == 0) { @@ -14089,11 +14770,25 @@ int main (int argc, char **argv) size_t device_extensions_size; - hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_EXTENSIONS, 0, NULL, &device_extensions_size); + CL_err = hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_EXTENSIONS, 0, NULL, &device_extensions_size); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetDeviceInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } char *device_extensions = mymalloc (device_extensions_size + 1); - hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_EXTENSIONS, device_extensions_size, device_extensions, NULL); + CL_err = hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_EXTENSIONS, device_extensions_size, device_extensions, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetDeviceInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } if (strstr (device_extensions, "base_atomics") == 0) { @@ -14115,7 +14810,14 @@ int main (int argc, char **argv) cl_ulong device_local_mem_size; - hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_LOCAL_MEM_SIZE, sizeof (device_local_mem_size), &device_local_mem_size, NULL); + CL_err = hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_LOCAL_MEM_SIZE, sizeof (device_local_mem_size), &device_local_mem_size, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetDeviceInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } if (device_local_mem_size < 32768) { @@ -14149,16 +14851,30 @@ int main (int argc, char **argv) // skipped - device_param->skipped |= ((devices_filter & (1 << device_id)) == 0); + device_param->skipped |= ((devices_filter & (1u << device_id)) == 0); device_param->skipped |= ((device_types_filter & (device_type)) == 0); // driver_version - hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DRIVER_VERSION, 0, NULL, ¶m_value_size); + CL_err = hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DRIVER_VERSION, 0, NULL, ¶m_value_size); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetDeviceInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } char *driver_version = (char *) mymalloc (param_value_size); - hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DRIVER_VERSION, param_value_size, driver_version, NULL); + CL_err = hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DRIVER_VERSION, param_value_size, driver_version, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetDeviceInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } device_param->driver_version = driver_version; @@ -14193,7 +14909,7 @@ int main (int argc, char **argv) { need_nvml = 1; - #ifdef LINUX + #ifdef __linux__ need_xnvctrl = 1; #endif @@ -14211,7 +14927,14 @@ int main (int argc, char **argv) #define CL_DEVICE_KERNEL_EXEC_TIMEOUT_NV 0x4005 - hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_KERNEL_EXEC_TIMEOUT_NV, sizeof (kernel_exec_timeout), &kernel_exec_timeout, NULL); + CL_err = hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_KERNEL_EXEC_TIMEOUT_NV, sizeof (kernel_exec_timeout), &kernel_exec_timeout, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetDeviceInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } device_param->kernel_exec_timeout = kernel_exec_timeout; @@ -14221,8 +14944,23 @@ int main (int argc, char **argv) #define CL_DEVICE_COMPUTE_CAPABILITY_MAJOR_NV 0x4000 #define CL_DEVICE_COMPUTE_CAPABILITY_MINOR_NV 0x4001 - hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_COMPUTE_CAPABILITY_MINOR_NV, sizeof (sm_minor), &sm_minor, NULL); - hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_COMPUTE_CAPABILITY_MAJOR_NV, sizeof (sm_major), &sm_major, NULL); + CL_err = hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_COMPUTE_CAPABILITY_MINOR_NV, sizeof (sm_minor), &sm_minor, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetDeviceInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } + + CL_err = hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_COMPUTE_CAPABILITY_MAJOR_NV, sizeof (sm_major), &sm_major, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetDeviceInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } device_param->sm_minor = sm_minor; device_param->sm_major = sm_major; @@ -14312,7 +15050,7 @@ int main (int argc, char **argv) log_info ("You can use --force to override this but do not post error reports if you do so"); log_info (""); - return (-1); + return -1; } if (catalyst_warn == 1) @@ -14327,7 +15065,7 @@ int main (int argc, char **argv) log_info ("You can use --force to override this but do not post error reports if you do so"); log_info (""); - return (-1); + return -1; } } else if (platform_vendor_id == VENDOR_ID_NV) @@ -14354,7 +15092,7 @@ int main (int argc, char **argv) log_info ("A good alternative is the free pocl >= v0.13, but make sure to use a LLVM >= v3.8"); log_info (""); - return (-1); + return -1; } } } @@ -14438,7 +15176,7 @@ int main (int argc, char **argv) { log_error ("ERROR: No devices found/left"); - return (-1); + return -1; } // additional check to see if the user has chosen a device that is not within the range of available devices (i.e. larger than devices_cnt) @@ -14451,7 +15189,7 @@ int main (int argc, char **argv) { log_error ("ERROR: The device specified by the --opencl-devices parameter is larger than the number of available devices (%d)", devices_cnt); - return (-1); + return -1; } } @@ -14464,10 +15202,15 @@ int main (int argc, char **argv) */ #ifdef HAVE_HWMON - hm_attrs_t hm_adapters_adl[DEVICES_MAX] = { { 0 } }; - hm_attrs_t hm_adapters_nvapi[DEVICES_MAX] = { { 0 } }; - hm_attrs_t hm_adapters_nvml[DEVICES_MAX] = { { 0 } }; - hm_attrs_t hm_adapters_xnvctrl[DEVICES_MAX] = { { 0 } }; + hm_attrs_t hm_adapters_adl[DEVICES_MAX]; + hm_attrs_t hm_adapters_nvapi[DEVICES_MAX]; + hm_attrs_t hm_adapters_nvml[DEVICES_MAX]; + hm_attrs_t hm_adapters_xnvctrl[DEVICES_MAX]; + + memset (hm_adapters_adl, 0, sizeof (hm_adapters_adl)); + memset (hm_adapters_nvapi, 0, sizeof (hm_adapters_nvapi)); + memset (hm_adapters_nvml, 0, sizeof (hm_adapters_nvml)); + memset (hm_adapters_xnvctrl, 0, sizeof (hm_adapters_xnvctrl)); if (gpu_temp_disable == 0) { @@ -14573,13 +15316,13 @@ int main (int argc, char **argv) int hm_adapters_num; - if (get_adapters_num_adl (data.hm_adl, &hm_adapters_num) != 0) return (-1); + if (get_adapters_num_adl (data.hm_adl, &hm_adapters_num) != 0) return -1; // adapter info LPAdapterInfo lpAdapterInfo = hm_get_adapter_info_adl (data.hm_adl, hm_adapters_num); - if (lpAdapterInfo == NULL) return (-1); + if (lpAdapterInfo == NULL) return -1; // get a list (of ids of) valid/usable adapters @@ -14638,7 +15381,7 @@ int main (int argc, char **argv) { log_error ("ERROR: Invalid values for gpu-temp-abort. Parameter gpu-temp-abort is less than gpu-temp-retain."); - return (-1); + return -1; } } @@ -14675,6 +15418,8 @@ int main (int argc, char **argv) log_info ("Rules: %u", kernel_rules_cnt); } + log_info ("Mangle: %d", mangle); + if (opti_type) { log_info ("Applicable Optimizers:"); @@ -14797,7 +15542,7 @@ int main (int argc, char **argv) { log_error ("ERROR: Failed to get ADL PowerControl Capabilities"); - return (-1); + return -1; } // first backup current value, we will restore it later @@ -14817,14 +15562,14 @@ int main (int argc, char **argv) { log_error ("ERROR: Failed to get current ADL PowerControl settings"); - return (-1); + return -1; } if ((ADL_rc = hm_ADL_Overdrive_PowerControl_Set (data.hm_adl, data.hm_device[device_id].adl, powertune.iMaxValue)) != ADL_OK) { log_error ("ERROR: Failed to set new ADL PowerControl values"); - return (-1); + return -1; } // clocks @@ -14837,7 +15582,7 @@ int main (int argc, char **argv) { log_error ("ERROR: Failed to get ADL memory and engine clock frequency"); - return (-1); + return -1; } // Query capabilities only to see if profiles were not "damaged", if so output a warning but do accept the users profile settings @@ -14848,7 +15593,7 @@ int main (int argc, char **argv) { log_error ("ERROR: Failed to get ADL device capabilities"); - return (-1); + return -1; } int engine_clock_max = caps.sEngineClockRange.iMax * 0.6666; @@ -14885,7 +15630,7 @@ int main (int argc, char **argv) { log_info ("ERROR: Failed to set ADL performance state"); - return (-1); + return -1; } local_free (performance_state); @@ -14902,14 +15647,14 @@ int main (int argc, char **argv) { log_error ("ERROR: Failed to get current ADL PowerControl settings"); - return (-1); + return -1; } if ((ADL_rc = hm_ADL_Overdrive_PowerControl_Set (data.hm_adl, data.hm_device[device_id].adl, powertune.iMaxValue)) != ADL_OK) { log_error ("ERROR: Failed to set new ADL PowerControl values"); - return (-1); + return -1; } } } @@ -14964,6 +15709,8 @@ int main (int argc, char **argv) for (uint device_id = 0; device_id < data.devices_cnt; device_id++) { + cl_int CL_err = CL_SUCCESS; + /** * host buffer */ @@ -14983,7 +15730,20 @@ int main (int argc, char **argv) * create context for each device */ - device_param->context = hc_clCreateContext (data.ocl, NULL, 1, &device_param->device, NULL, NULL); + cl_context_properties properties[3]; + + properties[0] = CL_CONTEXT_PLATFORM; + properties[1] = (cl_context_properties) device_param->platform; + properties[2] = 0; + + CL_err = hc_clCreateContext (data.ocl, properties, 1, &device_param->device, NULL, NULL, &device_param->context); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clCreateContext(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } /** * create command-queue @@ -14992,7 +15752,14 @@ int main (int argc, char **argv) // not supported with NV // device_param->command_queue = hc_clCreateCommandQueueWithProperties (device_param->context, device_param->device, NULL); - device_param->command_queue = hc_clCreateCommandQueue (data.ocl, device_param->context, device_param->device, CL_QUEUE_PROFILING_ENABLE); + CL_err = hc_clCreateCommandQueue (data.ocl, device_param->context, device_param->device, CL_QUEUE_PROFILING_ENABLE, &device_param->command_queue); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clCreateCommandQueue(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } /** * kernel threads: some algorithms need a fixed kernel-threads count @@ -15002,6 +15769,9 @@ int main (int argc, char **argv) uint kernel_threads = MIN (KERNEL_THREADS_MAX, device_param->device_maxworkgroup_size); + if (hash_mode == 8900) kernel_threads = 64; // Scrypt + if (hash_mode == 9300) kernel_threads = 64; // Scrypt + if (device_param->device_type & CL_DEVICE_TYPE_CPU) { kernel_threads = KERNEL_THREADS_MAX_CPU; @@ -15011,9 +15781,7 @@ int main (int argc, char **argv) if (hash_mode == 3000) kernel_threads = 64; // DES if (hash_mode == 3200) kernel_threads = 8; // Blowfish if (hash_mode == 7500) kernel_threads = 64; // RC4 - if (hash_mode == 8900) kernel_threads = 64; // Scrypt if (hash_mode == 9000) kernel_threads = 8; // Blowfish - if (hash_mode == 9300) kernel_threads = 64; // Scrypt if (hash_mode == 9700) kernel_threads = 64; // RC4 if (hash_mode == 9710) kernel_threads = 64; // RC4 if (hash_mode == 9800) kernel_threads = 64; // RC4 @@ -15117,7 +15885,7 @@ int main (int argc, char **argv) } } - data.scrypt_tmp_size = (128 * scrypt_r); + data.scrypt_tmp_size = (128 * scrypt_r * scrypt_p); device_param->kernel_accel_min = 1; device_param->kernel_accel_max = 8; @@ -15128,7 +15896,7 @@ int main (int argc, char **argv) { size_scrypt = (128 * scrypt_r) * scrypt_N; - size_scrypt /= 1 << tmto; + size_scrypt /= 1u << tmto; size_scrypt *= device_param->device_processors * device_param->kernel_threads * device_param->kernel_accel_max; @@ -15161,7 +15929,7 @@ int main (int argc, char **argv) return -1; } - if (quiet == 0) log_info ("SCRYPT tmto optimizer value set to: %u, mem: %u\n", data.scrypt_tmto_final, size_scrypt); + if (quiet == 0) log_info ("SCRYPT tmto optimizer value set to: %u, mem: %llu\n", data.scrypt_tmto_final, (unsigned long long int) size_scrypt); } size_t size_scrypt4 = size_scrypt / 4; @@ -15433,46 +16201,19 @@ int main (int argc, char **argv) * default building options */ - char cpath[1024] = { 0 }; - - char build_opts[1024] = { 0 }; - - #if _WIN - - snprintf (cpath, sizeof (cpath) - 1, "%s\\OpenCL\\", shared_dir); - - char *cpath_real = mymalloc (MAX_PATH); - - if (GetFullPathName (cpath, MAX_PATH, cpath_real, NULL) == 0) + if (chdir (cpath_real) == -1) { - log_error ("ERROR: %s: %s", cpath, "GetFullPathName()"); + log_error ("ERROR: %s: %s", cpath_real, strerror (errno)); return -1; } - naive_replace (cpath_real, '\\', '/'); - - // not escaping here, windows has quotes + char build_opts[1024] = { 0 }; + #if _WIN snprintf (build_opts, sizeof (build_opts) - 1, "-I \"%s\"", cpath_real); - #else - - snprintf (cpath, sizeof (cpath) - 1, "%s/OpenCL/", shared_dir); - - char *cpath_real = mymalloc (PATH_MAX); - - if (realpath (cpath, cpath_real) == NULL) - { - log_error ("ERROR: %s: %s", cpath, strerror (errno)); - - return -1; - } - - naive_escape (cpath_real, PATH_MAX, ' ', '\\'); - snprintf (build_opts, sizeof (build_opts) - 1, "-I %s", cpath_real); - #endif // include check @@ -15502,15 +16243,11 @@ int main (int argc, char **argv) for (int i = 0; i < files_cnt; i++) { - char path[1024] = { 0 }; - - snprintf (path, sizeof (path) - 1, "%s/%s", cpath_real, files_names[i]); - - FILE *fd = fopen (path, "r"); + FILE *fd = fopen (files_names[i], "r"); if (fd == NULL) { - log_error ("ERROR: %s: fopen(): %s", path, strerror (errno)); + log_error ("ERROR: %s: fopen(): %s", files_names[i], strerror (errno)); return -1; } @@ -15521,7 +16258,7 @@ int main (int argc, char **argv) if (n != 1) { - log_error ("ERROR: %s: fread(): %s", path, strerror (errno)); + log_error ("ERROR: %s: fread(): %s", files_names[i], strerror (errno)); return -1; } @@ -15529,8 +16266,6 @@ int main (int argc, char **argv) fclose (fd); } - myfree (cpath_real); - // we don't have sm_* on vendors not NV but it doesn't matter char build_opts_new[1024] = { 0 }; @@ -15554,7 +16289,7 @@ int main (int argc, char **argv) char source_file[256] = { 0 }; - generate_source_kernel_filename (attack_exec, attack_kern, kern_type, shared_dir, source_file); + generate_source_kernel_filename (attack_exec, attack_kern, kern_type, mangle, shared_dir, source_file); struct stat sst; @@ -15571,7 +16306,7 @@ int main (int argc, char **argv) char cached_file[256] = { 0 }; - generate_cached_kernel_filename (attack_exec, attack_kern, kern_type, profile_dir, device_name_chksum, cached_file); + generate_cached_kernel_filename (attack_exec, attack_kern, kern_type, mangle, profile_dir, device_name_chksum, cached_file); int cached = 1; @@ -15598,30 +16333,56 @@ int main (int argc, char **argv) load_kernel (source_file, 1, kernel_lengths, kernel_sources); - device_param->program = hc_clCreateProgramWithSource (data.ocl, device_param->context, 1, (const char **) kernel_sources, NULL); + CL_err = hc_clCreateProgramWithSource (data.ocl, device_param->context, 1, (const char **) kernel_sources, NULL, &device_param->program); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clCreateProgramWithSource(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } + + CL_err = 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 (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clBuildProgram(): %s\n", val2cstr_cl (CL_err)); + + //return -1; + } #ifdef DEBUG size_t build_log_size = 0; - hc_clGetProgramBuildInfo (data.ocl, device_param->program, device_param->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &build_log_size); + CL_err = hc_clGetProgramBuildInfo (data.ocl, device_param->program, device_param->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &build_log_size); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetProgramBuildInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } if (build_log_size > 1) { - char *build_log = (char *) malloc (build_log_size + 1); + char *build_log = (char *) mymalloc (build_log_size + 1); - memset (build_log, 0, build_log_size + 1); + CL_err = hc_clGetProgramBuildInfo (data.ocl, device_param->program, device_param->device, CL_PROGRAM_BUILD_LOG, build_log_size, build_log, NULL); - hc_clGetProgramBuildInfo (data.ocl, device_param->program, device_param->device, CL_PROGRAM_BUILD_LOG, build_log_size, build_log, NULL); + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetProgramBuildInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } puts (build_log); - free (build_log); + myfree (build_log); } #endif - if (rc != 0) + if (CL_err != CL_SUCCESS) { device_param->skipped = true; @@ -15632,11 +16393,25 @@ int main (int argc, char **argv) size_t binary_size; - hc_clGetProgramInfo (data.ocl, device_param->program, CL_PROGRAM_BINARY_SIZES, sizeof (size_t), &binary_size, NULL); + CL_err = hc_clGetProgramInfo (data.ocl, device_param->program, CL_PROGRAM_BINARY_SIZES, sizeof (size_t), &binary_size, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetProgramInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } u8 *binary = (u8 *) mymalloc (binary_size); - hc_clGetProgramInfo (data.ocl, device_param->program, CL_PROGRAM_BINARIES, sizeof (binary), &binary, NULL); + CL_err = hc_clGetProgramInfo (data.ocl, device_param->program, CL_PROGRAM_BINARIES, sizeof (binary), &binary, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetProgramInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } writeProgramBin (cached_file, binary, binary_size); @@ -15650,9 +16425,23 @@ int main (int argc, char **argv) load_kernel (cached_file, 1, kernel_lengths, kernel_sources); - device_param->program = hc_clCreateProgramWithBinary (data.ocl, device_param->context, 1, &device_param->device, kernel_lengths, (const u8 **) kernel_sources, NULL); + CL_err = hc_clCreateProgramWithBinary (data.ocl, device_param->context, 1, &device_param->device, kernel_lengths, (const u8 **) kernel_sources, NULL, &device_param->program); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clCreateProgramWithBinary(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } + + CL_err = hc_clBuildProgram (data.ocl, device_param->program, 1, &device_param->device, build_opts, NULL, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clBuildProgram(): %s\n", val2cstr_cl (CL_err)); - hc_clBuildProgram (data.ocl, device_param->program, 1, &device_param->device, build_opts, NULL, NULL, true); + return -1; + } } } else @@ -15663,7 +16452,14 @@ int main (int argc, char **argv) load_kernel (source_file, 1, kernel_lengths, kernel_sources); - device_param->program = hc_clCreateProgramWithSource (data.ocl, device_param->context, 1, (const char **) kernel_sources, NULL); + CL_err = hc_clCreateProgramWithSource (data.ocl, device_param->context, 1, (const char **) kernel_sources, NULL, &device_param->program); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clCreateProgramWithSource(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } char build_opts_update[1024] = { 0 }; @@ -15680,28 +16476,47 @@ int main (int argc, char **argv) snprintf (build_opts_update, sizeof (build_opts_update) - 1, "%s", build_opts); } - int rc = hc_clBuildProgram (data.ocl, device_param->program, 1, &device_param->device, build_opts_update, NULL, NULL, false); + CL_err = hc_clBuildProgram (data.ocl, device_param->program, 1, &device_param->device, build_opts_update, NULL, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clBuildProgram(): %s\n", val2cstr_cl (CL_err)); + + //return -1; + } #ifdef DEBUG size_t build_log_size = 0; - hc_clGetProgramBuildInfo (data.ocl, device_param->program, device_param->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &build_log_size); + CL_err = hc_clGetProgramBuildInfo (data.ocl, device_param->program, device_param->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &build_log_size); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetProgramBuildInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } if (build_log_size > 1) { - char *build_log = (char *) malloc (build_log_size + 1); + char *build_log = (char *) mymalloc (build_log_size + 1); - memset (build_log, 0, build_log_size + 1); + CL_err = hc_clGetProgramBuildInfo (data.ocl, device_param->program, device_param->device, CL_PROGRAM_BUILD_LOG, build_log_size, build_log, NULL); - hc_clGetProgramBuildInfo (data.ocl, device_param->program, device_param->device, CL_PROGRAM_BUILD_LOG, build_log_size, build_log, NULL); + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetProgramBuildInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } puts (build_log); - free (build_log); + myfree (build_log); } #endif - if (rc != 0) + if (CL_err != CL_SUCCESS) { device_param->skipped = true; @@ -15769,11 +16584,25 @@ int main (int argc, char **argv) load_kernel (source_file, 1, kernel_lengths, kernel_sources); - device_param->program_mp = hc_clCreateProgramWithSource (data.ocl, device_param->context, 1, (const char **) kernel_sources, NULL); + CL_err = hc_clCreateProgramWithSource (data.ocl, device_param->context, 1, (const char **) kernel_sources, NULL, &device_param->program_mp); - int rc = hc_clBuildProgram (data.ocl, device_param->program_mp, 1, &device_param->device, build_opts, NULL, NULL, false); + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clCreateProgramWithSource(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } - if (rc != 0) + CL_err = hc_clBuildProgram (data.ocl, device_param->program_mp, 1, &device_param->device, build_opts, NULL, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clBuildProgram(): %s\n", val2cstr_cl (CL_err)); + + //return -1; + } + + if (CL_err != CL_SUCCESS) { device_param->skipped = true; @@ -15784,11 +16613,25 @@ int main (int argc, char **argv) size_t binary_size; - hc_clGetProgramInfo (data.ocl, device_param->program_mp, CL_PROGRAM_BINARY_SIZES, sizeof (size_t), &binary_size, NULL); + CL_err = hc_clGetProgramInfo (data.ocl, device_param->program_mp, CL_PROGRAM_BINARY_SIZES, sizeof (size_t), &binary_size, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetProgramInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } u8 *binary = (u8 *) mymalloc (binary_size); - hc_clGetProgramInfo (data.ocl, device_param->program_mp, CL_PROGRAM_BINARIES, sizeof (binary), &binary, NULL); + CL_err = hc_clGetProgramInfo (data.ocl, device_param->program_mp, CL_PROGRAM_BINARIES, sizeof (binary), &binary, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetProgramInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } writeProgramBin (cached_file, binary, binary_size); @@ -15802,9 +16645,23 @@ int main (int argc, char **argv) load_kernel (cached_file, 1, kernel_lengths, kernel_sources); - device_param->program_mp = hc_clCreateProgramWithBinary (data.ocl, device_param->context, 1, &device_param->device, kernel_lengths, (const u8 **) kernel_sources, NULL); + CL_err = hc_clCreateProgramWithBinary (data.ocl, device_param->context, 1, &device_param->device, kernel_lengths, (const u8 **) kernel_sources, NULL, &device_param->program_mp); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clCreateProgramWithBinary(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } + + CL_err = hc_clBuildProgram (data.ocl, device_param->program_mp, 1, &device_param->device, build_opts, NULL, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clBuildProgram(): %s\n", val2cstr_cl (CL_err)); - hc_clBuildProgram (data.ocl, device_param->program_mp, 1, &device_param->device, build_opts, NULL, NULL, true); + return -1; + } } local_free (kernel_lengths); @@ -15871,11 +16728,25 @@ int main (int argc, char **argv) load_kernel (source_file, 1, kernel_lengths, kernel_sources); - device_param->program_amp = hc_clCreateProgramWithSource (data.ocl, device_param->context, 1, (const char **) kernel_sources, NULL); + CL_err = hc_clCreateProgramWithSource (data.ocl, device_param->context, 1, (const char **) kernel_sources, NULL, &device_param->program_amp); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clCreateProgramWithSource(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } + + CL_err = hc_clBuildProgram (data.ocl, device_param->program_amp, 1, &device_param->device, build_opts, NULL, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clBuildProgram(): %s\n", val2cstr_cl (CL_err)); - int rc = hc_clBuildProgram (data.ocl, device_param->program_amp, 1, &device_param->device, build_opts, NULL, NULL, false); + //return -1; + } - if (rc != 0) + if (CL_err != CL_SUCCESS) { device_param->skipped = true; @@ -15886,11 +16757,25 @@ int main (int argc, char **argv) size_t binary_size; - hc_clGetProgramInfo (data.ocl, device_param->program_amp, CL_PROGRAM_BINARY_SIZES, sizeof (size_t), &binary_size, NULL); + CL_err = hc_clGetProgramInfo (data.ocl, device_param->program_amp, CL_PROGRAM_BINARY_SIZES, sizeof (size_t), &binary_size, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetProgramInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } u8 *binary = (u8 *) mymalloc (binary_size); - hc_clGetProgramInfo (data.ocl, device_param->program_amp, CL_PROGRAM_BINARIES, sizeof (binary), &binary, NULL); + CL_err = hc_clGetProgramInfo (data.ocl, device_param->program_amp, CL_PROGRAM_BINARIES, sizeof (binary), &binary, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetProgramInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } writeProgramBin (cached_file, binary, binary_size); @@ -15904,9 +16789,23 @@ int main (int argc, char **argv) load_kernel (cached_file, 1, kernel_lengths, kernel_sources); - device_param->program_amp = hc_clCreateProgramWithBinary (data.ocl, device_param->context, 1, &device_param->device, kernel_lengths, (const u8 **) kernel_sources, NULL); + CL_err = hc_clCreateProgramWithBinary (data.ocl, device_param->context, 1, &device_param->device, kernel_lengths, (const u8 **) kernel_sources, NULL, &device_param->program_amp); - hc_clBuildProgram (data.ocl, device_param->program_amp, 1, &device_param->device, build_opts, NULL, NULL, true); + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clCreateProgramWithBinary(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } + + CL_err = hc_clBuildProgram (data.ocl, device_param->program_amp, 1, &device_param->device, build_opts, NULL, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clBuildProgram(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } } local_free (kernel_lengths); @@ -15914,6 +16813,15 @@ int main (int argc, char **argv) local_free (kernel_sources); } + // return back to the folder we came from initially (workaround) + + if (chdir (cwd) == -1) + { + log_error ("ERROR: %s: %s", cwd, strerror (errno)); + + return -1; + } + // some algorithm collide too fast, make that impossible if (benchmark == 1) @@ -15928,39 +16836,53 @@ int main (int argc, char **argv) * global buffers */ - device_param->d_pws_buf = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_pws, NULL); - device_param->d_pws_amp_buf = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_pws, NULL); - device_param->d_tmps = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE, size_tmps, NULL); - device_param->d_hooks = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE, size_hooks, NULL); - device_param->d_bitmap_s1_a = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL); - device_param->d_bitmap_s1_b = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL); - device_param->d_bitmap_s1_c = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL); - device_param->d_bitmap_s1_d = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL); - device_param->d_bitmap_s2_a = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL); - device_param->d_bitmap_s2_b = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL); - device_param->d_bitmap_s2_c = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL); - device_param->d_bitmap_s2_d = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL); - device_param->d_plain_bufs = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE, size_plains, NULL); - device_param->d_digests_buf = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_digests, NULL); - device_param->d_digests_shown = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE, size_shown, NULL); - device_param->d_salt_bufs = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_salts, NULL); - device_param->d_result = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE, size_results, NULL); - device_param->d_scryptV0_buf = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE, size_scrypt4, NULL); - device_param->d_scryptV1_buf = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE, size_scrypt4, NULL); - device_param->d_scryptV2_buf = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE, size_scrypt4, NULL); - device_param->d_scryptV3_buf = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE, size_scrypt4, NULL); - - hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s1_a, CL_TRUE, 0, bitmap_size, bitmap_s1_a, 0, NULL, NULL); - hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s1_b, CL_TRUE, 0, bitmap_size, bitmap_s1_b, 0, NULL, NULL); - hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s1_c, CL_TRUE, 0, bitmap_size, bitmap_s1_c, 0, NULL, NULL); - hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s1_d, CL_TRUE, 0, bitmap_size, bitmap_s1_d, 0, NULL, NULL); - hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s2_a, CL_TRUE, 0, bitmap_size, bitmap_s2_a, 0, NULL, NULL); - hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s2_b, CL_TRUE, 0, bitmap_size, bitmap_s2_b, 0, NULL, NULL); - hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s2_c, CL_TRUE, 0, bitmap_size, bitmap_s2_c, 0, NULL, NULL); - hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s2_d, CL_TRUE, 0, bitmap_size, bitmap_s2_d, 0, NULL, NULL); - hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_digests_buf, CL_TRUE, 0, size_digests, data.digests_buf, 0, NULL, NULL); - hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_digests_shown, CL_TRUE, 0, size_shown, data.digests_shown, 0, NULL, NULL); - hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_salt_bufs, CL_TRUE, 0, size_salts, data.salts_buf, 0, NULL, NULL); + CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_pws, NULL, &device_param->d_pws_buf); + CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_pws, NULL, &device_param->d_pws_amp_buf); + CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE, size_tmps, NULL, &device_param->d_tmps); + CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE, size_hooks, NULL, &device_param->d_hooks); + CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL, &device_param->d_bitmap_s1_a); + CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL, &device_param->d_bitmap_s1_b); + CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL, &device_param->d_bitmap_s1_c); + CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL, &device_param->d_bitmap_s1_d); + CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL, &device_param->d_bitmap_s2_a); + CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL, &device_param->d_bitmap_s2_b); + CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL, &device_param->d_bitmap_s2_c); + CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL, &device_param->d_bitmap_s2_d); + CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE, size_plains, NULL, &device_param->d_plain_bufs); + CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_digests, NULL, &device_param->d_digests_buf); + CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE, size_shown, NULL, &device_param->d_digests_shown); + CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_salts, NULL, &device_param->d_salt_bufs); + CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE, size_results, NULL, &device_param->d_result); + CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE, size_scrypt4, NULL, &device_param->d_scryptV0_buf); + CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE, size_scrypt4, NULL, &device_param->d_scryptV1_buf); + CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE, size_scrypt4, NULL, &device_param->d_scryptV2_buf); + CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE, size_scrypt4, NULL, &device_param->d_scryptV3_buf); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clCreateBuffer(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } + + CL_err |= hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s1_a, CL_TRUE, 0, bitmap_size, bitmap_s1_a, 0, NULL, NULL); + CL_err |= hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s1_b, CL_TRUE, 0, bitmap_size, bitmap_s1_b, 0, NULL, NULL); + CL_err |= hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s1_c, CL_TRUE, 0, bitmap_size, bitmap_s1_c, 0, NULL, NULL); + CL_err |= hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s1_d, CL_TRUE, 0, bitmap_size, bitmap_s1_d, 0, NULL, NULL); + CL_err |= hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s2_a, CL_TRUE, 0, bitmap_size, bitmap_s2_a, 0, NULL, NULL); + CL_err |= hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s2_b, CL_TRUE, 0, bitmap_size, bitmap_s2_b, 0, NULL, NULL); + CL_err |= hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s2_c, CL_TRUE, 0, bitmap_size, bitmap_s2_c, 0, NULL, NULL); + CL_err |= hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s2_d, CL_TRUE, 0, bitmap_size, bitmap_s2_d, 0, NULL, NULL); + CL_err |= hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_digests_buf, CL_TRUE, 0, size_digests, data.digests_buf, 0, NULL, NULL); + CL_err |= hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_digests_shown, CL_TRUE, 0, size_shown, data.digests_shown, 0, NULL, NULL); + CL_err |= hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_salt_bufs, CL_TRUE, 0, size_salts, data.salts_buf, 0, NULL, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clEnqueueWriteBuffer(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } /** * special buffers @@ -15968,32 +16890,74 @@ int main (int argc, char **argv) if (attack_kern == ATTACK_KERN_STRAIGHT) { - device_param->d_rules = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_rules, NULL); - device_param->d_rules_c = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_rules_c, NULL); + CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_rules, NULL, &device_param->d_rules); + CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_rules_c, NULL, &device_param->d_rules_c); - hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_rules, CL_TRUE, 0, size_rules, kernel_rules_buf, 0, NULL, NULL); + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clCreateBuffer(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } + + CL_err = hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_rules, CL_TRUE, 0, size_rules, kernel_rules_buf, 0, NULL, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clEnqueueWriteBuffer(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } } else if (attack_kern == ATTACK_KERN_COMBI) { - device_param->d_combs = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_combs, NULL); - device_param->d_combs_c = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_combs, NULL); - device_param->d_root_css_buf = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_root_css, NULL); - device_param->d_markov_css_buf = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_markov_css, NULL); + CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_combs, NULL, &device_param->d_combs); + CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_combs, NULL, &device_param->d_combs_c); + CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_root_css, NULL, &device_param->d_root_css_buf); + CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_markov_css, NULL, &device_param->d_markov_css_buf); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clCreateBuffer(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } } else if (attack_kern == ATTACK_KERN_BF) { - device_param->d_bfs = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_bfs, NULL); - device_param->d_bfs_c = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_bfs, NULL); - device_param->d_tm_c = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_tm, NULL); - device_param->d_root_css_buf = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_root_css, NULL); - device_param->d_markov_css_buf = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_markov_css, NULL); + CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_bfs, NULL, &device_param->d_bfs); + CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_bfs, NULL, &device_param->d_bfs_c); + CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_tm, NULL, &device_param->d_tm_c); + CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_root_css, NULL, &device_param->d_root_css_buf); + CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_markov_css, NULL, &device_param->d_markov_css_buf); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clCreateBuffer(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } } if (size_esalts) { - device_param->d_esalt_bufs = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_esalts, NULL); + CL_err = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_esalts, NULL, &device_param->d_esalt_bufs); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clCreateBuffer(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } + + CL_err = hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_esalt_bufs, CL_TRUE, 0, size_esalts, data.esalts_buf, 0, NULL, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clEnqueueWriteBuffer(): %s\n", val2cstr_cl (CL_err)); - hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_esalt_bufs, CL_TRUE, 0, size_esalts, data.esalts_buf, 0, NULL, NULL); + return -1; + } } /** @@ -16154,29 +17118,71 @@ int main (int argc, char **argv) { snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_s%02d", kern_type, 4); - device_param->kernel1 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name); + CL_err = hc_clCreateKernel (data.ocl, device_param->program, kernel_name, &device_param->kernel1); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clCreateKernel(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_s%02d", kern_type, 8); - device_param->kernel2 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name); + CL_err = hc_clCreateKernel (data.ocl, device_param->program, kernel_name, &device_param->kernel2); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clCreateKernel(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_s%02d", kern_type, 16); - device_param->kernel3 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name); + CL_err = hc_clCreateKernel (data.ocl, device_param->program, kernel_name, &device_param->kernel3); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clCreateKernel(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } } else { snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_m%02d", kern_type, 4); - device_param->kernel1 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name); + CL_err = hc_clCreateKernel (data.ocl, device_param->program, kernel_name, &device_param->kernel1); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clCreateKernel(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_m%02d", kern_type, 8); - device_param->kernel2 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name); + CL_err = hc_clCreateKernel (data.ocl, device_param->program, kernel_name, &device_param->kernel2); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clCreateKernel(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_m%02d", kern_type, 16); - device_param->kernel3 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name); + CL_err = hc_clCreateKernel (data.ocl, device_param->program, kernel_name, &device_param->kernel3); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clCreateKernel(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } } if (data.attack_mode == ATTACK_MODE_BF) @@ -16185,9 +17191,23 @@ int main (int argc, char **argv) { snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_tm", kern_type); - device_param->kernel_tm = hc_clCreateKernel (data.ocl, device_param->program, kernel_name); + CL_err = hc_clCreateKernel (data.ocl, device_param->program, kernel_name, &device_param->kernel_tm); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clCreateKernel(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } + + CL_err = hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel_tm, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetKernelWorkGroupInfo(): %s\n", val2cstr_cl (CL_err)); - hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel_tm, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp); + return -1; + } } } } @@ -16195,96 +17215,236 @@ int main (int argc, char **argv) { snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_init", kern_type); - device_param->kernel1 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name); + CL_err = hc_clCreateKernel (data.ocl, device_param->program, kernel_name, &device_param->kernel1); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clCreateKernel(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_loop", kern_type); - device_param->kernel2 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name); + CL_err = hc_clCreateKernel (data.ocl, device_param->program, kernel_name, &device_param->kernel2); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clCreateKernel(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_comp", kern_type); - device_param->kernel3 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name); + CL_err = hc_clCreateKernel (data.ocl, device_param->program, kernel_name, &device_param->kernel3); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clCreateKernel(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } if (opts_type & OPTS_TYPE_HOOK12) { snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_hook12", kern_type); - device_param->kernel12 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name); + CL_err = hc_clCreateKernel (data.ocl, device_param->program, kernel_name, &device_param->kernel12); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clCreateKernel(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } + + CL_err = hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel12, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetKernelWorkGroupInfo(): %s\n", val2cstr_cl (CL_err)); - hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel12, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp); + return -1; + } } if (opts_type & OPTS_TYPE_HOOK23) { snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_hook23", kern_type); - device_param->kernel23 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name); + CL_err = hc_clCreateKernel (data.ocl, device_param->program, kernel_name, &device_param->kernel23); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clCreateKernel(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } - hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel23, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp); + CL_err = hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel23, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetKernelWorkGroupInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } } } - hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel1, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp); - hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel2, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp); - hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel3, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp); + CL_err |= hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel1, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp); + CL_err |= hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel2, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp); + CL_err |= hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel3, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetKernelWorkGroupInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } for (uint i = 0; i <= 23; i++) { - hc_clSetKernelArg (data.ocl, device_param->kernel1, i, sizeof (cl_mem), device_param->kernel_params[i]); - hc_clSetKernelArg (data.ocl, device_param->kernel2, i, sizeof (cl_mem), device_param->kernel_params[i]); - hc_clSetKernelArg (data.ocl, device_param->kernel3, i, sizeof (cl_mem), device_param->kernel_params[i]); + CL_err |= hc_clSetKernelArg (data.ocl, device_param->kernel1, i, sizeof (cl_mem), device_param->kernel_params[i]); + CL_err |= hc_clSetKernelArg (data.ocl, device_param->kernel2, i, sizeof (cl_mem), device_param->kernel_params[i]); + CL_err |= hc_clSetKernelArg (data.ocl, device_param->kernel3, i, sizeof (cl_mem), device_param->kernel_params[i]); + + if (opts_type & OPTS_TYPE_HOOK12) CL_err |= hc_clSetKernelArg (data.ocl, device_param->kernel12, i, sizeof (cl_mem), device_param->kernel_params[i]); + if (opts_type & OPTS_TYPE_HOOK23) CL_err |= hc_clSetKernelArg (data.ocl, device_param->kernel23, i, sizeof (cl_mem), device_param->kernel_params[i]); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clSetKernelArg(): %s\n", val2cstr_cl (CL_err)); - if (opts_type & OPTS_TYPE_HOOK12) hc_clSetKernelArg (data.ocl, device_param->kernel12, i, sizeof (cl_mem), device_param->kernel_params[i]); - if (opts_type & OPTS_TYPE_HOOK23) hc_clSetKernelArg (data.ocl, device_param->kernel23, i, sizeof (cl_mem), device_param->kernel_params[i]); + return -1; + } } for (uint i = 24; i <= 34; i++) { - hc_clSetKernelArg (data.ocl, device_param->kernel1, i, sizeof (cl_uint), device_param->kernel_params[i]); - hc_clSetKernelArg (data.ocl, device_param->kernel2, i, sizeof (cl_uint), device_param->kernel_params[i]); - hc_clSetKernelArg (data.ocl, device_param->kernel3, i, sizeof (cl_uint), device_param->kernel_params[i]); + CL_err |= hc_clSetKernelArg (data.ocl, device_param->kernel1, i, sizeof (cl_uint), device_param->kernel_params[i]); + CL_err |= hc_clSetKernelArg (data.ocl, device_param->kernel2, i, sizeof (cl_uint), device_param->kernel_params[i]); + CL_err |= hc_clSetKernelArg (data.ocl, device_param->kernel3, i, sizeof (cl_uint), device_param->kernel_params[i]); + + if (opts_type & OPTS_TYPE_HOOK12) CL_err |= hc_clSetKernelArg (data.ocl, device_param->kernel12, i, sizeof (cl_uint), device_param->kernel_params[i]); + if (opts_type & OPTS_TYPE_HOOK23) CL_err |= hc_clSetKernelArg (data.ocl, device_param->kernel23, i, sizeof (cl_uint), device_param->kernel_params[i]); - if (opts_type & OPTS_TYPE_HOOK12) hc_clSetKernelArg (data.ocl, device_param->kernel12, i, sizeof (cl_uint), device_param->kernel_params[i]); - if (opts_type & OPTS_TYPE_HOOK23) hc_clSetKernelArg (data.ocl, device_param->kernel23, i, sizeof (cl_uint), device_param->kernel_params[i]); + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clSetKernelArg(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } } // GPU memset - device_param->kernel_memset = hc_clCreateKernel (data.ocl, device_param->program, "gpu_memset"); + CL_err = hc_clCreateKernel (data.ocl, device_param->program, "gpu_memset", &device_param->kernel_memset); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clCreateKernel(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } + + CL_err = hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel_memset, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetKernelWorkGroupInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } + + CL_err |= hc_clSetKernelArg (data.ocl, device_param->kernel_memset, 0, sizeof (cl_mem), device_param->kernel_params_memset[0]); + CL_err |= hc_clSetKernelArg (data.ocl, device_param->kernel_memset, 1, sizeof (cl_uint), device_param->kernel_params_memset[1]); + CL_err |= hc_clSetKernelArg (data.ocl, device_param->kernel_memset, 2, sizeof (cl_uint), device_param->kernel_params_memset[2]); - hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel_memset, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp); + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clSetKernelArg(): %s\n", val2cstr_cl (CL_err)); - hc_clSetKernelArg (data.ocl, device_param->kernel_memset, 0, sizeof (cl_mem), device_param->kernel_params_memset[0]); - hc_clSetKernelArg (data.ocl, device_param->kernel_memset, 1, sizeof (cl_uint), device_param->kernel_params_memset[1]); - hc_clSetKernelArg (data.ocl, device_param->kernel_memset, 2, sizeof (cl_uint), device_param->kernel_params_memset[2]); + return -1; + } // MP start if (attack_mode == ATTACK_MODE_BF) { - device_param->kernel_mp_l = hc_clCreateKernel (data.ocl, device_param->program_mp, "l_markov"); - device_param->kernel_mp_r = hc_clCreateKernel (data.ocl, device_param->program_mp, "r_markov"); + CL_err |= hc_clCreateKernel (data.ocl, device_param->program_mp, "l_markov", &device_param->kernel_mp_l); + CL_err |= hc_clCreateKernel (data.ocl, device_param->program_mp, "r_markov", &device_param->kernel_mp_r); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clCreateKernel(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } - hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel_mp_l, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp); - hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel_mp_r, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp); + CL_err |= hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel_mp_l, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp); + CL_err |= hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel_mp_r, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetKernelWorkGroupInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } if (opts_type & OPTS_TYPE_PT_BITSLICE) { - hc_clSetKernelArg (data.ocl, device_param->kernel_tm, 0, sizeof (cl_mem), device_param->kernel_params_tm[0]); - hc_clSetKernelArg (data.ocl, device_param->kernel_tm, 1, sizeof (cl_mem), device_param->kernel_params_tm[1]); + CL_err |= hc_clSetKernelArg (data.ocl, device_param->kernel_tm, 0, sizeof (cl_mem), device_param->kernel_params_tm[0]); + CL_err |= hc_clSetKernelArg (data.ocl, device_param->kernel_tm, 1, sizeof (cl_mem), device_param->kernel_params_tm[1]); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clSetKernelArg(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } } } else if (attack_mode == ATTACK_MODE_HYBRID1) { - device_param->kernel_mp = hc_clCreateKernel (data.ocl, device_param->program_mp, "C_markov"); + CL_err = hc_clCreateKernel (data.ocl, device_param->program_mp, "C_markov", &device_param->kernel_mp); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clCreateKernel(): %s\n", val2cstr_cl (CL_err)); - hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel_mp, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp); + return -1; + } + + CL_err = hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel_mp, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetKernelWorkGroupInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } } else if (attack_mode == ATTACK_MODE_HYBRID2) { - device_param->kernel_mp = hc_clCreateKernel (data.ocl, device_param->program_mp, "C_markov"); + CL_err = hc_clCreateKernel (data.ocl, device_param->program_mp, "C_markov", &device_param->kernel_mp); - hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel_mp, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp); + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clCreateKernel(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } + + CL_err = hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel_mp, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetKernelWorkGroupInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } } if (attack_exec == ATTACK_EXEC_INSIDE_KERNEL) @@ -16293,9 +17453,23 @@ int main (int argc, char **argv) } else { - device_param->kernel_amp = hc_clCreateKernel (data.ocl, device_param->program_amp, "amp"); + CL_err = hc_clCreateKernel (data.ocl, device_param->program_amp, "amp", &device_param->kernel_amp); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clCreateKernel(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } + + CL_err = hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel_amp, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp); - hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel_amp, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp); + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetKernelWorkGroupInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } } if (attack_exec == ATTACK_EXEC_INSIDE_KERNEL) @@ -16306,12 +17480,26 @@ int main (int argc, char **argv) { for (uint i = 0; i < 5; i++) { - hc_clSetKernelArg (data.ocl, device_param->kernel_amp, i, sizeof (cl_mem), device_param->kernel_params_amp[i]); + CL_err = hc_clSetKernelArg (data.ocl, device_param->kernel_amp, i, sizeof (cl_mem), device_param->kernel_params_amp[i]); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clSetKernelArg(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } } for (uint i = 5; i < 7; i++) { - hc_clSetKernelArg (data.ocl, device_param->kernel_amp, i, sizeof (cl_uint), device_param->kernel_params_amp[i]); + CL_err = hc_clSetKernelArg (data.ocl, device_param->kernel_amp, i, sizeof (cl_uint), device_param->kernel_params_amp[i]); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clSetKernelArg(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } } } @@ -16385,7 +17573,7 @@ int main (int argc, char **argv) } else if (device_param->device_vendor_id == VENDOR_ID_NV) { - #ifdef LINUX + #ifdef __linux__ rc = set_fan_control (data.hm_xnvctrl, data.hm_device[device_id].xnvctrl, NV_CTRL_GPU_COOLER_MANUAL_CONTROL_TRUE); #endif @@ -16435,6 +17623,7 @@ int main (int argc, char **argv) char *hash_type = strhashtype (data.hash_mode); // not a bug log_info ("Hashtype: %s", hash_type); + log_info ("Mangle: %d", mangle); log_info (""); } } @@ -16456,21 +17645,21 @@ int main (int argc, char **argv) { log_error ("ERROR: %s: %s", "stdin", strerror (errno)); - return (-1); + return -1; } if (_setmode (_fileno (stdout), _O_BINARY) == -1) { log_error ("ERROR: %s: %s", "stdout", strerror (errno)); - return (-1); + return -1; } if (_setmode (_fileno (stderr), _O_BINARY) == -1) { log_error ("ERROR: %s: %s", "stderr", strerror (errno)); - return (-1); + return -1; } #endif @@ -16515,7 +17704,7 @@ int main (int argc, char **argv) { log_error ("ERROR: %s: %s", l0_filename, strerror (errno)); - return (-1); + return -1; } uint is_dir = S_ISDIR (l0_stat.st_mode); @@ -16536,7 +17725,7 @@ int main (int argc, char **argv) { log_error ("ERROR: Keyspace parameter is not allowed together with a directory"); - return (-1); + return -1; } char **dictionary_files = NULL; @@ -16557,7 +17746,7 @@ int main (int argc, char **argv) { log_error ("ERROR: %s: %s", l1_filename, strerror (errno)); - return (-1); + return -1; } if (S_ISREG (l1_stat.st_mode)) @@ -16579,7 +17768,7 @@ int main (int argc, char **argv) { log_error ("ERROR: No usable dictionary file found."); - return (-1); + return -1; } } else if (wordlist_mode == WL_MODE_STDIN) @@ -16605,7 +17794,7 @@ int main (int argc, char **argv) { log_error ("ERROR: %s: %s", dictfile1, strerror (errno)); - return (-1); + return -1; } if (stat (dictfile1, &tmp_stat) == -1) @@ -16614,7 +17803,7 @@ int main (int argc, char **argv) fclose (fp1); - return (-1); + return -1; } if (S_ISDIR (tmp_stat.st_mode)) @@ -16623,7 +17812,7 @@ int main (int argc, char **argv) fclose (fp1); - return (-1); + return -1; } if ((fp2 = fopen (dictfile2, "rb")) == NULL) @@ -16632,7 +17821,7 @@ int main (int argc, char **argv) fclose (fp1); - return (-1); + return -1; } if (stat (dictfile2, &tmp_stat) == -1) @@ -16642,7 +17831,7 @@ int main (int argc, char **argv) fclose (fp1); fclose (fp2); - return (-1); + return -1; } if (S_ISDIR (tmp_stat.st_mode)) @@ -16652,7 +17841,7 @@ int main (int argc, char **argv) fclose (fp1); fclose (fp2); - return (-1); + return -1; } data.combs_cnt = 1; @@ -16670,7 +17859,7 @@ int main (int argc, char **argv) fclose (fp1); fclose (fp2); - return (-1); + return -1; } data.combs_cnt = 1; @@ -16688,7 +17877,7 @@ int main (int argc, char **argv) fclose (fp1); fclose (fp2); - return (-1); + return -1; } fclose (fp1); @@ -16766,7 +17955,7 @@ int main (int argc, char **argv) { log_error ("ERROR: %s: %s", mask, strerror (errno)); - return (-1); + return -1; } } @@ -16780,7 +17969,7 @@ int main (int argc, char **argv) { log_error ("ERROR: %s: %s", mask, strerror (errno)); - return (-1); + return -1; } char *line_buf = (char *) mymalloc (HCBUFSIZ); @@ -16815,7 +18004,7 @@ int main (int argc, char **argv) { log_error ("ERROR: %s: unsupported file-type", mask); - return (-1); + return -1; } } @@ -16921,7 +18110,7 @@ int main (int argc, char **argv) { log_error ("ERROR: %s: %s", mask, strerror (errno)); - return (-1); + return -1; } char *line_buf = (char *) mymalloc (HCBUFSIZ); @@ -16978,7 +18167,7 @@ int main (int argc, char **argv) { log_error ("ERROR: %s: %s", filename, strerror (errno)); - return (-1); + return -1; } uint is_dir = S_ISDIR (file_stat.st_mode); @@ -16999,7 +18188,7 @@ int main (int argc, char **argv) { log_error ("ERROR: Keyspace parameter is not allowed together with a directory"); - return (-1); + return -1; } char **dictionary_files = NULL; @@ -17020,7 +18209,7 @@ int main (int argc, char **argv) { log_error ("ERROR: %s: %s", l1_filename, strerror (errno)); - return (-1); + return -1; } if (S_ISREG (l1_stat.st_mode)) @@ -17042,7 +18231,7 @@ int main (int argc, char **argv) { log_error ("ERROR: No usable dictionary file found."); - return (-1); + return -1; } if (increment) @@ -17100,7 +18289,7 @@ int main (int argc, char **argv) { log_error ("ERROR: %s: %s", mask, strerror (errno)); - return (-1); + return -1; } char *line_buf = (char *) mymalloc (HCBUFSIZ); @@ -17157,7 +18346,7 @@ int main (int argc, char **argv) { log_error ("ERROR: %s: %s", filename, strerror (errno)); - return (-1); + return -1; } uint is_dir = S_ISDIR (file_stat.st_mode); @@ -17178,7 +18367,7 @@ int main (int argc, char **argv) { log_error ("ERROR: Keyspace parameter is not allowed together with a directory"); - return (-1); + return -1; } char **dictionary_files = NULL; @@ -17199,7 +18388,7 @@ int main (int argc, char **argv) { log_error ("ERROR: %s: %s", l1_filename, strerror (errno)); - return (-1); + return -1; } if (S_ISREG (l1_stat.st_mode)) @@ -17221,7 +18410,7 @@ int main (int argc, char **argv) { log_error ("ERROR: No usable dictionary file found."); - return (-1); + return -1; } if (increment) @@ -17448,6 +18637,36 @@ int main (int argc, char **argv) mask = mask + str_pos + 1; } + + /** + * What follows is a very special case where "\," is within the mask field of a line in a .hcmask file only because otherwise (without the "\") + * it would be interpreted as a custom charset definition. + * + * We need to replace all "\," with just "," within the mask (but allow the special case "\\," which means "\" followed by ",") + * Note: "\\" is not needed to replace all "\" within the mask! The meaning of "\\" within a line containing the string "\\," is just to allow "\" followed by "," + */ + + uint mask_len_cur = strlen (mask); + + uint mask_out_pos = 0; + char mask_prev = 0; + + for (uint mask_iter = 0; mask_iter < mask_len_cur; mask_iter++, mask_out_pos++) + { + if (mask[mask_iter] == ',') + { + if (mask_prev == '\\') + { + mask_out_pos -= 1; // this means: skip the previous "\" + } + } + + mask_prev = mask[mask_iter]; + + mask[mask_out_pos] = mask[mask_iter]; + } + + mask[mask_out_pos] = '\0'; } if ((attack_mode == ATTACK_MODE_HYBRID1) || (attack_mode == ATTACK_MODE_HYBRID2)) @@ -17523,12 +18742,28 @@ int main (int argc, char **argv) device_param->kernel_params_mp_buf32[7] = 0; } - for (uint i = 0; i < 3; i++) hc_clSetKernelArg (data.ocl, device_param->kernel_mp, i, sizeof (cl_mem), (void *) device_param->kernel_params_mp[i]); - for (uint i = 3; i < 4; i++) hc_clSetKernelArg (data.ocl, device_param->kernel_mp, i, sizeof (cl_ulong), (void *) device_param->kernel_params_mp[i]); - for (uint i = 4; i < 8; i++) hc_clSetKernelArg (data.ocl, device_param->kernel_mp, i, sizeof (cl_uint), (void *) device_param->kernel_params_mp[i]); + cl_int CL_err = CL_SUCCESS; - hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_root_css_buf, CL_TRUE, 0, device_param->size_root_css, root_css_buf, 0, NULL, NULL); - hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_markov_css_buf, CL_TRUE, 0, device_param->size_markov_css, markov_css_buf, 0, NULL, NULL); + for (uint i = 0; i < 3; i++) CL_err |= hc_clSetKernelArg (data.ocl, device_param->kernel_mp, i, sizeof (cl_mem), (void *) device_param->kernel_params_mp[i]); + for (uint i = 3; i < 4; i++) CL_err |= hc_clSetKernelArg (data.ocl, device_param->kernel_mp, i, sizeof (cl_ulong), (void *) device_param->kernel_params_mp[i]); + for (uint i = 4; i < 8; i++) CL_err |= hc_clSetKernelArg (data.ocl, device_param->kernel_mp, i, sizeof (cl_uint), (void *) device_param->kernel_params_mp[i]); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clSetKernelArg(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } + + CL_err |= hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_root_css_buf, CL_TRUE, 0, device_param->size_root_css, root_css_buf, 0, NULL, NULL); + CL_err |= hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_markov_css_buf, CL_TRUE, 0, device_param->size_markov_css, markov_css_buf, 0, NULL, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clEnqueueWriteBuffer(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } } } else if (attack_mode == ATTACK_MODE_BF) @@ -17564,7 +18799,7 @@ int main (int argc, char **argv) { log_error ("ERROR: Mask is too small"); - return (-1); + return -1; } } } @@ -17597,7 +18832,7 @@ int main (int argc, char **argv) { log_error ("ERROR: --keyspace is not supported with --increment or mask files"); - return (-1); + return -1; } } @@ -17709,7 +18944,7 @@ int main (int argc, char **argv) { log_error ("ERROR: %s: %s", dictfile, strerror (errno)); - return (-1); + return -1; } data.words_cnt = count_words (wl_data, fd2, dictfile, dictstat_base, &dictstat_nmemb); @@ -17740,7 +18975,7 @@ int main (int argc, char **argv) { log_error ("ERROR: %s: %s", dictfile, strerror (errno)); - return (-1); + return -1; } data.words_cnt = count_words (wl_data, fd2, dictfile, dictstat_base, &dictstat_nmemb); @@ -17755,7 +18990,7 @@ int main (int argc, char **argv) { log_error ("ERROR: %s: %s", dictfile2, strerror (errno)); - return (-1); + return -1; } data.words_cnt = count_words (wl_data, fd2, dictfile2, dictstat_base, &dictstat_nmemb); @@ -17796,7 +19031,7 @@ int main (int argc, char **argv) { log_error ("ERROR: %s: %s", dictfile, strerror (errno)); - return (-1); + return -1; } data.words_cnt = count_words (wl_data, fd2, dictfile, dictstat_base, &dictstat_nmemb); @@ -18027,16 +19262,32 @@ int main (int argc, char **argv) device_param->kernel_params_mp_r_buf32[6] = 0; device_param->kernel_params_mp_r_buf32[7] = 0; - for (uint i = 0; i < 3; i++) hc_clSetKernelArg (data.ocl, device_param->kernel_mp_l, i, sizeof (cl_mem), (void *) device_param->kernel_params_mp_l[i]); - for (uint i = 3; i < 4; i++) hc_clSetKernelArg (data.ocl, device_param->kernel_mp_l, i, sizeof (cl_ulong), (void *) device_param->kernel_params_mp_l[i]); - for (uint i = 4; i < 9; i++) hc_clSetKernelArg (data.ocl, device_param->kernel_mp_l, i, sizeof (cl_uint), (void *) device_param->kernel_params_mp_l[i]); + cl_int CL_err = CL_SUCCESS; + + for (uint i = 0; i < 3; i++) CL_err |= hc_clSetKernelArg (data.ocl, device_param->kernel_mp_l, i, sizeof (cl_mem), (void *) device_param->kernel_params_mp_l[i]); + for (uint i = 3; i < 4; i++) CL_err |= hc_clSetKernelArg (data.ocl, device_param->kernel_mp_l, i, sizeof (cl_ulong), (void *) device_param->kernel_params_mp_l[i]); + for (uint i = 4; i < 9; i++) CL_err |= hc_clSetKernelArg (data.ocl, device_param->kernel_mp_l, i, sizeof (cl_uint), (void *) device_param->kernel_params_mp_l[i]); + + for (uint i = 0; i < 3; i++) CL_err |= hc_clSetKernelArg (data.ocl, device_param->kernel_mp_r, i, sizeof (cl_mem), (void *) device_param->kernel_params_mp_r[i]); + for (uint i = 3; i < 4; i++) CL_err |= hc_clSetKernelArg (data.ocl, device_param->kernel_mp_r, i, sizeof (cl_ulong), (void *) device_param->kernel_params_mp_r[i]); + for (uint i = 4; i < 8; i++) CL_err |= hc_clSetKernelArg (data.ocl, device_param->kernel_mp_r, i, sizeof (cl_uint), (void *) device_param->kernel_params_mp_r[i]); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clSetKernelArg(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } + + CL_err |= hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_root_css_buf, CL_TRUE, 0, device_param->size_root_css, root_css_buf, 0, NULL, NULL); + CL_err |= hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_markov_css_buf, CL_TRUE, 0, device_param->size_markov_css, markov_css_buf, 0, NULL, NULL); - for (uint i = 0; i < 3; i++) hc_clSetKernelArg (data.ocl, device_param->kernel_mp_r, i, sizeof (cl_mem), (void *) device_param->kernel_params_mp_r[i]); - for (uint i = 3; i < 4; i++) hc_clSetKernelArg (data.ocl, device_param->kernel_mp_r, i, sizeof (cl_ulong), (void *) device_param->kernel_params_mp_r[i]); - for (uint i = 4; i < 8; i++) hc_clSetKernelArg (data.ocl, device_param->kernel_mp_r, i, sizeof (cl_uint), (void *) device_param->kernel_params_mp_r[i]); + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clEnqueueWriteBuffer(): %s\n", val2cstr_cl (CL_err)); - hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_root_css_buf, CL_TRUE, 0, device_param->size_root_css, root_css_buf, 0, NULL, NULL); - hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_markov_css_buf, CL_TRUE, 0, device_param->size_markov_css, markov_css_buf, 0, NULL, NULL); + return -1; + } } } @@ -18070,14 +19321,14 @@ int main (int argc, char **argv) { log_info ("%llu", (unsigned long long int) words_base); - return (0); + return 0; } if (data.words_cur > data.words_base) { log_error ("ERROR: Restore value greater keyspace"); - return (-1); + return -1; } if (data.words_cur) @@ -18237,6 +19488,8 @@ int main (int argc, char **argv) data.runtime_start = runtime_start; + data.prepare_time += runtime_start - prepare_start; + for (uint device_id = 0; device_id < data.devices_cnt; device_id++) { hc_device_param_t *device_param = &devices_param[device_id]; @@ -18311,6 +19564,10 @@ int main (int argc, char **argv) if (induction_dictionaries_cnt) { qsort (induction_dictionaries, induction_dictionaries_cnt, sizeof (char *), sort_by_mtime); + + // yeah, this next statement is a little hack to make sure that --loopback runs correctly (because with it we guarantee that the loop iterates one more time) + + dictpos--; } time_t runtime_stop; @@ -18322,6 +19579,8 @@ int main (int argc, char **argv) logfile_sub_uint (runtime_start); logfile_sub_uint (runtime_stop); + time (&prepare_start); + logfile_sub_msg ("STOP"); global_free (subid); @@ -18448,6 +19707,8 @@ int main (int argc, char **argv) if (device_param->skipped) continue; + cl_int CL_err = CL_SUCCESS; + local_free (device_param->combs_buf); local_free (device_param->hooks_buf); local_free (device_param->device_name); @@ -18455,57 +19716,94 @@ int main (int argc, char **argv) local_free (device_param->device_version); local_free (device_param->driver_version); - if (device_param->pws_buf) myfree (device_param->pws_buf); - if (device_param->d_pws_buf) hc_clReleaseMemObject (data.ocl, device_param->d_pws_buf); - if (device_param->d_pws_amp_buf) hc_clReleaseMemObject (data.ocl, device_param->d_pws_amp_buf); - if (device_param->d_rules) hc_clReleaseMemObject (data.ocl, device_param->d_rules); - if (device_param->d_rules_c) hc_clReleaseMemObject (data.ocl, device_param->d_rules_c); - if (device_param->d_combs) hc_clReleaseMemObject (data.ocl, device_param->d_combs); - if (device_param->d_combs_c) hc_clReleaseMemObject (data.ocl, device_param->d_combs_c); - if (device_param->d_bfs) hc_clReleaseMemObject (data.ocl, device_param->d_bfs); - if (device_param->d_bfs_c) hc_clReleaseMemObject (data.ocl, device_param->d_bfs_c); - if (device_param->d_bitmap_s1_a) hc_clReleaseMemObject (data.ocl, device_param->d_bitmap_s1_a); - if (device_param->d_bitmap_s1_b) hc_clReleaseMemObject (data.ocl, device_param->d_bitmap_s1_b); - if (device_param->d_bitmap_s1_c) hc_clReleaseMemObject (data.ocl, device_param->d_bitmap_s1_c); - if (device_param->d_bitmap_s1_d) hc_clReleaseMemObject (data.ocl, device_param->d_bitmap_s1_d); - if (device_param->d_bitmap_s2_a) hc_clReleaseMemObject (data.ocl, device_param->d_bitmap_s2_a); - if (device_param->d_bitmap_s2_b) hc_clReleaseMemObject (data.ocl, device_param->d_bitmap_s2_b); - if (device_param->d_bitmap_s2_c) hc_clReleaseMemObject (data.ocl, device_param->d_bitmap_s2_c); - if (device_param->d_bitmap_s2_d) hc_clReleaseMemObject (data.ocl, device_param->d_bitmap_s2_d); - if (device_param->d_plain_bufs) hc_clReleaseMemObject (data.ocl, device_param->d_plain_bufs); - if (device_param->d_digests_buf) hc_clReleaseMemObject (data.ocl, device_param->d_digests_buf); - if (device_param->d_digests_shown) hc_clReleaseMemObject (data.ocl, device_param->d_digests_shown); - if (device_param->d_salt_bufs) hc_clReleaseMemObject (data.ocl, device_param->d_salt_bufs); - if (device_param->d_esalt_bufs) hc_clReleaseMemObject (data.ocl, device_param->d_esalt_bufs); - if (device_param->d_tmps) hc_clReleaseMemObject (data.ocl, device_param->d_tmps); - if (device_param->d_hooks) hc_clReleaseMemObject (data.ocl, device_param->d_hooks); - if (device_param->d_result) hc_clReleaseMemObject (data.ocl, device_param->d_result); - if (device_param->d_scryptV0_buf) hc_clReleaseMemObject (data.ocl, device_param->d_scryptV0_buf); - if (device_param->d_scryptV1_buf) hc_clReleaseMemObject (data.ocl, device_param->d_scryptV1_buf); - if (device_param->d_scryptV2_buf) hc_clReleaseMemObject (data.ocl, device_param->d_scryptV2_buf); - if (device_param->d_scryptV3_buf) hc_clReleaseMemObject (data.ocl, device_param->d_scryptV3_buf); - if (device_param->d_root_css_buf) hc_clReleaseMemObject (data.ocl, device_param->d_root_css_buf); - if (device_param->d_markov_css_buf) hc_clReleaseMemObject (data.ocl, device_param->d_markov_css_buf); - if (device_param->d_tm_c) hc_clReleaseMemObject (data.ocl, device_param->d_tm_c); - - if (device_param->kernel1) hc_clReleaseKernel (data.ocl, device_param->kernel1); - if (device_param->kernel12) hc_clReleaseKernel (data.ocl, device_param->kernel12); - if (device_param->kernel2) hc_clReleaseKernel (data.ocl, device_param->kernel2); - if (device_param->kernel23) hc_clReleaseKernel (data.ocl, device_param->kernel23); - if (device_param->kernel3) hc_clReleaseKernel (data.ocl, device_param->kernel3); - if (device_param->kernel_mp) hc_clReleaseKernel (data.ocl, device_param->kernel_mp); - if (device_param->kernel_mp_l) hc_clReleaseKernel (data.ocl, device_param->kernel_mp_l); - if (device_param->kernel_mp_r) hc_clReleaseKernel (data.ocl, device_param->kernel_mp_r); - if (device_param->kernel_tm) hc_clReleaseKernel (data.ocl, device_param->kernel_tm); - if (device_param->kernel_amp) hc_clReleaseKernel (data.ocl, device_param->kernel_amp); - if (device_param->kernel_memset) hc_clReleaseKernel (data.ocl, device_param->kernel_memset); - - if (device_param->program) hc_clReleaseProgram (data.ocl, device_param->program); - if (device_param->program_mp) hc_clReleaseProgram (data.ocl, device_param->program_mp); - if (device_param->program_amp) hc_clReleaseProgram (data.ocl, device_param->program_amp); - - if (device_param->command_queue) hc_clReleaseCommandQueue (data.ocl, device_param->command_queue); - if (device_param->context) hc_clReleaseContext (data.ocl, device_param->context); + if (device_param->pws_buf) myfree (device_param->pws_buf); + + if (device_param->d_pws_buf) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_pws_buf); + if (device_param->d_pws_amp_buf) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_pws_amp_buf); + if (device_param->d_rules) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_rules); + if (device_param->d_rules_c) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_rules_c); + if (device_param->d_combs) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_combs); + if (device_param->d_combs_c) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_combs_c); + if (device_param->d_bfs) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_bfs); + if (device_param->d_bfs_c) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_bfs_c); + if (device_param->d_bitmap_s1_a) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_bitmap_s1_a); + if (device_param->d_bitmap_s1_b) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_bitmap_s1_b); + if (device_param->d_bitmap_s1_c) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_bitmap_s1_c); + if (device_param->d_bitmap_s1_d) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_bitmap_s1_d); + if (device_param->d_bitmap_s2_a) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_bitmap_s2_a); + if (device_param->d_bitmap_s2_b) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_bitmap_s2_b); + if (device_param->d_bitmap_s2_c) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_bitmap_s2_c); + if (device_param->d_bitmap_s2_d) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_bitmap_s2_d); + if (device_param->d_plain_bufs) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_plain_bufs); + if (device_param->d_digests_buf) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_digests_buf); + if (device_param->d_digests_shown) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_digests_shown); + if (device_param->d_salt_bufs) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_salt_bufs); + if (device_param->d_esalt_bufs) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_esalt_bufs); + if (device_param->d_tmps) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_tmps); + if (device_param->d_hooks) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_hooks); + if (device_param->d_result) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_result); + if (device_param->d_scryptV0_buf) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_scryptV0_buf); + if (device_param->d_scryptV1_buf) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_scryptV1_buf); + if (device_param->d_scryptV2_buf) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_scryptV2_buf); + if (device_param->d_scryptV3_buf) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_scryptV3_buf); + if (device_param->d_root_css_buf) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_root_css_buf); + if (device_param->d_markov_css_buf) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_markov_css_buf); + if (device_param->d_tm_c) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_tm_c); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clReleaseMemObject(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } + + if (device_param->kernel1) CL_err |= hc_clReleaseKernel (data.ocl, device_param->kernel1); + if (device_param->kernel12) CL_err |= hc_clReleaseKernel (data.ocl, device_param->kernel12); + if (device_param->kernel2) CL_err |= hc_clReleaseKernel (data.ocl, device_param->kernel2); + if (device_param->kernel23) CL_err |= hc_clReleaseKernel (data.ocl, device_param->kernel23); + if (device_param->kernel3) CL_err |= hc_clReleaseKernel (data.ocl, device_param->kernel3); + if (device_param->kernel_mp) CL_err |= hc_clReleaseKernel (data.ocl, device_param->kernel_mp); + if (device_param->kernel_mp_l) CL_err |= hc_clReleaseKernel (data.ocl, device_param->kernel_mp_l); + if (device_param->kernel_mp_r) CL_err |= hc_clReleaseKernel (data.ocl, device_param->kernel_mp_r); + if (device_param->kernel_tm) CL_err |= hc_clReleaseKernel (data.ocl, device_param->kernel_tm); + if (device_param->kernel_amp) CL_err |= hc_clReleaseKernel (data.ocl, device_param->kernel_amp); + if (device_param->kernel_memset) CL_err |= hc_clReleaseKernel (data.ocl, device_param->kernel_memset); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clReleaseKernel(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } + + if (device_param->program) CL_err |= hc_clReleaseProgram (data.ocl, device_param->program); + if (device_param->program_mp) CL_err |= hc_clReleaseProgram (data.ocl, device_param->program_mp); + if (device_param->program_amp) CL_err |= hc_clReleaseProgram (data.ocl, device_param->program_amp); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clReleaseProgram(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } + + if (device_param->command_queue) CL_err |= hc_clReleaseCommandQueue (data.ocl, device_param->command_queue); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clReleaseCommandQueue(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } + + if (device_param->context) CL_err |= hc_clReleaseContext (data.ocl, device_param->context); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: hc_clReleaseContext(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } } // reset default fan speed @@ -18533,7 +19831,7 @@ int main (int argc, char **argv) } else if (device_param->device_vendor_id == VENDOR_ID_NV) { - #ifdef LINUX + #ifdef __linux__ rc = set_fan_control (data.hm_xnvctrl, data.hm_device[device_id].xnvctrl, NV_CTRL_GPU_COOLER_MANUAL_CONTROL_FALSE); #endif @@ -18574,7 +19872,7 @@ int main (int argc, char **argv) { log_error ("ERROR: Failed to get ADL PowerControl Capabilities"); - return (-1); + return -1; } if (powertune_supported != 0) @@ -18585,7 +19883,7 @@ int main (int argc, char **argv) { log_info ("ERROR: Failed to restore the ADL PowerControl values"); - return (-1); + return -1; } // clocks @@ -18603,7 +19901,7 @@ int main (int argc, char **argv) { log_info ("ERROR: Failed to restore ADL performance state"); - return (-1); + return -1; } local_free (performance_state); @@ -18788,7 +20086,7 @@ int main (int argc, char **argv) { log_error ("ERROR: %s: %s", induction_directory, strerror (errno)); - return (-1); + return -1; } } @@ -18814,7 +20112,7 @@ int main (int argc, char **argv) { log_error ("ERROR: %s: %s", outfile_check_directory, strerror (errno)); - return (-1); + return -1; } }