* License.....: MIT
*/
-#if defined(DARWIN) || defined(__FreeBSD__)
+#ifdef __APPLE__
+#include <stdio.h>
+#endif
+
+#ifdef __FreeBSD__
#include <stdio.h>
#endif
#define INCREMENT 0
#define INCREMENT_MIN 1
#define INCREMENT_MAX PW_MAX
+#define MANGLE 0
#define SEPARATOR ':'
#define BITMAP_MIN 16
#define BITMAP_MAX 24
" -i, --increment | | Enable mask increment mode |",
" --increment-min | Num | Start mask incrementing at X | --increment-min=4",
" --increment-max | Num | Stop mask incrementing at X | --increment-max=8",
+ " --mangle | | Mangle password before hashing |",
"",
"- [ Hash modes ] -",
"",
#ifdef WIN
- __time64_t runtime_left = data.proc_start + data.runtime - runtime_cur;
+ __time64_t runtime_left = data.proc_start + data.runtime + data.prepare_time - runtime_cur;
tmp = _gmtime64 (&runtime_left);
#else
- time_t runtime_left = data.proc_start + data.runtime - runtime_cur;
+ time_t runtime_left = data.proc_start + data.runtime + data.prepare_time - runtime_cur;
tmp = gmtime (&runtime_left);
* hashcat -only- functions
*/
-static void generate_source_kernel_filename (const uint attack_exec, const uint attack_kern, const uint kern_type, char *shared_dir, char *source_file)
+static void generate_source_kernel_filename (const uint attack_exec, const uint attack_kern, const uint kern_type, uint mangle, char *shared_dir, char *source_file)
{
+ char const * postfix = "";
+ if (mangle == 1)
+ {
+ postfix = "_m";
+ }
+
if (attack_exec == ATTACK_EXEC_INSIDE_KERNEL)
{
if (attack_kern == ATTACK_KERN_STRAIGHT)
- snprintf (source_file, 255, "%s/OpenCL/m%05d_a0.cl", shared_dir, (int) kern_type);
+ snprintf (source_file, 255, "%s/OpenCL/m%05d_a0%s.cl", shared_dir, (int) kern_type, postfix);
else if (attack_kern == ATTACK_KERN_COMBI)
- snprintf (source_file, 255, "%s/OpenCL/m%05d_a1.cl", shared_dir, (int) kern_type);
+ snprintf (source_file, 255, "%s/OpenCL/m%05d_a1%s.cl", shared_dir, (int) kern_type, postfix);
else if (attack_kern == ATTACK_KERN_BF)
- snprintf (source_file, 255, "%s/OpenCL/m%05d_a3.cl", shared_dir, (int) kern_type);
+ snprintf (source_file, 255, "%s/OpenCL/m%05d_a3%s.cl", shared_dir, (int) kern_type, postfix);
}
else
snprintf (source_file, 255, "%s/OpenCL/m%05d.cl", shared_dir, (int) kern_type);
}
-static void generate_cached_kernel_filename (const uint attack_exec, const uint attack_kern, const uint kern_type, char *profile_dir, const char *device_name_chksum, char *cached_file)
+static void generate_cached_kernel_filename (const uint attack_exec, const uint attack_kern, const uint kern_type, uint mangle, char *profile_dir, const char *device_name_chksum, char *cached_file)
{
+ char const * postfix = "";
+ if (mangle == 1)
+ {
+ postfix = "_m";
+ }
+
if (attack_exec == ATTACK_EXEC_INSIDE_KERNEL)
{
if (attack_kern == ATTACK_KERN_STRAIGHT)
- snprintf (cached_file, 255, "%s/kernels/m%05d_a0.%s.kernel", profile_dir, (int) kern_type, device_name_chksum);
+ snprintf (cached_file, 255, "%s/kernels/m%05d_a0%s.%s.kernel", profile_dir, (int) kern_type, postfix, device_name_chksum);
else if (attack_kern == ATTACK_KERN_COMBI)
- snprintf (cached_file, 255, "%s/kernels/m%05d_a1.%s.kernel", profile_dir, (int) kern_type, device_name_chksum);
+ snprintf (cached_file, 255, "%s/kernels/m%05d_a1%s.%s.kernel", profile_dir, (int) kern_type, postfix, device_name_chksum);
else if (attack_kern == ATTACK_KERN_BF)
- snprintf (cached_file, 255, "%s/kernels/m%05d_a3.%s.kernel", profile_dir, (int) kern_type, device_name_chksum);
+ snprintf (cached_file, 255, "%s/kernels/m%05d_a3%s.%s.kernel", profile_dir, (int) kern_type, postfix, device_name_chksum);
}
else
{
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 ;)
unlink (old_hashfile);
}
-static void run_kernel (const uint kern_run, hc_device_param_t *device_param, const uint num, const uint event_update, const uint iteration)
+static int run_kernel (const uint kern_run, hc_device_param_t *device_param, const uint num, const uint event_update, const uint iteration)
{
+ cl_int CL_err = CL_SUCCESS;
+
uint num_elements = num;
device_param->kernel_params_buf32[33] = data.combs_mode;
case KERN_RUN_3: kernel = device_param->kernel3; break;
}
- hc_clSetKernelArg (data.ocl, kernel, 24, sizeof (cl_uint), device_param->kernel_params[24]);
- hc_clSetKernelArg (data.ocl, kernel, 25, sizeof (cl_uint), device_param->kernel_params[25]);
- hc_clSetKernelArg (data.ocl, kernel, 26, sizeof (cl_uint), device_param->kernel_params[26]);
- hc_clSetKernelArg (data.ocl, kernel, 27, sizeof (cl_uint), device_param->kernel_params[27]);
- hc_clSetKernelArg (data.ocl, kernel, 28, sizeof (cl_uint), device_param->kernel_params[28]);
- hc_clSetKernelArg (data.ocl, kernel, 29, sizeof (cl_uint), device_param->kernel_params[29]);
- hc_clSetKernelArg (data.ocl, kernel, 30, sizeof (cl_uint), device_param->kernel_params[30]);
- hc_clSetKernelArg (data.ocl, kernel, 31, sizeof (cl_uint), device_param->kernel_params[31]);
- hc_clSetKernelArg (data.ocl, kernel, 32, sizeof (cl_uint), device_param->kernel_params[32]);
- hc_clSetKernelArg (data.ocl, kernel, 33, sizeof (cl_uint), device_param->kernel_params[33]);
- hc_clSetKernelArg (data.ocl, kernel, 34, sizeof (cl_uint), device_param->kernel_params[34]);
+ CL_err |= hc_clSetKernelArg (data.ocl, kernel, 24, sizeof (cl_uint), device_param->kernel_params[24]);
+ CL_err |= hc_clSetKernelArg (data.ocl, kernel, 25, sizeof (cl_uint), device_param->kernel_params[25]);
+ CL_err |= hc_clSetKernelArg (data.ocl, kernel, 26, sizeof (cl_uint), device_param->kernel_params[26]);
+ CL_err |= hc_clSetKernelArg (data.ocl, kernel, 27, sizeof (cl_uint), device_param->kernel_params[27]);
+ CL_err |= hc_clSetKernelArg (data.ocl, kernel, 28, sizeof (cl_uint), device_param->kernel_params[28]);
+ CL_err |= hc_clSetKernelArg (data.ocl, kernel, 29, sizeof (cl_uint), device_param->kernel_params[29]);
+ CL_err |= hc_clSetKernelArg (data.ocl, kernel, 30, sizeof (cl_uint), device_param->kernel_params[30]);
+ CL_err |= hc_clSetKernelArg (data.ocl, kernel, 31, sizeof (cl_uint), device_param->kernel_params[31]);
+ CL_err |= hc_clSetKernelArg (data.ocl, kernel, 32, sizeof (cl_uint), device_param->kernel_params[32]);
+ CL_err |= hc_clSetKernelArg (data.ocl, kernel, 33, sizeof (cl_uint), device_param->kernel_params[33]);
+ CL_err |= hc_clSetKernelArg (data.ocl, kernel, 34, sizeof (cl_uint), device_param->kernel_params[34]);
+
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: clSetKernelArg(): %s\n", val2cstr_cl (CL_err));
+
+ return -1;
+ }
cl_event event;
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);
- 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;
+ }
+
+ 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);
- hc_clFlush (data.ocl, device_param->command_queue);
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: clFinish(): %s\n", val2cstr_cl (CL_err));
+
+ return -1;
+ }
- hc_clFinish (data.ocl, device_param->command_queue);
+ return 0;
}
-static void run_kernel_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);
+
+ 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_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;
+ log_error ("ERROR: clEnqueueWriteBuffer(): %s\n", val2cstr_cl (CL_err));
- const size_t fillsz = MIN (FILLSZ, left);
-
- hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, buf, CL_TRUE, i, fillsz, tmp, 0, NULL, NULL);
+ return -1;
}
-
- myfree (tmp);
}
- */
+
+ return 0;
+}
+
+static int run_kernel_bzero (hc_device_param_t *device_param, cl_mem buf, const size_t size)
+{
+ return run_kernel_memset (device_param, buf, 0, size);
}
-static void choose_kernel (hc_device_param_t *device_param, const uint attack_exec, const uint attack_mode, const uint opts_type, const salt_t *salt_buf, const uint highest_pw_len, const uint pws_cnt, const uint fast_iteration)
+static int choose_kernel (hc_device_param_t *device_param, const uint attack_exec, const uint attack_mode, const uint opts_type, const salt_t *salt_buf, const uint highest_pw_len, const uint pws_cnt, const uint fast_iteration)
{
+ cl_int CL_err = CL_SUCCESS;
+
if (data.hash_mode == 2000)
{
process_stdout (device_param, pws_cnt);
- return;
+ return 0;
}
if (attack_exec == ATTACK_EXEC_INSIDE_KERNEL)
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;
{
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)
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
{
for (int i = 0; i < STEPS_CNT; i++)
{
- const u32 kernel_accel_try = 1 << i;
+ const u32 kernel_accel_try = 1u << i;
if (kernel_accel_try < kernel_accel_min) continue;
if (kernel_accel_try > kernel_accel_max) break;
}
#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);
if (data.attack_mode == ATTACK_MODE_STRAIGHT)
{
- hc_clEnqueueCopyBuffer (data.ocl, device_param->command_queue, device_param->d_rules, device_param->d_rules_c, innerloop_pos * sizeof (kernel_rule_t), 0, innerloop_left * sizeof (kernel_rule_t), 0, NULL, NULL);
+ cl_int CL_err = hc_clEnqueueCopyBuffer (data.ocl, device_param->command_queue, device_param->d_rules, device_param->d_rules_c, innerloop_pos * sizeof (kernel_rule_t), 0, innerloop_left * sizeof (kernel_rule_t), 0, NULL, NULL);
+
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: clEnqueueCopyBuffer(): %s\n", val2cstr_cl (CL_err));
+
+ return -1;
+ }
}
else if (data.attack_mode == ATTACK_MODE_COMBI)
{
- hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_combs_c, CL_TRUE, 0, innerloop_left * sizeof (comb_t), device_param->combs_buf, 0, NULL, NULL);
+ cl_int CL_err = hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_combs_c, CL_TRUE, 0, innerloop_left * sizeof (comb_t), device_param->combs_buf, 0, NULL, NULL);
+
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: clEnqueueWriteBuffer(): %s\n", val2cstr_cl (CL_err));
+
+ return -1;
+ }
}
else if (data.attack_mode == ATTACK_MODE_BF)
{
- hc_clEnqueueCopyBuffer (data.ocl, device_param->command_queue, device_param->d_bfs, device_param->d_bfs_c, 0, 0, innerloop_left * sizeof (bf_t), 0, NULL, NULL);
+ cl_int CL_err = hc_clEnqueueCopyBuffer (data.ocl, device_param->command_queue, device_param->d_bfs, device_param->d_bfs_c, 0, 0, innerloop_left * sizeof (bf_t), 0, NULL, NULL);
+
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: clEnqueueCopyBuffer(): %s\n", val2cstr_cl (CL_err));
+
+ return -1;
+ }
}
else if (data.attack_mode == ATTACK_MODE_HYBRID1)
{
- hc_clEnqueueCopyBuffer (data.ocl, device_param->command_queue, device_param->d_combs, device_param->d_combs_c, 0, 0, innerloop_left * sizeof (comb_t), 0, NULL, NULL);
+ cl_int CL_err = hc_clEnqueueCopyBuffer (data.ocl, device_param->command_queue, device_param->d_combs, device_param->d_combs_c, 0, 0, innerloop_left * sizeof (comb_t), 0, NULL, NULL);
+
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: clEnqueueCopyBuffer(): %s\n", val2cstr_cl (CL_err));
+
+ return -1;
+ }
}
else if (data.attack_mode == ATTACK_MODE_HYBRID2)
{
- hc_clEnqueueCopyBuffer (data.ocl, device_param->command_queue, device_param->d_combs, device_param->d_combs_c, 0, 0, innerloop_left * sizeof (comb_t), 0, NULL, NULL);
+ cl_int CL_err = hc_clEnqueueCopyBuffer (data.ocl, device_param->command_queue, device_param->d_combs, device_param->d_combs_c, 0, 0, innerloop_left * sizeof (comb_t), 0, NULL, NULL);
+
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: clEnqueueCopyBuffer(): %s\n", val2cstr_cl (CL_err));
+
+ return -1;
+ }
}
if (data.benchmark == 1)
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 ();
device_param->speed_pos = speed_pos;
myfree (line_buf);
+
+ return 0;
}
static void load_segment (wl_data_t *wl_data, FILE *fd)
{
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
}
time (&runtime_cur);
- int runtime_left = data.proc_start + data.runtime - runtime_cur;
+ int runtime_left = data.proc_start + data.runtime + data.prepare_time - runtime_cur;
if (runtime_left <= 0)
{
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
*/
data.proc_start = proc_start;
+ time_t prepare_start;
+
+ time (&prepare_start);
+
int myargc = argc;
char **myargv = argv;
uint increment = INCREMENT;
uint increment_min = INCREMENT_MIN;
uint increment_max = INCREMENT_MAX;
+ uint mangle = MANGLE;
char *cpu_affinity = NULL;
OCL_PTR *ocl = NULL;
char *opencl_devices = NULL;
#define IDX_INCREMENT 'i'
#define IDX_INCREMENT_MIN 0xff12
#define IDX_INCREMENT_MAX 0xff13
+ #define IDX_MANGLE 0xff80
#define IDX_OUTFILE 'o'
#define IDX_OUTFILE_FORMAT 0xff14
#define IDX_OUTFILE_AUTOHEX_DISABLE 0xff39
{"increment", no_argument, 0, IDX_INCREMENT},
{"increment-min", required_argument, 0, IDX_INCREMENT_MIN},
{"increment-max", required_argument, 0, IDX_INCREMENT_MAX},
+ {"mangle", no_argument, 0, IDX_MANGLE},
{"custom-charset1", required_argument, 0, IDX_CUSTOM_CHARSET_1},
{"custom-charset2", required_argument, 0, IDX_CUSTOM_CHARSET_2},
{"custom-charset3", required_argument, 0, IDX_CUSTOM_CHARSET_3},
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 ();
- #if defined(LINUX) || defined(__APPLE__) || defined(__FreeBSD__)
+ #if defined(__linux__) || defined(__APPLE__) || defined(__FreeBSD__)
char *resolved_install_folder = realpath (INSTALL_FOLDER, NULL);
char *resolved_exec_path = realpath (exec_path, NULL);
{
log_error ("ERROR: %s: %s", resolved_install_folder, strerror (errno));
- return (-1);
+ return -1;
}
if (resolved_exec_path == NULL)
{
log_error ("ERROR: %s: %s", resolved_exec_path, strerror (errno));
- return (-1);
+ return -1;
}
char *install_dir = get_install_dir (resolved_exec_path);
myfree (exec_path);
+ /**
+ * There's alot of problem related to bad support -I parameters when building the kernel.
+ * Each OpenCL runtime handles it slightly different.
+ * The most problematic is with new AMD drivers on Windows, which can not handle quote characters!
+ * The best workaround found so far is to modify the TMP variable (only inside hashcat process) before the runtime is load
+ */
+
+ char cpath[1024] = { 0 };
+
+ #if _WIN
+
+ snprintf (cpath, sizeof (cpath) - 1, "%s\\OpenCL\\", shared_dir);
+
+ char *cpath_real = mymalloc (MAX_PATH);
+
+ if (GetFullPathName (cpath, MAX_PATH, cpath_real, NULL) == 0)
+ {
+ log_error ("ERROR: %s: %s", cpath, "GetFullPathName()");
+
+ return -1;
+ }
+
+ #else
+
+ snprintf (cpath, sizeof (cpath) - 1, "%s/OpenCL/", shared_dir);
+
+ char *cpath_real = mymalloc (PATH_MAX);
+
+ if (realpath (cpath, cpath_real) == NULL)
+ {
+ log_error ("ERROR: %s: %s", cpath, strerror (errno));
+
+ return -1;
+ }
+
+ #endif
+
+ //if (getenv ("TMP") == NULL)
+ if (1)
+ {
+ char tmp[1000];
+
+ snprintf (tmp, sizeof (tmp) - 1, "TMP=%s", cpath_real);
+
+ putenv (tmp);
+ }
+
+ #if _WIN
+
+ naive_replace (cpath_real, '\\', '/');
+
+ // not escaping here, windows using quotes later
+ // naive_escape (cpath_real, PATH_MAX, ' ', '\\');
+
+ #else
+
+ naive_escape (cpath_real, PATH_MAX, ' ', '\\');
+
+ #endif
+
/**
* kernel cache, we need to make sure folder exist
*/
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;
increment_min_chgd = 1; break;
case IDX_INCREMENT_MAX: increment_max = atoi (optarg);
increment_max_chgd = 1; break;
+ case IDX_MANGLE: mangle = 1; break;
case IDX_CUSTOM_CHARSET_1: custom_charset_1 = optarg; break;
case IDX_CUSTOM_CHARSET_2: custom_charset_2 = optarg; break;
case IDX_CUSTOM_CHARSET_3: custom_charset_3 = optarg; break;
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;
}
}
}
logfile_top_uint (increment);
logfile_top_uint (increment_max);
logfile_top_uint (increment_min);
+ logfile_top_uint (mangle);
logfile_top_uint (keyspace);
logfile_top_uint (left);
logfile_top_uint (logfile_disable);
dgst_pos1 = 3;
dgst_pos2 = 2;
dgst_pos3 = 1;
+// mangle = 0;
+ break;
+
+ case 1: hash_type = HASH_TYPE_MD5;
+ salt_type = SALT_TYPE_NONE;
+ attack_exec = ATTACK_EXEC_INSIDE_KERNEL;
+ opts_type = OPTS_TYPE_PT_GENERATE_LE
+ | OPTS_TYPE_PT_ADD80
+ | OPTS_TYPE_PT_ADDBITS14;
+ kern_type = KERN_TYPE_MD5;
+ dgst_size = DGST_SIZE_4_4;
+ parse_func = md5_parse_hash;
+ sort_by_digest = sort_by_digest_4_4;
+ opti_type = OPTI_TYPE_ZERO_BYTE
+ | OPTI_TYPE_PRECOMPUTE_INIT
+ | OPTI_TYPE_PRECOMPUTE_MERKLE
+ | OPTI_TYPE_MEET_IN_MIDDLE
+ | OPTI_TYPE_EARLY_SKIP
+ | OPTI_TYPE_NOT_ITERATED
+ | OPTI_TYPE_NOT_SALTED
+ | OPTI_TYPE_RAW_HASH;
+ dgst_pos0 = 0;
+ dgst_pos1 = 3;
+ dgst_pos2 = 2;
+ dgst_pos3 = 1;
+// mangle = 1;
break;
case 10: hash_type = HASH_TYPE_MD5;
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
}
if (bitmap_max < bitmap_min) bitmap_max = bitmap_min;
- uint *bitmap_s1_a = (uint *) mymalloc ((1 << bitmap_max) * sizeof (uint));
- uint *bitmap_s1_b = (uint *) mymalloc ((1 << bitmap_max) * sizeof (uint));
- uint *bitmap_s1_c = (uint *) mymalloc ((1 << bitmap_max) * sizeof (uint));
- uint *bitmap_s1_d = (uint *) mymalloc ((1 << bitmap_max) * sizeof (uint));
- uint *bitmap_s2_a = (uint *) mymalloc ((1 << bitmap_max) * sizeof (uint));
- uint *bitmap_s2_b = (uint *) mymalloc ((1 << bitmap_max) * sizeof (uint));
- uint *bitmap_s2_c = (uint *) mymalloc ((1 << bitmap_max) * sizeof (uint));
- uint *bitmap_s2_d = (uint *) mymalloc ((1 << bitmap_max) * sizeof (uint));
+ uint *bitmap_s1_a = (uint *) mymalloc ((1u << bitmap_max) * sizeof (uint));
+ uint *bitmap_s1_b = (uint *) mymalloc ((1u << bitmap_max) * sizeof (uint));
+ uint *bitmap_s1_c = (uint *) mymalloc ((1u << bitmap_max) * sizeof (uint));
+ uint *bitmap_s1_d = (uint *) mymalloc ((1u << bitmap_max) * sizeof (uint));
+ uint *bitmap_s2_a = (uint *) mymalloc ((1u << bitmap_max) * sizeof (uint));
+ uint *bitmap_s2_b = (uint *) mymalloc ((1u << bitmap_max) * sizeof (uint));
+ uint *bitmap_s2_c = (uint *) mymalloc ((1u << bitmap_max) * sizeof (uint));
+ uint *bitmap_s2_d = (uint *) mymalloc ((1u << bitmap_max) * sizeof (uint));
uint bitmap_bits;
uint bitmap_nums;
{
if (data.quiet == 0) log_info_nn ("Generating bitmap tables with %u bits...", bitmap_bits);
- bitmap_nums = 1 << bitmap_bits;
+ bitmap_nums = 1u << bitmap_bits;
bitmap_mask = bitmap_nums - 1;
break;
}
- bitmap_nums = 1 << bitmap_bits;
+ bitmap_nums = 1u << bitmap_bits;
bitmap_mask = bitmap_nums - 1;
{
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;
}
}
}
for (uint platform_id = 0; platform_id < platforms_cnt; platform_id++)
{
- if ((opencl_platforms_filter & (1 << platform_id)) == 0) continue;
+ if ((opencl_platforms_filter & (1u << platform_id)) == 0) continue;
cl_platform_id platform = platforms[platform_id];
- hc_clGetDeviceIDs (data.ocl, platform, CL_DEVICE_TYPE_ALL, DEVICES_MAX, platform_devices, &platform_devices_cnt);
+ cl_int CL_err = hc_clGetDeviceIDs (data.ocl, platform, CL_DEVICE_TYPE_ALL, DEVICES_MAX, platform_devices, &platform_devices_cnt);
+
+ if (CL_err != CL_SUCCESS)
+ {
+ //log_error ("ERROR: clGetDeviceIDs(): %s\n", val2cstr_cl (CL_err));
+
+ //return -1;
+
+ // Silently ignore at this point, it will be reused later and create a note for the user at that point
+
+ continue;
+ }
for (uint platform_devices_id = 0; platform_devices_id < platform_devices_cnt; platform_devices_id++)
{
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 & (1u << platform_id)) == 0);
+
+ CL_err = hc_clGetDeviceIDs (data.ocl, platform, CL_DEVICE_TYPE_ALL, DEVICES_MAX, platform_devices, &platform_devices_cnt);
+
+ if (CL_err != CL_SUCCESS)
+ {
+ //log_error ("ERROR: clGetDeviceIDs(): %s\n", val2cstr_cl (CL_err));
+
+ //return -1;
+
+ platform_skipped = 2;
+ }
if ((benchmark == 1 || quiet == 0) && (algorithm_pos == 0))
{
log_info (line);
}
- else
+ else if (platform_skipped == 1)
{
log_info ("OpenCL Platform #%u: %s, skipped", platform_id + 1, platform_vendor);
log_info ("");
}
+ else if (platform_skipped == 2)
+ {
+ log_info ("OpenCL Platform #%u: %s, skipped! No OpenCL compatible devices found", platform_id + 1, platform_vendor);
+ log_info ("");
+ }
}
}
if (platform_skipped == 1) continue;
+ if (platform_skipped == 2) continue;
for (uint platform_devices_id = 0; platform_devices_id < platform_devices_cnt; platform_devices_id++)
{
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)
{
// skipped
- device_param->skipped |= ((devices_filter & (1 << device_id)) == 0);
+ device_param->skipped |= ((devices_filter & (1u << device_id)) == 0);
device_param->skipped |= ((device_types_filter & (device_type)) == 0);
// driver_version
- hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DRIVER_VERSION, 0, NULL, ¶m_value_size);
+ CL_err = hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DRIVER_VERSION, 0, NULL, ¶m_value_size);
+
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: clGetDeviceInfo(): %s\n", val2cstr_cl (CL_err));
+
+ return -1;
+ }
char *driver_version = (char *) mymalloc (param_value_size);
- hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DRIVER_VERSION, param_value_size, driver_version, NULL);
+ CL_err = hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DRIVER_VERSION, param_value_size, driver_version, NULL);
+
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: clGetDeviceInfo(): %s\n", val2cstr_cl (CL_err));
+
+ return -1;
+ }
device_param->driver_version = driver_version;
{
need_nvml = 1;
- #ifdef LINUX
+ #ifdef __linux__
need_xnvctrl = 1;
#endif
#define CL_DEVICE_KERNEL_EXEC_TIMEOUT_NV 0x4005
- hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_KERNEL_EXEC_TIMEOUT_NV, sizeof (kernel_exec_timeout), &kernel_exec_timeout, NULL);
+ CL_err = hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_KERNEL_EXEC_TIMEOUT_NV, sizeof (kernel_exec_timeout), &kernel_exec_timeout, NULL);
+
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: clGetDeviceInfo(): %s\n", val2cstr_cl (CL_err));
+
+ return -1;
+ }
device_param->kernel_exec_timeout = kernel_exec_timeout;
#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
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_info ("Rules: %u", kernel_rules_cnt);
}
+ log_info ("Mangle: %d", mangle);
+
if (opti_type)
{
log_info ("Applicable Optimizers:");
{
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
*/
* 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
{
size_scrypt = (128 * scrypt_r) * scrypt_N;
- size_scrypt /= 1 << tmto;
+ size_scrypt /= 1u << tmto;
size_scrypt *= device_param->device_processors * device_param->kernel_threads * device_param->kernel_accel_max;
* default building options
*/
- char cpath[1024] = { 0 };
-
- char build_opts[1024] = { 0 };
-
- #if _WIN
-
- snprintf (cpath, sizeof (cpath) - 1, "%s\\OpenCL\\", shared_dir);
-
- char *cpath_real = mymalloc (MAX_PATH);
-
- if (GetFullPathName (cpath, MAX_PATH, cpath_real, NULL) == 0)
+ if (chdir (cpath_real) == -1)
{
- log_error ("ERROR: %s: %s", cpath, "GetFullPathName()");
+ log_error ("ERROR: %s: %s", cpath_real, strerror (errno));
return -1;
}
- naive_replace (cpath_real, '\\', '/');
-
- // not escaping here, windows has quotes
+ char build_opts[1024] = { 0 };
+ #if _WIN
snprintf (build_opts, sizeof (build_opts) - 1, "-I \"%s\"", cpath_real);
-
#else
-
- snprintf (cpath, sizeof (cpath) - 1, "%s/OpenCL/", shared_dir);
-
- char *cpath_real = mymalloc (PATH_MAX);
-
- if (realpath (cpath, cpath_real) == NULL)
- {
- log_error ("ERROR: %s: %s", cpath, strerror (errno));
-
- return -1;
- }
-
- naive_escape (cpath_real, PATH_MAX, ' ', '\\');
-
snprintf (build_opts, sizeof (build_opts) - 1, "-I %s", cpath_real);
-
#endif
// include check
for (int i = 0; i < files_cnt; i++)
{
- char path[1024] = { 0 };
-
- snprintf (path, sizeof (path) - 1, "%s/%s", cpath_real, files_names[i]);
-
- FILE *fd = fopen (path, "r");
+ FILE *fd = fopen (files_names[i], "r");
if (fd == NULL)
{
- log_error ("ERROR: %s: fopen(): %s", path, strerror (errno));
+ log_error ("ERROR: %s: fopen(): %s", files_names[i], strerror (errno));
return -1;
}
if (n != 1)
{
- log_error ("ERROR: %s: fread(): %s", path, strerror (errno));
+ log_error ("ERROR: %s: fread(): %s", files_names[i], strerror (errno));
return -1;
}
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 };
char source_file[256] = { 0 };
- generate_source_kernel_filename (attack_exec, attack_kern, kern_type, shared_dir, source_file);
+ generate_source_kernel_filename (attack_exec, attack_kern, kern_type, mangle, shared_dir, source_file);
struct stat sst;
char cached_file[256] = { 0 };
- generate_cached_kernel_filename (attack_exec, attack_kern, kern_type, profile_dir, device_name_chksum, cached_file);
+ generate_cached_kernel_filename (attack_exec, attack_kern, kern_type, mangle, profile_dir, device_name_chksum, cached_file);
int cached = 1;
load_kernel (source_file, 1, kernel_lengths, kernel_sources);
- device_param->program = hc_clCreateProgramWithSource (data.ocl, device_param->context, 1, (const char **) kernel_sources, NULL);
+ CL_err = hc_clCreateProgramWithSource (data.ocl, device_param->context, 1, (const char **) kernel_sources, NULL, &device_param->program);
+
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: clCreateProgramWithSource(): %s\n", val2cstr_cl (CL_err));
+
+ return -1;
+ }
+
+ CL_err = hc_clBuildProgram (data.ocl, device_param->program, 1, &device_param->device, build_opts, NULL, NULL);
+
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: clBuildProgram(): %s\n", val2cstr_cl (CL_err));
- int rc = hc_clBuildProgram (data.ocl, device_param->program, 1, &device_param->device, build_opts, NULL, NULL, false);
+ //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;
+ }
- hc_clBuildProgram (data.ocl, device_param->program, 1, &device_param->device, build_opts, NULL, NULL, true);
+ 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;
+ }
}
}
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 };
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));
+
+ return -1;
+ }
- int rc = hc_clBuildProgram (data.ocl, device_param->program_mp, 1, &device_param->device, build_opts, NULL, NULL, false);
+ CL_err = hc_clBuildProgram (data.ocl, device_param->program_mp, 1, &device_param->device, build_opts, NULL, NULL);
- if (rc != 0)
+ 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);
- hc_clBuildProgram (data.ocl, device_param->program_mp, 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_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);
+
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: clCreateProgramWithSource(): %s\n", val2cstr_cl (CL_err));
+
+ return -1;
+ }
- int rc = hc_clBuildProgram (data.ocl, device_param->program_amp, 1, &device_param->device, build_opts, NULL, NULL, false);
+ CL_err = hc_clBuildProgram (data.ocl, device_param->program_amp, 1, &device_param->device, build_opts, NULL, NULL);
- if (rc != 0)
+ 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_scryptV0_buf = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE, size_scrypt4, NULL);
- device_param->d_scryptV1_buf = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE, size_scrypt4, NULL);
- device_param->d_scryptV2_buf = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE, size_scrypt4, NULL);
- device_param->d_scryptV3_buf = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE, size_scrypt4, NULL);
-
- hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s1_a, CL_TRUE, 0, bitmap_size, bitmap_s1_a, 0, NULL, NULL);
- hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s1_b, CL_TRUE, 0, bitmap_size, bitmap_s1_b, 0, NULL, NULL);
- hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s1_c, CL_TRUE, 0, bitmap_size, bitmap_s1_c, 0, NULL, NULL);
- hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s1_d, CL_TRUE, 0, bitmap_size, bitmap_s1_d, 0, NULL, NULL);
- hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s2_a, CL_TRUE, 0, bitmap_size, bitmap_s2_a, 0, NULL, NULL);
- hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s2_b, CL_TRUE, 0, bitmap_size, bitmap_s2_b, 0, NULL, NULL);
- hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s2_c, CL_TRUE, 0, bitmap_size, bitmap_s2_c, 0, NULL, NULL);
- hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s2_d, CL_TRUE, 0, bitmap_size, bitmap_s2_d, 0, NULL, NULL);
- hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_digests_buf, CL_TRUE, 0, size_digests, data.digests_buf, 0, NULL, NULL);
- hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_digests_shown, CL_TRUE, 0, size_shown, data.digests_shown, 0, NULL, NULL);
- hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_salt_bufs, CL_TRUE, 0, size_salts, data.salts_buf, 0, NULL, NULL);
+ CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_pws, NULL, &device_param->d_pws_buf);
+ CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_pws, NULL, &device_param->d_pws_amp_buf);
+ CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE, size_tmps, NULL, &device_param->d_tmps);
+ CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE, size_hooks, NULL, &device_param->d_hooks);
+ CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL, &device_param->d_bitmap_s1_a);
+ CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL, &device_param->d_bitmap_s1_b);
+ CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL, &device_param->d_bitmap_s1_c);
+ CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL, &device_param->d_bitmap_s1_d);
+ CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL, &device_param->d_bitmap_s2_a);
+ CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL, &device_param->d_bitmap_s2_b);
+ CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL, &device_param->d_bitmap_s2_c);
+ CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL, &device_param->d_bitmap_s2_d);
+ CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE, size_plains, NULL, &device_param->d_plain_bufs);
+ CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_digests, NULL, &device_param->d_digests_buf);
+ CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE, size_shown, NULL, &device_param->d_digests_shown);
+ CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_salts, NULL, &device_param->d_salt_bufs);
+ CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE, size_results, NULL, &device_param->d_result);
+ CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE, size_scrypt4, NULL, &device_param->d_scryptV0_buf);
+ CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE, size_scrypt4, NULL, &device_param->d_scryptV1_buf);
+ CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE, size_scrypt4, NULL, &device_param->d_scryptV2_buf);
+ CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE, size_scrypt4, NULL, &device_param->d_scryptV3_buf);
+
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: clCreateBuffer(): %s\n", val2cstr_cl (CL_err));
+
+ return -1;
+ }
+
+ CL_err |= hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s1_a, CL_TRUE, 0, bitmap_size, bitmap_s1_a, 0, NULL, NULL);
+ CL_err |= hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s1_b, CL_TRUE, 0, bitmap_size, bitmap_s1_b, 0, NULL, NULL);
+ CL_err |= hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s1_c, CL_TRUE, 0, bitmap_size, bitmap_s1_c, 0, NULL, NULL);
+ CL_err |= hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s1_d, CL_TRUE, 0, bitmap_size, bitmap_s1_d, 0, NULL, NULL);
+ CL_err |= hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s2_a, CL_TRUE, 0, bitmap_size, bitmap_s2_a, 0, NULL, NULL);
+ CL_err |= hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s2_b, CL_TRUE, 0, bitmap_size, bitmap_s2_b, 0, NULL, NULL);
+ CL_err |= hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s2_c, CL_TRUE, 0, bitmap_size, bitmap_s2_c, 0, NULL, NULL);
+ CL_err |= hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s2_d, CL_TRUE, 0, bitmap_size, bitmap_s2_d, 0, NULL, NULL);
+ CL_err |= hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_digests_buf, CL_TRUE, 0, size_digests, data.digests_buf, 0, NULL, NULL);
+ CL_err |= hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_digests_shown, CL_TRUE, 0, size_shown, data.digests_shown, 0, NULL, NULL);
+ CL_err |= hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_salt_bufs, CL_TRUE, 0, size_salts, data.salts_buf, 0, NULL, NULL);
+
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: clEnqueueWriteBuffer(): %s\n", val2cstr_cl (CL_err));
+
+ return -1;
+ }
/**
* special buffers
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);
+
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: clCreateBuffer(): %s\n", val2cstr_cl (CL_err));
+
+ return -1;
+ }
- hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_rules, CL_TRUE, 0, size_rules, kernel_rules_buf, 0, NULL, NULL);
+ CL_err = hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_rules, CL_TRUE, 0, size_rules, kernel_rules_buf, 0, NULL, NULL);
+
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: clEnqueueWriteBuffer(): %s\n", val2cstr_cl (CL_err));
+
+ return -1;
+ }
}
else if (attack_kern == ATTACK_KERN_COMBI)
{
- device_param->d_combs = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_combs, NULL);
- device_param->d_combs_c = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_combs, NULL);
- device_param->d_root_css_buf = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_root_css, NULL);
- device_param->d_markov_css_buf = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_markov_css, NULL);
+ CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_combs, NULL, &device_param->d_combs);
+ CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_combs, NULL, &device_param->d_combs_c);
+ CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_root_css, NULL, &device_param->d_root_css_buf);
+ CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_markov_css, NULL, &device_param->d_markov_css_buf);
+
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: clCreateBuffer(): %s\n", val2cstr_cl (CL_err));
+
+ return -1;
+ }
}
else if (attack_kern == ATTACK_KERN_BF)
{
- device_param->d_bfs = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_bfs, NULL);
- device_param->d_bfs_c = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_bfs, NULL);
- device_param->d_tm_c = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_tm, NULL);
- device_param->d_root_css_buf = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_root_css, NULL);
- device_param->d_markov_css_buf = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_markov_css, NULL);
+ CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_bfs, NULL, &device_param->d_bfs);
+ CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_bfs, NULL, &device_param->d_bfs_c);
+ CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_tm, NULL, &device_param->d_tm_c);
+ CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_root_css, NULL, &device_param->d_root_css_buf);
+ CL_err |= hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_markov_css, NULL, &device_param->d_markov_css_buf);
+
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: clCreateBuffer(): %s\n", val2cstr_cl (CL_err));
+
+ return -1;
+ }
}
if (size_esalts)
{
- device_param->d_esalt_bufs = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_esalts, NULL);
+ CL_err = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_esalts, NULL, &device_param->d_esalt_bufs);
+
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: clCreateBuffer(): %s\n", val2cstr_cl (CL_err));
+
+ return -1;
+ }
+
+ CL_err = hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_esalt_bufs, CL_TRUE, 0, size_esalts, data.esalts_buf, 0, NULL, NULL);
+
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: clEnqueueWriteBuffer(): %s\n", val2cstr_cl (CL_err));
- hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_esalt_bufs, CL_TRUE, 0, size_esalts, data.esalts_buf, 0, NULL, NULL);
+ return -1;
+ }
}
/**
{
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);
- 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: clCreateKernel(): %s\n", val2cstr_cl (CL_err));
+
+ return -1;
+ }
+
+ CL_err = hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel_tm, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp);
+
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: clGetKernelWorkGroupInfo(): %s\n", val2cstr_cl (CL_err));
+
+ 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;
+ }
- 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);
+ CL_err = hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel12, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp);
+
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: clGetKernelWorkGroupInfo(): %s\n", val2cstr_cl (CL_err));
+
+ 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);
- 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: 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));
+
+ return -1;
+ }
}
}
- hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel1, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp);
- hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel2, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp);
- hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel3, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp);
+ CL_err |= hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel1, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp);
+ CL_err |= hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel2, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp);
+ CL_err |= hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel3, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp);
+
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: clGetKernelWorkGroupInfo(): %s\n", val2cstr_cl (CL_err));
+
+ return -1;
+ }
for (uint i = 0; i <= 23; i++)
{
- hc_clSetKernelArg (data.ocl, device_param->kernel1, i, sizeof (cl_mem), device_param->kernel_params[i]);
- hc_clSetKernelArg (data.ocl, device_param->kernel2, i, sizeof (cl_mem), device_param->kernel_params[i]);
- hc_clSetKernelArg (data.ocl, device_param->kernel3, i, sizeof (cl_mem), device_param->kernel_params[i]);
+ CL_err |= hc_clSetKernelArg (data.ocl, device_param->kernel1, i, sizeof (cl_mem), device_param->kernel_params[i]);
+ CL_err |= hc_clSetKernelArg (data.ocl, device_param->kernel2, i, sizeof (cl_mem), device_param->kernel_params[i]);
+ CL_err |= hc_clSetKernelArg (data.ocl, device_param->kernel3, i, sizeof (cl_mem), device_param->kernel_params[i]);
+
+ if (opts_type & OPTS_TYPE_HOOK12) CL_err |= hc_clSetKernelArg (data.ocl, device_param->kernel12, i, sizeof (cl_mem), device_param->kernel_params[i]);
+ if (opts_type & OPTS_TYPE_HOOK23) CL_err |= hc_clSetKernelArg (data.ocl, device_param->kernel23, i, sizeof (cl_mem), device_param->kernel_params[i]);
+
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: clSetKernelArg(): %s\n", val2cstr_cl (CL_err));
- if (opts_type & OPTS_TYPE_HOOK12) hc_clSetKernelArg (data.ocl, device_param->kernel12, i, sizeof (cl_mem), device_param->kernel_params[i]);
- if (opts_type & OPTS_TYPE_HOOK23) hc_clSetKernelArg (data.ocl, device_param->kernel23, i, sizeof (cl_mem), device_param->kernel_params[i]);
+ return -1;
+ }
}
for (uint i = 24; i <= 34; i++)
{
- hc_clSetKernelArg (data.ocl, device_param->kernel1, i, sizeof (cl_uint), device_param->kernel_params[i]);
- hc_clSetKernelArg (data.ocl, device_param->kernel2, i, sizeof (cl_uint), device_param->kernel_params[i]);
- hc_clSetKernelArg (data.ocl, device_param->kernel3, i, sizeof (cl_uint), device_param->kernel_params[i]);
+ CL_err |= hc_clSetKernelArg (data.ocl, device_param->kernel1, i, sizeof (cl_uint), device_param->kernel_params[i]);
+ CL_err |= hc_clSetKernelArg (data.ocl, device_param->kernel2, i, sizeof (cl_uint), device_param->kernel_params[i]);
+ CL_err |= hc_clSetKernelArg (data.ocl, device_param->kernel3, i, sizeof (cl_uint), device_param->kernel_params[i]);
+
+ if (opts_type & OPTS_TYPE_HOOK12) CL_err |= hc_clSetKernelArg (data.ocl, device_param->kernel12, i, sizeof (cl_uint), device_param->kernel_params[i]);
+ if (opts_type & OPTS_TYPE_HOOK23) CL_err |= hc_clSetKernelArg (data.ocl, device_param->kernel23, i, sizeof (cl_uint), device_param->kernel_params[i]);
- if (opts_type & OPTS_TYPE_HOOK12) hc_clSetKernelArg (data.ocl, device_param->kernel12, i, sizeof (cl_uint), device_param->kernel_params[i]);
- if (opts_type & OPTS_TYPE_HOOK23) hc_clSetKernelArg (data.ocl, device_param->kernel23, i, sizeof (cl_uint), device_param->kernel_params[i]);
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: clSetKernelArg(): %s\n", val2cstr_cl (CL_err));
+
+ return -1;
+ }
}
// GPU memset
- device_param->kernel_memset = hc_clCreateKernel (data.ocl, device_param->program, "gpu_memset");
+ CL_err = hc_clCreateKernel (data.ocl, device_param->program, "gpu_memset", &device_param->kernel_memset);
+
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: clCreateKernel(): %s\n", val2cstr_cl (CL_err));
+
+ return -1;
+ }
+
+ CL_err = hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel_memset, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp);
+
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: clGetKernelWorkGroupInfo(): %s\n", val2cstr_cl (CL_err));
+
+ return -1;
+ }
+
+ CL_err |= hc_clSetKernelArg (data.ocl, device_param->kernel_memset, 0, sizeof (cl_mem), device_param->kernel_params_memset[0]);
+ CL_err |= hc_clSetKernelArg (data.ocl, device_param->kernel_memset, 1, sizeof (cl_uint), device_param->kernel_params_memset[1]);
+ CL_err |= hc_clSetKernelArg (data.ocl, device_param->kernel_memset, 2, sizeof (cl_uint), device_param->kernel_params_memset[2]);
- hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel_memset, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp);
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: clSetKernelArg(): %s\n", val2cstr_cl (CL_err));
- hc_clSetKernelArg (data.ocl, device_param->kernel_memset, 0, sizeof (cl_mem), device_param->kernel_params_memset[0]);
- hc_clSetKernelArg (data.ocl, device_param->kernel_memset, 1, sizeof (cl_uint), device_param->kernel_params_memset[1]);
- hc_clSetKernelArg (data.ocl, device_param->kernel_memset, 2, sizeof (cl_uint), device_param->kernel_params_memset[2]);
+ return -1;
+ }
// MP start
if (attack_mode == ATTACK_MODE_BF)
{
- device_param->kernel_mp_l = hc_clCreateKernel (data.ocl, device_param->program_mp, "l_markov");
- device_param->kernel_mp_r = hc_clCreateKernel (data.ocl, device_param->program_mp, "r_markov");
+ CL_err |= hc_clCreateKernel (data.ocl, device_param->program_mp, "l_markov", &device_param->kernel_mp_l);
+ CL_err |= hc_clCreateKernel (data.ocl, device_param->program_mp, "r_markov", &device_param->kernel_mp_r);
+
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: clCreateKernel(): %s\n", val2cstr_cl (CL_err));
- 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);
+ 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;
+ }
+
+ 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));
- 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;
+ }
}
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));
+
+ 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);
- 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);
- 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: 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));
+
+ 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
char *hash_type = strhashtype (data.hash_mode); // not a bug
log_info ("Hashtype: %s", hash_type);
+ log_info ("Mangle: %d", mangle);
log_info ("");
}
}
{
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)
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;
+ }
- 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);
+ CL_err |= hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_root_css_buf, CL_TRUE, 0, device_param->size_root_css, root_css_buf, 0, NULL, NULL);
+ CL_err |= hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_markov_css_buf, CL_TRUE, 0, device_param->size_markov_css, markov_css_buf, 0, NULL, NULL);
+
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: clEnqueueWriteBuffer(): %s\n", val2cstr_cl (CL_err));
+
+ return -1;
+ }
}
}
else if (attack_mode == ATTACK_MODE_BF)
{
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;
}
}
{
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++) CL_err |= hc_clSetKernelArg (data.ocl, device_param->kernel_mp_l, i, sizeof (cl_mem), (void *) device_param->kernel_params_mp_l[i]);
+ for (uint i = 3; i < 4; i++) CL_err |= hc_clSetKernelArg (data.ocl, device_param->kernel_mp_l, i, sizeof (cl_ulong), (void *) device_param->kernel_params_mp_l[i]);
+ for (uint i = 4; i < 9; i++) CL_err |= hc_clSetKernelArg (data.ocl, device_param->kernel_mp_l, i, sizeof (cl_uint), (void *) device_param->kernel_params_mp_l[i]);
+
+ for (uint i = 0; i < 3; i++) CL_err |= hc_clSetKernelArg (data.ocl, device_param->kernel_mp_r, i, sizeof (cl_mem), (void *) device_param->kernel_params_mp_r[i]);
+ for (uint i = 3; i < 4; i++) CL_err |= hc_clSetKernelArg (data.ocl, device_param->kernel_mp_r, i, sizeof (cl_ulong), (void *) device_param->kernel_params_mp_r[i]);
+ for (uint i = 4; i < 8; i++) CL_err |= hc_clSetKernelArg (data.ocl, device_param->kernel_mp_r, i, sizeof (cl_uint), (void *) device_param->kernel_params_mp_r[i]);
+
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: clSetKernelArg(): %s\n", val2cstr_cl (CL_err));
+
+ return -1;
+ }
- 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]);
+ 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;
+ }
}
}
{
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)
data.runtime_start = runtime_start;
+ data.prepare_time += runtime_start - prepare_start;
+
for (uint device_id = 0; device_id < data.devices_cnt; device_id++)
{
hc_device_param_t *device_param = &devices_param[device_id];
logfile_sub_uint (runtime_start);
logfile_sub_uint (runtime_stop);
+ time (&prepare_start);
+
logfile_sub_msg ("STOP");
global_free (subid);
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_scryptV0_buf) hc_clReleaseMemObject (data.ocl, device_param->d_scryptV0_buf);
- if (device_param->d_scryptV1_buf) hc_clReleaseMemObject (data.ocl, device_param->d_scryptV1_buf);
- if (device_param->d_scryptV2_buf) hc_clReleaseMemObject (data.ocl, device_param->d_scryptV2_buf);
- if (device_param->d_scryptV3_buf) hc_clReleaseMemObject (data.ocl, device_param->d_scryptV3_buf);
- if (device_param->d_root_css_buf) hc_clReleaseMemObject (data.ocl, device_param->d_root_css_buf);
- if (device_param->d_markov_css_buf) hc_clReleaseMemObject (data.ocl, device_param->d_markov_css_buf);
- if (device_param->d_tm_c) hc_clReleaseMemObject (data.ocl, device_param->d_tm_c);
-
- if (device_param->kernel1) hc_clReleaseKernel (data.ocl, device_param->kernel1);
- if (device_param->kernel12) hc_clReleaseKernel (data.ocl, device_param->kernel12);
- if (device_param->kernel2) hc_clReleaseKernel (data.ocl, device_param->kernel2);
- if (device_param->kernel23) hc_clReleaseKernel (data.ocl, device_param->kernel23);
- if (device_param->kernel3) hc_clReleaseKernel (data.ocl, device_param->kernel3);
- if (device_param->kernel_mp) hc_clReleaseKernel (data.ocl, device_param->kernel_mp);
- if (device_param->kernel_mp_l) hc_clReleaseKernel (data.ocl, device_param->kernel_mp_l);
- if (device_param->kernel_mp_r) hc_clReleaseKernel (data.ocl, device_param->kernel_mp_r);
- if (device_param->kernel_tm) hc_clReleaseKernel (data.ocl, device_param->kernel_tm);
- if (device_param->kernel_amp) hc_clReleaseKernel (data.ocl, device_param->kernel_amp);
- if (device_param->kernel_memset) hc_clReleaseKernel (data.ocl, device_param->kernel_memset);
-
- if (device_param->program) hc_clReleaseProgram (data.ocl, device_param->program);
- if (device_param->program_mp) hc_clReleaseProgram (data.ocl, device_param->program_mp);
- if (device_param->program_amp) hc_clReleaseProgram (data.ocl, device_param->program_amp);
-
- if (device_param->command_queue) hc_clReleaseCommandQueue (data.ocl, device_param->command_queue);
- if (device_param->context) hc_clReleaseContext (data.ocl, device_param->context);
+ if (device_param->pws_buf) myfree (device_param->pws_buf);
+
+ if (device_param->d_pws_buf) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_pws_buf);
+ if (device_param->d_pws_amp_buf) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_pws_amp_buf);
+ if (device_param->d_rules) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_rules);
+ if (device_param->d_rules_c) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_rules_c);
+ if (device_param->d_combs) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_combs);
+ if (device_param->d_combs_c) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_combs_c);
+ if (device_param->d_bfs) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_bfs);
+ if (device_param->d_bfs_c) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_bfs_c);
+ if (device_param->d_bitmap_s1_a) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_bitmap_s1_a);
+ if (device_param->d_bitmap_s1_b) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_bitmap_s1_b);
+ if (device_param->d_bitmap_s1_c) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_bitmap_s1_c);
+ if (device_param->d_bitmap_s1_d) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_bitmap_s1_d);
+ if (device_param->d_bitmap_s2_a) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_bitmap_s2_a);
+ if (device_param->d_bitmap_s2_b) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_bitmap_s2_b);
+ if (device_param->d_bitmap_s2_c) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_bitmap_s2_c);
+ if (device_param->d_bitmap_s2_d) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_bitmap_s2_d);
+ if (device_param->d_plain_bufs) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_plain_bufs);
+ if (device_param->d_digests_buf) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_digests_buf);
+ if (device_param->d_digests_shown) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_digests_shown);
+ if (device_param->d_salt_bufs) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_salt_bufs);
+ if (device_param->d_esalt_bufs) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_esalt_bufs);
+ if (device_param->d_tmps) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_tmps);
+ if (device_param->d_hooks) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_hooks);
+ if (device_param->d_result) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_result);
+ if (device_param->d_scryptV0_buf) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_scryptV0_buf);
+ if (device_param->d_scryptV1_buf) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_scryptV1_buf);
+ if (device_param->d_scryptV2_buf) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_scryptV2_buf);
+ if (device_param->d_scryptV3_buf) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_scryptV3_buf);
+ if (device_param->d_root_css_buf) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_root_css_buf);
+ if (device_param->d_markov_css_buf) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_markov_css_buf);
+ if (device_param->d_tm_c) CL_err |= hc_clReleaseMemObject (data.ocl, device_param->d_tm_c);
+
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: clReleaseMemObject(): %s\n", val2cstr_cl (CL_err));
+
+ return -1;
+ }
+
+ if (device_param->kernel1) CL_err |= hc_clReleaseKernel (data.ocl, device_param->kernel1);
+ if (device_param->kernel12) CL_err |= hc_clReleaseKernel (data.ocl, device_param->kernel12);
+ if (device_param->kernel2) CL_err |= hc_clReleaseKernel (data.ocl, device_param->kernel2);
+ if (device_param->kernel23) CL_err |= hc_clReleaseKernel (data.ocl, device_param->kernel23);
+ if (device_param->kernel3) CL_err |= hc_clReleaseKernel (data.ocl, device_param->kernel3);
+ if (device_param->kernel_mp) CL_err |= hc_clReleaseKernel (data.ocl, device_param->kernel_mp);
+ if (device_param->kernel_mp_l) CL_err |= hc_clReleaseKernel (data.ocl, device_param->kernel_mp_l);
+ if (device_param->kernel_mp_r) CL_err |= hc_clReleaseKernel (data.ocl, device_param->kernel_mp_r);
+ if (device_param->kernel_tm) CL_err |= hc_clReleaseKernel (data.ocl, device_param->kernel_tm);
+ if (device_param->kernel_amp) CL_err |= hc_clReleaseKernel (data.ocl, device_param->kernel_amp);
+ if (device_param->kernel_memset) CL_err |= hc_clReleaseKernel (data.ocl, device_param->kernel_memset);
+
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: clReleaseKernel(): %s\n", val2cstr_cl (CL_err));
+
+ return -1;
+ }
+
+ if (device_param->program) CL_err |= hc_clReleaseProgram (data.ocl, device_param->program);
+ if (device_param->program_mp) CL_err |= hc_clReleaseProgram (data.ocl, device_param->program_mp);
+ if (device_param->program_amp) CL_err |= hc_clReleaseProgram (data.ocl, device_param->program_amp);
+
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: clReleaseProgram(): %s\n", val2cstr_cl (CL_err));
+
+ return -1;
+ }
+
+ if (device_param->command_queue) CL_err |= hc_clReleaseCommandQueue (data.ocl, device_param->command_queue);
+
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: clReleaseCommandQueue(): %s\n", val2cstr_cl (CL_err));
+
+ return -1;
+ }
+
+ if (device_param->context) CL_err |= hc_clReleaseContext (data.ocl, device_param->context);
+
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: hc_clReleaseContext(): %s\n", val2cstr_cl (CL_err));
+
+ return -1;
+ }
}
// reset default fan speed
}
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
{
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);
{
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;
}
}
projects / hashcat.git / blobdiff