- /**
+/**
* Authors.....: Jens Steube <jens.steube@gmail.com>
* Gabriele Gristina <matrix@hashcat.net>
* magnum <john.magnum@hushmail.com>
* License.....: MIT
*/
-#ifdef OSX
+#ifdef __APPLE__
+#include <stdio.h>
+#endif
+
+#ifdef __FreeBSD__
#include <stdio.h>
#endif
#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
" 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",
* flush
*/
- #ifdef _WIN
- fputc ('\r', out);
- fputc ('\n', out);
- #endif
-
- #ifdef _POSIX
- fputc ('\n', out);
- #endif
-
+ fputs (EOL, out);
fflush (out);
}
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 ();
}
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));
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 - runtime_cur;
+
+ tmp = _gmtime64 (&runtime_left);
+
+ #else
+
+ time_t runtime_left = data.proc_start + data.runtime - 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);
+ }
}
}
}
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 ();
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)
}
}
-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)
{
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;
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 ;)
uint plain_len = 0;
- const uint il_cnt = device_param->kernel_params_buf32[27]; // ugly, i know
+ const uint il_cnt = device_param->kernel_params_buf32[30]; // ugly, i know
if (data.attack_mode == ATTACK_MODE_STRAIGHT)
{
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[30] = data.combs_mode;
- device_param->kernel_params_buf32[31] = num;
+ device_param->kernel_params_buf32[33] = data.combs_mode;
+ device_param->kernel_params_buf32[34] = num;
uint kernel_threads = device_param->kernel_threads;
case KERN_RUN_3: kernel = device_param->kernel3; break;
}
- hc_clSetKernelArg (data.ocl, kernel, 21, sizeof (cl_uint), device_param->kernel_params[21]);
- hc_clSetKernelArg (data.ocl, kernel, 22, sizeof (cl_uint), device_param->kernel_params[22]);
- hc_clSetKernelArg (data.ocl, kernel, 23, sizeof (cl_uint), device_param->kernel_params[23]);
- 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]);
+ 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;
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
{
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)
{
}
}
- 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;
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)
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;
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;
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);
- 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));
+
+ 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;
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;
-
- const size_t fillsz = MIN (FILLSZ, left);
+ log_error ("ERROR: clEnqueueWriteBuffer(): %s\n", val2cstr_cl (CL_err));
- 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)
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;
+ }
}
}
{
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;
loop_left = MIN (loop_left, loop_step);
- device_param->kernel_params_buf32[25] = loop_pos;
- device_param->kernel_params_buf32[26] = loop_left;
+ device_param->kernel_params_buf32[28] = loop_pos;
+ device_param->kernel_params_buf32[29] = loop_left;
run_kernel (KERN_RUN_2, device_param, pws_cnt, true, slow_iteration);
if (data.devices_status == STATUS_CRACKED) break;
if (data.devices_status == STATUS_ABORTED) break;
if (data.devices_status == STATUS_QUIT) break;
+ if (data.devices_status == STATUS_BYPASS) break;
/**
* speed
{
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)
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)
{
}
}
- 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)
{
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)
{
const u32 kernel_power_try = device_param->device_processors * device_param->kernel_threads * kernel_accel;
- device_param->kernel_params_buf32[25] = 0;
- device_param->kernel_params_buf32[26] = kernel_loops; // not a bug, both need to be set
- device_param->kernel_params_buf32[27] = kernel_loops; // because there's two variables for inner iters for slow and fast hashes
+ device_param->kernel_params_buf32[28] = 0;
+ device_param->kernel_params_buf32[29] = kernel_loops; // not a bug, both need to be set
+ device_param->kernel_params_buf32[30] = kernel_loops; // because there's two variables for inner iters for slow and fast hashes
if (data.attack_exec == ATTACK_EXEC_INSIDE_KERNEL)
{
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];
device_param->kernel_power = kernel_power;
- return;
+ return 0;
}
// from here it's clear we are allowed to autotune
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
}
#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);
salt_t *salt_buf = &data.salts_buf[salt_pos];
- device_param->kernel_params_buf32[24] = salt_pos;
- device_param->kernel_params_buf32[28] = salt_buf->digests_cnt;
- device_param->kernel_params_buf32[29] = salt_buf->digests_offset;
+ device_param->kernel_params_buf32[27] = salt_pos;
+ device_param->kernel_params_buf32[31] = salt_buf->digests_cnt;
+ device_param->kernel_params_buf32[32] = salt_buf->digests_offset;
FILE *combs_fp = device_param->combs_fp;
device_param->innerloop_pos = innerloop_pos;
device_param->innerloop_left = innerloop_left;
- device_param->kernel_params_buf32[27] = innerloop_left;
+ device_param->kernel_params_buf32[30] = innerloop_left;
// i think we can get rid of this
if (innerloop_left == 0)
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);
- }
- else if (data.attack_mode == ATTACK_MODE_BF)
+ 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)
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 ();
if (data.devices_status == STATUS_CRACKED) break;
if (data.devices_status == STATUS_ABORTED) break;
if (data.devices_status == STATUS_QUIT) break;
+ if (data.devices_status == STATUS_BYPASS) break;
/**
* result
device_param->speed_pos = speed_pos;
myfree (line_buf);
+
+ return 0;
}
static void load_segment (wl_data_t *wl_data, FILE *fd)
return (p);
}
- while ((data.devices_status != STATUS_EXHAUSTED) && (data.devices_status != STATUS_CRACKED) && (data.devices_status != STATUS_ABORTED) && (data.devices_status != STATUS_QUIT))
+ while (data.shutdown_inner == 0)
{
hc_sleep (sleep_time);
{
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);
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
}
uint check_left = outfile_check_timer; // or 1 if we want to check it at startup
- while ((data.devices_status != STATUS_EXHAUSTED) && (data.devices_status != STATUS_CRACKED) && (data.devices_status != STATUS_ABORTED) && (data.devices_status != STATUS_QUIT))
+ while (data.shutdown_inner == 0)
{
hc_sleep (1);
const uint attack_kern = data.attack_kern;
- while ((data.devices_status != STATUS_EXHAUSTED) && (data.devices_status != STATUS_CRACKED) && (data.devices_status != STATUS_ABORTED) && (data.devices_status != STATUS_QUIT))
+ while ((data.devices_status != STATUS_CRACKED) && (data.devices_status != STATUS_ABORTED) && (data.devices_status != STATUS_QUIT))
{
hc_thread_mutex_lock (mux_dispatcher);
if (attack_mode == ATTACK_MODE_BF)
{
- while ((data.devices_status != STATUS_EXHAUSTED) && (data.devices_status != STATUS_CRACKED) && (data.devices_status != STATUS_ABORTED) && (data.devices_status != STATUS_QUIT))
+ while ((data.devices_status != STATUS_CRACKED) && (data.devices_status != STATUS_ABORTED) && (data.devices_status != STATUS_QUIT))
{
const uint work = get_work (device_param, -1);
u64 words_cur = 0;
- while ((data.devices_status != STATUS_EXHAUSTED) && (data.devices_status != STATUS_CRACKED) && (data.devices_status != STATUS_ABORTED) && (data.devices_status != STATUS_QUIT))
+ while ((data.devices_status != STATUS_CRACKED) && (data.devices_status != STATUS_ABORTED) && (data.devices_status != STATUS_QUIT))
{
u64 words_off = 0;
u64 words_fin = 0;
salt_t *salt_buf = &data.salts_buf[salt_pos];
- device_param->kernel_params_buf32[24] = salt_pos;
- device_param->kernel_params_buf32[27] = 1;
- device_param->kernel_params_buf32[28] = salt_buf->digests_cnt;
- device_param->kernel_params_buf32[29] = salt_buf->digests_offset;
- device_param->kernel_params_buf32[30] = 0;
- device_param->kernel_params_buf32[31] = 1;
+ device_param->kernel_params_buf32[27] = salt_pos;
+ device_param->kernel_params_buf32[30] = 1;
+ device_param->kernel_params_buf32[31] = salt_buf->digests_cnt;
+ device_param->kernel_params_buf32[32] = salt_buf->digests_offset;
+ device_param->kernel_params_buf32[33] = 0;
+ device_param->kernel_params_buf32[34] = 1;
char *dictfile_old = data.dictfile;
loop_left = MIN (loop_left, loop_step);
- device_param->kernel_params_buf32[25] = loop_pos;
- device_param->kernel_params_buf32[26] = loop_left;
+ device_param->kernel_params_buf32[28] = loop_pos;
+ device_param->kernel_params_buf32[29] = loop_left;
run_kernel (KERN_RUN_2, device_param, 1, false, 0);
}
* cleanup
*/
- device_param->kernel_params_buf32[24] = 0;
- device_param->kernel_params_buf32[25] = 0;
- device_param->kernel_params_buf32[26] = 0;
device_param->kernel_params_buf32[27] = 0;
device_param->kernel_params_buf32[28] = 0;
device_param->kernel_params_buf32[29] = 0;
device_param->kernel_params_buf32[30] = 0;
device_param->kernel_params_buf32[31] = 0;
+ device_param->kernel_params_buf32[32] = 0;
+ device_param->kernel_params_buf32[33] = 0;
+ device_param->kernel_params_buf32[34] = 0;
data.dictfile = dictfile_old;
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
*/
case IDX_SESSION: session = optarg; break;
case IDX_SHOW: show = 1; break;
case IDX_LEFT: left = 1; break;
- case '?': return (-1);
+ case '?': return -1;
}
}
{
log_error ("ERROR: Invalid argument specified");
- return (-1);
+ return -1;
}
/**
{
log_info ("%s", VERSION_TAG);
- return (0);
+ return 0;
}
if (usage)
{
usage_big_print (PROGNAME);
- return (0);
+ return 0;
}
/**
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;
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)
{
log_error ("ERROR: Incompatible restore-file version");
- return (-1);
+ return -1;
}
myargc = rd->argc;
default:
log_error ("ERROR: Invalid argument specified");
- return (-1);
+ return -1;
}
}
{
log_error ("ERROR: Invalid argument specified");
- return (-1);
+ return -1;
}
/**
{
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
{
log_error ("Old -m specified, use -m %d instead", n);
- return (-1);
+ return -1;
}
}
{
log_error ("ERROR: Mixing support for user names and hashes of type %s is not supported", strhashtype (hash_mode));
- return (-1);
+ return -1;
}
}
{
log_error ("ERROR: Invalid outfile-format specified");
- return (-1);
+ return -1;
}
if (left == 1)
{
log_error ("ERROR: Mixing outfile-format > 1 with left parameter is not allowed");
- return (-1);
+ return -1;
}
}
else
{
log_error ("ERROR: Mixing outfile-format > 7 with show parameter is not allowed");
- return (-1);
+ return -1;
}
}
}
{
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)
{
log_error ("ERROR: Use of rules-file or rules-generate only allowed in attack-mode 0");
- return (-1);
+ return -1;
}
}
{
log_error ("ERROR: Invalid rp-gen-func-min specified");
- return (-1);
+ return -1;
}
if (kernel_accel_chgd == 1)
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;
}
}
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;
}
}
{
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)
{
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;
}
}
{
log_error ("ERROR: Invalid argument for benchmark mode specified");
- return (-1);
+ return -1;
}
if (attack_mode_chgd == 1)
{
log_error ("ERROR: Only attack-mode 3 allowed in benchmark mode");
- return (-1);
+ return -1;
}
}
}
{
usage_mini_print (myargv[0]);
- return (-1);
+ return -1;
}
}
else if (attack_kern == ATTACK_KERN_STRAIGHT)
{
usage_mini_print (myargv[0]);
- return (-1);
+ return -1;
}
}
else if (attack_kern == ATTACK_KERN_COMBI)
{
usage_mini_print (myargv[0]);
- return (-1);
+ return -1;
}
}
else if (attack_kern == ATTACK_KERN_BF)
{
usage_mini_print (myargv[0]);
- return (-1);
+ return -1;
}
}
else
{
usage_mini_print (myargv[0]);
- return (-1);
+ return -1;
}
}
{
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;
{
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;
}
}
{
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;
}
}
{
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;
}
}
{
log_error ("ERROR: Invalid debug-mode specified");
- return (-1);
+ return -1;
}
if (debug_file != NULL)
{
log_error ("ERROR: Parameter debug-file requires parameter debug-mode to be set");
- return (-1);
+ return -1;
}
}
{
log_error ("ERROR: Parameter induction-dir not allowed with brute-force attacks");
- return (-1);
+ return -1;
}
}
{
log_error ("ERROR: setting --weak-hash-threshold allowed only in straight-attack mode");
- return (-1);
+ return -1;
}
weak_hash_threshold = 0;
{
log_error ("ERROR: setting --nvidia-spin-damp must be between 0 and 100 (inclusive)");
- return (-1);
+ return -1;
}
{
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;
}
}
{
log_error ("ERROR: %s: %s", induction_directory, strerror (errno));
- return (-1);
+ return -1;
}
}
}
{
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)
{
log_error ("ERROR: %s: %s", outfile_check_directory, strerror (errno));
- return (-1);
+ return -1;
}
}
}
hc_thread_t *outer_threads = (hc_thread_t *) mycalloc (10, sizeof (hc_thread_t));
+ data.shutdown_outer = 0;
+
if (keyspace == 0 && benchmark == 0 && stdout_flag == 0)
{
if ((data.wordlist_mode == WL_MODE_FILE) || (data.wordlist_mode == WL_MODE_MASK))
outer_threads_cnt++;
}
-
- hc_thread_create (outer_threads[outer_threads_cnt], thread_monitor, NULL);
-
- outer_threads_cnt++;
}
/**
dgst_pos3 = 6;
break;
- default: usage_mini_print (PROGNAME); return (-1);
+ default: usage_mini_print (PROGNAME); return -1;
}
/**
{
log_error ("ERROR: Parameter hex-salt not valid for hash-type %u", hash_mode);
- return (-1);
+ return -1;
}
}
{
log_error ("ERROR: %s: %s", potfile, strerror (errno));
- return (-1);
+ return -1;
}
if (outfile != NULL)
fclose (pot_fp);
- return (-1);
+ return -1;
}
}
else
{
log_error ("ERROR: %s: %s", potfile, strerror (errno));
- return (-1);
+ return -1;
}
data.pot_fp = pot_fp;
{
log_error ("ERROR: %s: %s", data.hashfile, strerror (errno));
- return (-1);
+ return -1;
}
hashes_avail = st.st_size / sizeof (hccap_t);
{
log_error ("ERROR: %s: %s", hashfile, strerror (errno));
- return (-1);
+ return -1;
}
if (data.quiet == 0) log_info_nn ("Counting lines in %s", hashfile);
fclose (fp);
- return (-1);
+ return -1;
}
hashlist_format = hlfmt_detect (fp, 100); // 100 = max numbers to "scan". could be hashes_avail, too
fclose (fp);
- return (-1);
+ return -1;
}
fclose (fp);
{
log_error ("ERROR: hccap file not specified");
- return (-1);
+ return -1;
}
hashlist_mode = HL_MODE_FILE;
{
log_error ("ERROR: %s: %s", hash_buf, strerror (errno));
- return (-1);
+ return -1;
}
if (hashes_avail < 1)
fclose (fp);
- return (-1);
+ return -1;
}
uint hccap_size = sizeof (hccap_t);
{
log_error ("ERROR: %s: %s", hashfile, strerror (errno));
- return (-1);
+ return -1;
}
uint line_num = 0;
if (data.quiet == 0) log_info_nn ("");
- return (0);
+ return 0;
}
if ((keyspace == 0) && (stdout_flag == 0))
{
log_error ("ERROR: No hashes loaded");
- return (-1);
+ return -1;
}
}
{
log_error ("ERROR: Hashfile and Outfile are not allowed to point to the same file");
- return (-1);
+ return -1;
}
#endif
{
log_error ("ERROR: Hashfile and Outfile are not allowed to point to the same file");
- return (-1);
+ return -1;
}
#endif
}
// 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, ':');
{
log_error ("ERROR: %s: %s", rp_file, strerror (errno));
- return (-1);
+ return -1;
}
while (!feof (fp))
{
log_error ("ERROR: No valid rules left");
- return (-1);
+ return -1;
}
/**
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)
{
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)
{
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;
}
}
}
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++)
{
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;
}
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
platform_vendor_id = VENDOR_ID_GENERIC;
}
- const uint platform_skipped = ((opencl_platforms_filter & (1 << platform_id)) == 0);
+ uint platform_skipped = ((opencl_platforms_filter & (1 << 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))
{
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 ("");
+ }
}
}
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;
// 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;
// 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';
{
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
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;
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);
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;
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;
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;
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)
{
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)
{
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 (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: clGetDeviceInfo(): %s\n", val2cstr_cl (CL_err));
+
+ return -1;
+ }
if (device_compiler_available == CL_FALSE)
{
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)
{
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)
{
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)
{
// 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;
device_param->device_name_chksum = device_name_chksum;
- // device_processor_cores
+ // vendor specific
if (device_param->device_type & CL_DEVICE_TYPE_GPU)
{
{
need_nvml = 1;
- #ifdef LINUX
+ #ifdef __linux__
need_xnvctrl = 1;
#endif
}
}
- // device_processor_cores
-
- if (device_type & CL_DEVICE_TYPE_CPU)
- {
- cl_uint device_processor_cores = 1;
-
- device_param->device_processor_cores = device_processor_cores;
- }
-
if (device_type & CL_DEVICE_TYPE_GPU)
{
- if (device_vendor_id == VENDOR_ID_AMD)
- {
- cl_uint device_processor_cores = 0;
-
- #define CL_DEVICE_WAVEFRONT_WIDTH_AMD 0x4043
-
- hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_WAVEFRONT_WIDTH_AMD, sizeof (device_processor_cores), &device_processor_cores, NULL);
-
- device_param->device_processor_cores = device_processor_cores;
- }
- else if (device_vendor_id == VENDOR_ID_NV)
+ if (device_vendor_id == VENDOR_ID_NV)
{
cl_uint kernel_exec_timeout = 0;
#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);
-
- device_param->kernel_exec_timeout = kernel_exec_timeout;
-
- cl_uint device_processor_cores = 0;
+ CL_err = hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_KERNEL_EXEC_TIMEOUT_NV, sizeof (kernel_exec_timeout), &kernel_exec_timeout, NULL);
- #define CL_DEVICE_WARP_SIZE_NV 0x4003
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: clGetDeviceInfo(): %s\n", val2cstr_cl (CL_err));
- hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_WARP_SIZE_NV, sizeof (device_processor_cores), &device_processor_cores, NULL);
+ return -1;
+ }
- device_param->device_processor_cores = device_processor_cores;
+ device_param->kernel_exec_timeout = kernel_exec_timeout;
cl_uint sm_minor = 0;
cl_uint sm_major = 0;
#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);
- device_param->sm_minor = sm_minor;
- device_param->sm_major = sm_major;
+ 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;
// CPU burning loop damper
// Value is given as number between 0-100
device_param->nvidia_spin_damp /= 100;
}
- else
- {
- cl_uint device_processor_cores = 1;
-
- device_param->device_processor_cores = device_processor_cores;
- }
}
// display results
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)
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)
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;
}
}
}
{
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)
{
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;
}
}
*/
#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)
{
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
{
log_error ("ERROR: Invalid values for gpu-temp-abort. Parameter gpu-temp-abort is less than gpu-temp-retain.");
- return (-1);
+ return -1;
}
}
{
log_error ("ERROR: Failed to get ADL PowerControl Capabilities");
- return (-1);
+ return -1;
}
// first backup current value, we will restore it later
{
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
{
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
{
log_error ("ERROR: Failed to get ADL device capabilities");
- return (-1);
+ return -1;
}
int engine_clock_max = caps.sEngineClockRange.iMax * 0.6666;
{
log_info ("ERROR: Failed to set ADL performance state");
- return (-1);
+ return -1;
}
local_free (performance_state);
{
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;
}
}
}
for (uint device_id = 0; device_id < data.devices_cnt; device_id++)
{
+ cl_int CL_err = CL_SUCCESS;
+
/**
* host buffer
*/
const char *device_name_chksum = device_param->device_name_chksum;
const u32 device_processors = device_param->device_processors;
- const u32 device_processor_cores = device_param->device_processor_cores;
/**
* 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
// 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
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;
// scryptV stuff
- size_t size_scryptV = 1;
+ size_t size_scrypt = 4;
if ((hash_mode == 8900) || (hash_mode == 9300))
{
+ // we need to check that all hashes have the same scrypt settings
+
+ const u32 scrypt_N = data.salts_buf[0].scrypt_N;
+ const u32 scrypt_r = data.salts_buf[0].scrypt_r;
+ const u32 scrypt_p = data.salts_buf[0].scrypt_p;
+
+ for (uint i = 1; i < salts_cnt; i++)
+ {
+ if ((data.salts_buf[i].scrypt_N != scrypt_N)
+ || (data.salts_buf[i].scrypt_r != scrypt_r)
+ || (data.salts_buf[i].scrypt_p != scrypt_p))
+ {
+ log_error ("ERROR: Mixed scrypt settings not supported");
+
+ return -1;
+ }
+ }
+
uint tmto_start = 0;
uint tmto_stop = 10;
else
{
// in case the user did not specify the tmto manually
- // use some values known to run best (tested on 290x for AMD and 980ti for NV)
- // but set the lower end only in case the user has a device with too less memory
+ // use some values known to run best (tested on 290x for AMD and GTX1080 for NV)
if (hash_mode == 8900)
{
if (device_param->device_vendor_id == VENDOR_ID_AMD)
{
- tmto_start = 1;
+ tmto_start = 3;
}
else if (device_param->device_vendor_id == VENDOR_ID_NV)
{
}
else if (device_param->device_vendor_id == VENDOR_ID_NV)
{
- tmto_start = 2;
+ tmto_start = 4;
}
}
}
- for (uint tmto = tmto_start; tmto < tmto_stop; tmto++)
+ data.scrypt_tmp_size = (128 * scrypt_r * scrypt_p);
+
+ device_param->kernel_accel_min = 1;
+ device_param->kernel_accel_max = 8;
+
+ uint tmto;
+
+ for (tmto = tmto_start; tmto < tmto_stop; tmto++)
{
- // TODO: in theory the following calculation needs to be done per salt, not global
- // we assume all hashes have the same scrypt settings
+ size_scrypt = (128 * scrypt_r) * scrypt_N;
- size_scryptV = (128 * data.salts_buf[0].scrypt_r) * data.salts_buf[0].scrypt_N;
+ size_scrypt /= 1 << tmto;
- size_scryptV /= 1 << tmto;
+ size_scrypt *= device_param->device_processors * device_param->kernel_threads * device_param->kernel_accel_max;
- size_scryptV *= device_processors * device_processor_cores;
+ if ((size_scrypt / 4) > device_param->device_maxmem_alloc)
+ {
+ if (quiet == 0) log_info ("WARNING: Not enough single-block device memory allocatable to use --scrypt-tmto %d, increasing...", tmto);
+
+ continue;
+ }
- if (size_scryptV > device_param->device_maxmem_alloc)
+ if (size_scrypt > device_param->device_global_mem)
{
- if (quiet == 0) log_info ("WARNING: Not enough device memory allocatable to use --scrypt-tmto %d, increasing...", tmto);
+ if (quiet == 0) log_info ("WARNING: Not enough total device memory allocatable to use --scrypt-tmto %d, increasing...", tmto);
continue;
}
for (uint salts_pos = 0; salts_pos < data.salts_cnt; salts_pos++)
{
- data.salts_buf[salts_pos].scrypt_tmto = tmto;
- data.salts_buf[salts_pos].scrypt_phy = device_processors * device_processor_cores;
+ data.scrypt_tmto_final = tmto;
}
break;
}
- if (data.salts_buf[0].scrypt_phy == 0)
+ if (tmto == tmto_stop)
{
log_error ("ERROR: Can't allocate enough device memory");
return -1;
}
- if (quiet == 0) log_info ("SCRYPT tmto optimizer value set to: %u, mem: %u\n", data.salts_buf[0].scrypt_tmto, size_scryptV);
+ 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;
+
/**
* some algorithms need a fixed kernel-loops count
*/
case 7900: size_tmps = kernel_power_max * sizeof (drupal7_tmp_t); break;
case 8200: size_tmps = kernel_power_max * sizeof (pbkdf2_sha512_tmp_t); break;
case 8800: size_tmps = kernel_power_max * sizeof (androidfde_tmp_t); break;
- case 8900: size_tmps = kernel_power_max * sizeof (scrypt_tmp_t); break;
+ case 8900: size_tmps = kernel_power_max * data.scrypt_tmp_size; break;
case 9000: size_tmps = kernel_power_max * sizeof (pwsafe2_tmp_t); break;
case 9100: size_tmps = kernel_power_max * sizeof (lotus8_tmp_t); break;
case 9200: size_tmps = kernel_power_max * sizeof (pbkdf2_sha256_tmp_t); break;
- case 9300: size_tmps = kernel_power_max * sizeof (scrypt_tmp_t); break;
+ case 9300: size_tmps = kernel_power_max * data.scrypt_tmp_size; break;
case 9400: size_tmps = kernel_power_max * sizeof (office2007_tmp_t); break;
case 9500: size_tmps = kernel_power_max * sizeof (office2010_tmp_t); break;
case 9600: size_tmps = kernel_power_max * sizeof (office2013_tmp_t); break;
+ size_rules
+ size_rules_c
+ size_salts
- + size_scryptV
+ + size_scrypt4
+ + size_scrypt4
+ + size_scrypt4
+ + size_scrypt4
+ size_shown
+ size_tm
+ size_tmps;
snprintf (build_opts, sizeof (build_opts) - 1, "-I \"%s\"", cpath_real);
- myfree (cpath_real);
-
#else
snprintf (cpath, sizeof (cpath) - 1, "%s/OpenCL/", shared_dir);
snprintf (build_opts, sizeof (build_opts) - 1, "-I %s", cpath_real);
- myfree (cpath_real);
-
#endif
+ // include check
+ // this test needs to be done manually because of osx opencl runtime
+ // if there's a problem with permission, its not reporting back and erroring out silently
+
+ #define files_cnt 15
+
+ const char *files_names[files_cnt] =
+ {
+ "inc_cipher_aes256.cl",
+ "inc_cipher_serpent256.cl",
+ "inc_cipher_twofish256.cl",
+ "inc_common.cl",
+ "inc_comp_multi_bs.cl",
+ "inc_comp_multi.cl",
+ "inc_comp_single_bs.cl",
+ "inc_comp_single.cl",
+ "inc_hash_constants.h",
+ "inc_hash_functions.cl",
+ "inc_rp.cl",
+ "inc_rp.h",
+ "inc_simd.cl",
+ "inc_types.cl",
+ "inc_vendor.cl",
+ };
+
+ if (chdir (cpath_real) == -1)
+ {
+ log_error ("ERROR: %s: %s", cpath_real, strerror (errno));
+
+ return -1;
+ }
+
+ for (int i = 0; i < files_cnt; i++)
+ {
+ FILE *fd = fopen (files_names[i], "r");
+
+ if (fd == NULL)
+ {
+ log_error ("ERROR: %s: fopen(): %s", files_names[i], strerror (errno));
+
+ return -1;
+ }
+
+ char buf[1];
+
+ size_t n = fread (buf, 1, 1, fd);
+
+ if (n != 1)
+ {
+ log_error ("ERROR: %s: fread(): %s", files_names[i], strerror (errno));
+
+ return -1;
+ }
+
+ 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 };
- snprintf (build_opts_new, sizeof (build_opts_new) - 1, "%s -D VENDOR_ID=%u -D CUDA_ARCH=%d -D VECT_SIZE=%u -D DEVICE_TYPE=%u -D KERN_TYPE=%u -D _unroll -cl-std=CL1.1", build_opts, device_param->device_vendor_id, (device_param->sm_major * 100) + device_param->sm_minor, device_param->vector_width, (u32) device_param->device_type, kern_type);
+ snprintf (build_opts_new, sizeof (build_opts_new) - 1, "%s -D VENDOR_ID=%u -D CUDA_ARCH=%d -D VECT_SIZE=%u -D DEVICE_TYPE=%u -D DGST_R0=%u -D DGST_R1=%u -D DGST_R2=%u -D DGST_R3=%u -D DGST_ELEM=%u -D KERN_TYPE=%u -D _unroll -cl-std=CL1.1", build_opts, device_param->device_vendor_id, (device_param->sm_major * 100) + device_param->sm_minor, device_param->vector_width, (u32) device_param->device_type, data.dgst_pos0, data.dgst_pos1, data.dgst_pos2, data.dgst_pos3, data.dgst_size / 4, kern_type);
strncpy (build_opts, build_opts_new, sizeof (build_opts));
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);
- 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: 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);
+
+ 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;
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);
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);
- hc_clBuildProgram (data.ocl, device_param->program, 1, &device_param->device, build_opts, NULL, NULL, true);
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: clBuildProgram(): %s\n", val2cstr_cl (CL_err));
+
+ return -1;
+ }
}
}
else
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 };
if (force_jit_compilation == 1500)
{
- snprintf (build_opts_update, sizeof (build_opts_update) - 1, "%s -DDESCRYPT_SALT=%d", build_opts, data.salts_buf[0].salt_buf[0]);
+ snprintf (build_opts_update, sizeof (build_opts_update) - 1, "%s -DDESCRYPT_SALT=%u", build_opts, data.salts_buf[0].salt_buf[0]);
}
else if (force_jit_compilation == 8900)
{
- snprintf (build_opts_update, sizeof (build_opts_update) - 1, "%s -DSCRYPT_N=%d -DSCRYPT_R=%d -DSCRYPT_P=%d -DSCRYPT_TMTO=%d", build_opts, data.salts_buf[0].scrypt_N, data.salts_buf[0].scrypt_r, data.salts_buf[0].scrypt_p, 1 << data.salts_buf[0].scrypt_tmto);
+ snprintf (build_opts_update, sizeof (build_opts_update) - 1, "%s -DSCRYPT_N=%u -DSCRYPT_R=%u -DSCRYPT_P=%u -DSCRYPT_TMTO=%u -DSCRYPT_TMP_ELEM=%u", build_opts, data.salts_buf[0].scrypt_N, data.salts_buf[0].scrypt_r, data.salts_buf[0].scrypt_p, 1 << data.scrypt_tmto_final, data.scrypt_tmp_size / 16);
}
else
{
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);
+
+ CL_err = hc_clGetProgramBuildInfo (data.ocl, device_param->program, device_param->device, CL_PROGRAM_BUILD_LOG, build_log_size, build_log, NULL);
- memset (build_log, 0, build_log_size + 1);
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: clGetProgramBuildInfo(): %s\n", val2cstr_cl (CL_err));
- hc_clGetProgramBuildInfo (data.ocl, device_param->program, device_param->device, CL_PROGRAM_BUILD_LOG, build_log_size, build_log, NULL);
+ return -1;
+ }
puts (build_log);
- free (build_log);
+ myfree (build_log);
}
#endif
- if (rc != 0)
+ if (CL_err != CL_SUCCESS)
{
device_param->skipped = true;
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);
+
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: clCreateProgramWithSource(): %s\n", val2cstr_cl (CL_err));
- int rc = hc_clBuildProgram (data.ocl, device_param->program_mp, 1, &device_param->device, build_opts, NULL, NULL, false);
+ 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;
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);
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));
- hc_clBuildProgram (data.ocl, device_param->program_mp, 1, &device_param->device, build_opts, NULL, NULL, true);
+ 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));
+
+ return -1;
+ }
}
local_free (kernel_lengths);
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);
- int rc = hc_clBuildProgram (data.ocl, device_param->program_amp, 1, &device_param->device, build_opts, NULL, NULL, false);
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: clCreateProgramWithSource(): %s\n", val2cstr_cl (CL_err));
- if (rc != 0)
+ 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;
+ }
+
+ if (CL_err != CL_SUCCESS)
{
device_param->skipped = true;
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);
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);
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)
* 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_scryptV_buf = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE, size_scryptV, 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
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));
- 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;
+ }
+
+ 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));
+
+ return -1;
+ }
}
/**
* kernel args
*/
- device_param->kernel_params_buf32[21] = bitmap_mask;
- device_param->kernel_params_buf32[22] = bitmap_shift1;
- device_param->kernel_params_buf32[23] = bitmap_shift2;
- device_param->kernel_params_buf32[24] = 0; // salt_pos
- device_param->kernel_params_buf32[25] = 0; // loop_pos
- device_param->kernel_params_buf32[26] = 0; // loop_cnt
- device_param->kernel_params_buf32[27] = 0; // kernel_rules_cnt
- device_param->kernel_params_buf32[28] = 0; // digests_cnt
- device_param->kernel_params_buf32[29] = 0; // digests_offset
- device_param->kernel_params_buf32[30] = 0; // combs_mode
- device_param->kernel_params_buf32[31] = 0; // gid_max
+ device_param->kernel_params_buf32[24] = bitmap_mask;
+ device_param->kernel_params_buf32[25] = bitmap_shift1;
+ device_param->kernel_params_buf32[26] = bitmap_shift2;
+ device_param->kernel_params_buf32[27] = 0; // salt_pos
+ device_param->kernel_params_buf32[28] = 0; // loop_pos
+ device_param->kernel_params_buf32[29] = 0; // loop_cnt
+ device_param->kernel_params_buf32[30] = 0; // kernel_rules_cnt
+ device_param->kernel_params_buf32[31] = 0; // digests_cnt
+ device_param->kernel_params_buf32[32] = 0; // digests_offset
+ device_param->kernel_params_buf32[33] = 0; // combs_mode
+ device_param->kernel_params_buf32[34] = 0; // gid_max
device_param->kernel_params[ 0] = (attack_exec == ATTACK_EXEC_INSIDE_KERNEL)
? &device_param->d_pws_buf
device_param->kernel_params[17] = &device_param->d_salt_bufs;
device_param->kernel_params[18] = &device_param->d_esalt_bufs;
device_param->kernel_params[19] = &device_param->d_result;
- device_param->kernel_params[20] = &device_param->d_scryptV_buf;
- device_param->kernel_params[21] = &device_param->kernel_params_buf32[21];
- device_param->kernel_params[22] = &device_param->kernel_params_buf32[22];
- device_param->kernel_params[23] = &device_param->kernel_params_buf32[23];
+ device_param->kernel_params[20] = &device_param->d_scryptV0_buf;
+ device_param->kernel_params[21] = &device_param->d_scryptV1_buf;
+ device_param->kernel_params[22] = &device_param->d_scryptV2_buf;
+ device_param->kernel_params[23] = &device_param->d_scryptV3_buf;
device_param->kernel_params[24] = &device_param->kernel_params_buf32[24];
device_param->kernel_params[25] = &device_param->kernel_params_buf32[25];
device_param->kernel_params[26] = &device_param->kernel_params_buf32[26];
device_param->kernel_params[29] = &device_param->kernel_params_buf32[29];
device_param->kernel_params[30] = &device_param->kernel_params_buf32[30];
device_param->kernel_params[31] = &device_param->kernel_params_buf32[31];
+ device_param->kernel_params[32] = &device_param->kernel_params_buf32[32];
+ device_param->kernel_params[33] = &device_param->kernel_params_buf32[33];
+ device_param->kernel_params[34] = &device_param->kernel_params_buf32[34];
device_param->kernel_params_mp_buf64[3] = 0;
device_param->kernel_params_mp_buf32[4] = 0;
{
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)
{
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;
+ }
- 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);
+ 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));
+
+ return -1;
+ }
}
}
}
{
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);
- 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));
+
+ 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;
+ }
+
+ 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));
- 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);
+ 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);
- for (uint i = 0; i <= 20; i++)
+ if (CL_err != CL_SUCCESS)
{
- 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]);
+ log_error ("ERROR: clGetKernelWorkGroupInfo(): %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 = 21; i <= 31; i++)
+ for (uint i = 0; i <= 23; 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_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_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]);
+ return -1;
+ }
+ }
+
+ for (uint i = 24; i <= 34; 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 (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);
- 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);
+ 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_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));
+
+ return -1;
+ }
- 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);
+ 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);
+
+ 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;
+ }
}
if (attack_exec == ATTACK_EXEC_INSIDE_KERNEL)
}
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);
+
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: clGetKernelWorkGroupInfo(): %s\n", val2cstr_cl (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);
+ return -1;
+ }
}
if (attack_exec == ATTACK_EXEC_INSIDE_KERNEL)
{
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;
+ }
}
}
}
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
{
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
{
log_error ("ERROR: %s: %s", l0_filename, strerror (errno));
- return (-1);
+ return -1;
}
uint is_dir = S_ISDIR (l0_stat.st_mode);
{
log_error ("ERROR: Keyspace parameter is not allowed together with a directory");
- return (-1);
+ return -1;
}
char **dictionary_files = NULL;
{
log_error ("ERROR: %s: %s", l1_filename, strerror (errno));
- return (-1);
+ return -1;
}
if (S_ISREG (l1_stat.st_mode))
{
log_error ("ERROR: No usable dictionary file found.");
- return (-1);
+ return -1;
}
}
else if (wordlist_mode == WL_MODE_STDIN)
{
log_error ("ERROR: %s: %s", dictfile1, strerror (errno));
- return (-1);
+ return -1;
}
if (stat (dictfile1, &tmp_stat) == -1)
fclose (fp1);
- return (-1);
+ return -1;
}
if (S_ISDIR (tmp_stat.st_mode))
fclose (fp1);
- return (-1);
+ return -1;
}
if ((fp2 = fopen (dictfile2, "rb")) == NULL)
fclose (fp1);
- return (-1);
+ return -1;
}
if (stat (dictfile2, &tmp_stat) == -1)
fclose (fp1);
fclose (fp2);
- return (-1);
+ return -1;
}
if (S_ISDIR (tmp_stat.st_mode))
fclose (fp1);
fclose (fp2);
- return (-1);
+ return -1;
}
data.combs_cnt = 1;
fclose (fp1);
fclose (fp2);
- return (-1);
+ return -1;
}
data.combs_cnt = 1;
fclose (fp1);
fclose (fp2);
- return (-1);
+ return -1;
}
fclose (fp1);
{
log_error ("ERROR: %s: %s", mask, strerror (errno));
- return (-1);
+ return -1;
}
}
{
log_error ("ERROR: %s: %s", mask, strerror (errno));
- return (-1);
+ return -1;
}
char *line_buf = (char *) mymalloc (HCBUFSIZ);
{
log_error ("ERROR: %s: unsupported file-type", mask);
- return (-1);
+ return -1;
}
}
{
log_error ("ERROR: %s: %s", mask, strerror (errno));
- return (-1);
+ return -1;
}
char *line_buf = (char *) mymalloc (HCBUFSIZ);
{
log_error ("ERROR: %s: %s", filename, strerror (errno));
- return (-1);
+ return -1;
}
uint is_dir = S_ISDIR (file_stat.st_mode);
{
log_error ("ERROR: Keyspace parameter is not allowed together with a directory");
- return (-1);
+ return -1;
}
char **dictionary_files = NULL;
{
log_error ("ERROR: %s: %s", l1_filename, strerror (errno));
- return (-1);
+ return -1;
}
if (S_ISREG (l1_stat.st_mode))
{
log_error ("ERROR: No usable dictionary file found.");
- return (-1);
+ return -1;
}
if (increment)
{
log_error ("ERROR: %s: %s", mask, strerror (errno));
- return (-1);
+ return -1;
}
char *line_buf = (char *) mymalloc (HCBUFSIZ);
{
log_error ("ERROR: %s: %s", filename, strerror (errno));
- return (-1);
+ return -1;
}
uint is_dir = S_ISDIR (file_stat.st_mode);
{
log_error ("ERROR: Keyspace parameter is not allowed together with a directory");
- return (-1);
+ return -1;
}
char **dictionary_files = NULL;
{
log_error ("ERROR: %s: %s", l1_filename, strerror (errno));
- return (-1);
+ return -1;
}
if (S_ISREG (l1_stat.st_mode))
{
log_error ("ERROR: No usable dictionary file found.");
- return (-1);
+ return -1;
}
if (increment)
* status and monitor threads
*/
- if ((data.devices_status != STATUS_CRACKED) && (data.devices_status != STATUS_ABORTED) && (data.devices_status != STATUS_QUIT))
+ if ((data.devices_status != STATUS_BYPASS) && (data.devices_status != STATUS_CRACKED) && (data.devices_status != STATUS_ABORTED) && (data.devices_status != STATUS_QUIT))
{
data.devices_status = STATUS_STARTING;
}
hc_thread_t *inner_threads = (hc_thread_t *) mycalloc (10, sizeof (hc_thread_t));
+ data.shutdown_inner = 0;
+
/**
* Outfile remove
*/
if (keyspace == 0 && benchmark == 0 && stdout_flag == 0)
{
+ hc_thread_create (inner_threads[inner_threads_cnt], thread_monitor, NULL);
+
+ inner_threads_cnt++;
+
if (outfile_check_timer != 0)
{
if (data.outfile_check_directory != NULL)
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))
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;
+
+ 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);
- 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);
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: clEnqueueWriteBuffer(): %s\n", val2cstr_cl (CL_err));
+
+ return -1;
+ }
}
}
else if (attack_mode == ATTACK_MODE_BF)
{
log_error ("ERROR: Mask is too small");
- return (-1);
+ return -1;
}
}
}
{
log_error ("ERROR: --keyspace is not supported with --increment or mask files");
- return (-1);
+ return -1;
}
}
logfile_sub_msg ("START");
- if ((data.devices_status != STATUS_CRACKED) && (data.devices_status != STATUS_ABORTED) && (data.devices_status != STATUS_QUIT))
+ if ((data.devices_status != STATUS_BYPASS) && (data.devices_status != STATUS_CRACKED) && (data.devices_status != STATUS_ABORTED) && (data.devices_status != STATUS_QUIT))
{
data.devices_status = STATUS_INIT;
}
{
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);
{
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);
{
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);
{
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);
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++) 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]);
+ 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));
- 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;
+ }
+
+ 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;
+ }
}
}
{
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)
hc_thread_t *c_threads = (hc_thread_t *) mycalloc (data.devices_cnt, sizeof (hc_thread_t));
- if ((data.devices_status != STATUS_CRACKED) && (data.devices_status != STATUS_ABORTED) && (data.devices_status != STATUS_QUIT))
+ if ((data.devices_status != STATUS_BYPASS) && (data.devices_status != STATUS_CRACKED) && (data.devices_status != STATUS_ABORTED) && (data.devices_status != STATUS_QUIT))
{
data.devices_status = STATUS_AUTOTUNE;
}
* create cracker threads
*/
- if ((data.devices_status != STATUS_CRACKED) && (data.devices_status != STATUS_ABORTED) && (data.devices_status != STATUS_QUIT))
+ if ((data.devices_status != STATUS_BYPASS) && (data.devices_status != STATUS_CRACKED) && (data.devices_status != STATUS_ABORTED) && (data.devices_status != STATUS_QUIT))
{
data.devices_status = STATUS_RUNNING;
}
local_free (c_threads);
+ if ((data.devices_status != STATUS_BYPASS) && (data.devices_status != STATUS_CRACKED) && (data.devices_status != STATUS_ABORTED) && (data.devices_status != STATUS_QUIT))
+ {
+ data.devices_status = STATUS_EXHAUSTED;
+ }
+
logfile_sub_var_uint ("status-after-work", data.devices_status);
data.restore = 0;
log_info ("");
- if (stdout_flag == 0) status_display ();
+ status_display ();
log_info ("");
}
+ else
+ {
+ if (status == 1)
+ {
+ status_display ();
+ }
+ }
}
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;
data.devices_status = STATUS_RUNNING;
}
+ // and overwrite benchmark aborts as well
+
+ if (data.benchmark == 1)
+ {
+ if (data.devices_status == STATUS_ABORTED)
+ {
+ data.devices_status = STATUS_RUNNING;
+ }
+ }
+
// finalize task
if (data.devices_status == STATUS_CRACKED) break;
}
}
- // wait for non-interactive threads
+ // wait for inner threads
- if ((data.devices_status != STATUS_CRACKED) && (data.devices_status != STATUS_ABORTED) && (data.devices_status != STATUS_QUIT))
- {
- data.devices_status = STATUS_EXHAUSTED;
- }
+ data.shutdown_inner = 1;
for (uint thread_idx = 0; thread_idx < inner_threads_cnt; thread_idx++)
{
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);
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_scryptV_buf) hc_clReleaseMemObject (data.ocl, device_param->d_scryptV_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
#ifdef HAVE_HWMON
if (gpu_temp_disable == 0)
{
- if (gpu_temp_retain != 0) // VENDOR_ID_AMD is implied here
+ if (gpu_temp_retain != 0)
{
hc_thread_mutex_lock (mux_adl);
}
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
#ifdef WIN
- rc = hm_set_fanspeed_with_device_id_nvapi (device_id, 100, 8);
+ rc = hm_set_fanspeed_with_device_id_nvapi (device_id, 100, 0);
#endif
}
// reset power tuning
- if (powertune_enable == 1) // VENDOR_ID_AMD is implied here
+ if (powertune_enable == 1)
{
hc_thread_mutex_lock (mux_adl);
{
log_error ("ERROR: Failed to get ADL PowerControl Capabilities");
- return (-1);
+ return -1;
}
if (powertune_supported != 0)
{
log_info ("ERROR: Failed to restore the ADL PowerControl values");
- return (-1);
+ return -1;
}
// clocks
{
log_info ("ERROR: Failed to restore ADL performance state");
- return (-1);
+ return -1;
}
local_free (performance_state);
if (data.devices_status == STATUS_QUIT) break;
}
- // wait for interactive threads
+ // wait for outer threads
+
+ data.shutdown_outer = 1;
for (uint thread_idx = 0; thread_idx < outer_threads_cnt; thread_idx++)
{
{
log_error ("ERROR: %s: %s", induction_directory, strerror (errno));
- return (-1);
+ return -1;
}
}
{
log_error ("ERROR: %s: %s", outfile_check_directory, strerror (errno));
- return (-1);
+ return -1;
}
}