/**
- * Author......: Jens Steube <jens.steube@gmail.com>
+ * Authors.....: Jens Steube <jens.steube@gmail.com>
+ * Gabriele Gristina <matrix@hashcat.net>
+ *
* License.....: MIT
*/
#include <rp_kernel_on_cpu.h>
#include <getopt.h>
-const char *PROGNAME = "oclHashcat";
-const uint VERSION_BIN = 210;
-const uint RESTORE_MIN = 210;
+const char *PROGNAME = "oclHashcat";
+const uint VERSION_BIN = 210;
+const uint RESTORE_MIN = 210;
-#define INCR_RULES 10000
-#define INCR_SALTS 100000
-#define INCR_MASKS 1000
-#define INCR_POT 1000
+#define INCR_RULES 10000
+#define INCR_SALTS 100000
+#define INCR_MASKS 1000
+#define INCR_POT 1000
#define USAGE 0
#define VERSION 0
5000,
10100,
6000,
- 6100,
- 6900,
- 11700,
- 11800,
+ 6100, // broken in osx
+ 6900, // broken in osx
+ 11700, // broken in osx
+ 11800, // broken in osx
400,
8900,
11900,
12100,
23,
2500,
- 5300,
- 5400,
- 5500,
- 5600,
+ 5300, // broken in osx
+ 5400, // broken in osx
+ 5500, // broken in osx
+ 5600, // broken in osx
7300,
- 7500,
+ 7500, // broken in osx
8300,
- 11100,
+ 11100, // broken in osx
11200,
- 11400,
+ 11400, // broken in osx
121,
- 2611,
- 2711,
- 2811,
- 8400,
+ 2611, // broken in osx
+ 2711, // broken in osx
+ 2811, // broken in osx
+ 8400, // broken in osx
11,
- 2612,
+ 2612, // broken in osx
7900,
21,
11000,
124,
10000,
- 3711,
- 7600,
+ 3711, // broken in osx
+ 7600, // broken in osx
12,
131,
132,
1731,
200,
300,
- 3100,
+ 3100, // broken in osx
112,
12300,
- 8000,
+ 8000, // broken in osx
141,
1441,
1600,
- 12600,
+ 12600, // broken in osx
1421,
101,
111,
1711,
- 3000,
+ 3000, // broken in osx
1000,
1100,
2100,
12800,
- 1500,
- 12400,
+ 1500, // broken in osx
+ 12400, // broken in osx
500,
3200,
7400,
501,
5800,
8100,
- 8500,
+ 8500, // broken in osx
7200,
9900,
7700,
7800,
10300,
- 8600,
- 8700,
- 9100,
+ 8600, // broken in osx
+ 8700, // broken in osx
+ 9100, // broken in osx
133,
- 11600,
- 12500,
+ 11600, // broken in osx
+ 12500, // broken in osx
13000,
6211,
6221,
- 6231,
+ 6231, // broken in osx
6241,
- 8800,
+ 8800, // broken in osx
12900,
12200,
- 9700,
- 9710,
- 9800,
- 9810,
- 9400,
- 9500,
- 9600,
- 10400,
- 10410,
+ 9700, // broken in osx
+ 9710, // broken in osx
+ 9800, // broken in osx
+ 9810, // broken in osx
+ 9400, // broken in osx
+ 9500, // broken in osx
+ 9600, // broken in osx
+ 10400, // broken in osx
+ 10410, // broken in osx
10500,
10600,
- 10700,
+ 10700, // broken in osx
9000,
5200,
- 6800,
- 6600,
+ 6800, // broken in osx
+ 6600, // broken in osx
8200,
- 11300,
- 12700
+ 11300, // broken in osx
+ 12700 // broken in osx
};
/**
" -w, --workload-profile=NUM Enable a specific workload profile, see references below",
" -n, --kernel-accel=NUM Workload tuning: 1, 8, 40, 80, 160",
" -u, --kernel-loops=NUM Workload fine-tuning: 8 - 1024",
- #ifdef HAVE_HWMON
" --gpu-temp-disable Disable temperature and fanspeed readings and triggers",
+ #ifdef HAVE_HWMON
" --gpu-temp-abort=NUM Abort session if GPU temperature reaches NUM degrees celsius",
" --gpu-temp-retain=NUM Try to retain GPU temperature at NUM degrees celsius (AMD only)",
#ifdef HAVE_ADL
static void gidd_to_pw_t (hc_device_param_t *device_param, const u64 gidd, pw_t *pw)
{
- hc_clEnqueueReadBuffer (device_param->command_queue, device_param->d_pws_buf, CL_TRUE, gidd * sizeof (pw_t), sizeof (pw_t), pw, 0, NULL, NULL);
+ 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);
}
static void check_hash (hc_device_param_t *device_param, const uint salt_pos, const uint digest_pos)
plain_t plain;
- hc_clEnqueueReadBuffer (device_param->command_queue, device_param->d_plain_bufs, CL_TRUE, idx * sizeof (plain_t), sizeof (plain_t), &plain, 0, NULL, NULL);
+ hc_clEnqueueReadBuffer (data.ocl, device_param->command_queue, device_param->d_plain_bufs, CL_TRUE, idx * sizeof (plain_t), sizeof (plain_t), &plain, 0, NULL, NULL);
uint gidvid = plain.gidvid;
uint il_pos = plain.il_pos;
for (int i = 0, j = gidm; i < 16; i++, j++)
{
- plain_buf[i] = pw.hi1[0][j];
+ plain_buf[i] = pw.h.hi1[0][j];
}
plain_len = pw.pw_len;
for (int i = 0, j = gidm; i < 16; i++, j++)
{
- plain_buf[i] = pw.hi1[0][j];
+ plain_buf[i] = pw.h.hi1[0][j];
}
plain_len = pw.pw_len;
for (int i = 0, j = gidm; i < 16; i++, j++)
{
- plain_buf[i] = pw.hi1[0][j];
+ plain_buf[i] = pw.h.hi1[0][j];
}
plain_len = pw.pw_len;
for (int i = 0, j = gidm; i < 16; i++, j++)
{
- plain_buf[i] = pw.hi1[0][j];
+ plain_buf[i] = pw.h.hi1[0][j];
}
plain_len = pw.pw_len;
int found = 0;
- hc_clEnqueueReadBuffer (device_param->command_queue, device_param->d_result, CL_TRUE, 0, device_param->size_results, device_param->result, 0, NULL, NULL);
+ hc_clEnqueueReadBuffer (data.ocl, device_param->command_queue, device_param->d_result, CL_TRUE, 0, device_param->size_results, device_param->result, 0, NULL, NULL);
for (uint i = 0; i < KERNEL_THREADS; i++) if (device_param->result[i] == 1) found = 1;
log_info_nn ("");
- hc_clEnqueueReadBuffer (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);
+ hc_clEnqueueReadBuffer (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);
uint cpt_cracked = 0;
memset (data.digests_shown_tmp, 0, salt_buf->digests_cnt * sizeof (uint));
- hc_clEnqueueWriteBuffer (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);
+ 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);
}
memset (device_param->result, 0, device_param->size_results);
- hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_result, CL_TRUE, 0, device_param->size_results, device_param->result, 0, NULL, NULL);
+ hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_result, CL_TRUE, 0, device_param->size_results, device_param->result, 0, NULL, NULL);
}
}
case KERN_RUN_3: kernel = device_param->kernel3; break;
}
- hc_clSetKernelArg (kernel, 21, sizeof (cl_uint), device_param->kernel_params[21]);
- hc_clSetKernelArg (kernel, 22, sizeof (cl_uint), device_param->kernel_params[22]);
- hc_clSetKernelArg (kernel, 23, sizeof (cl_uint), device_param->kernel_params[23]);
- hc_clSetKernelArg (kernel, 24, sizeof (cl_uint), device_param->kernel_params[24]);
- hc_clSetKernelArg (kernel, 25, sizeof (cl_uint), device_param->kernel_params[25]);
- hc_clSetKernelArg (kernel, 26, sizeof (cl_uint), device_param->kernel_params[26]);
- hc_clSetKernelArg (kernel, 27, sizeof (cl_uint), device_param->kernel_params[27]);
- hc_clSetKernelArg (kernel, 28, sizeof (cl_uint), device_param->kernel_params[28]);
- hc_clSetKernelArg (kernel, 29, sizeof (cl_uint), device_param->kernel_params[29]);
- hc_clSetKernelArg (kernel, 30, sizeof (cl_uint), device_param->kernel_params[30]);
- hc_clSetKernelArg (kernel, 31, sizeof (cl_uint), device_param->kernel_params[31]);
+ hc_clSetKernelArg (data.ocl, kernel, 21, sizeof (cl_uint), device_param->kernel_params[21]);
+ hc_clSetKernelArg (data.ocl, kernel, 22, sizeof (cl_uint), device_param->kernel_params[22]);
+ hc_clSetKernelArg (data.ocl, kernel, 23, sizeof (cl_uint), device_param->kernel_params[23]);
+ hc_clSetKernelArg (data.ocl, kernel, 24, sizeof (cl_uint), device_param->kernel_params[24]);
+ hc_clSetKernelArg (data.ocl, kernel, 25, sizeof (cl_uint), device_param->kernel_params[25]);
+ hc_clSetKernelArg (data.ocl, kernel, 26, sizeof (cl_uint), device_param->kernel_params[26]);
+ hc_clSetKernelArg (data.ocl, kernel, 27, sizeof (cl_uint), device_param->kernel_params[27]);
+ hc_clSetKernelArg (data.ocl, kernel, 28, sizeof (cl_uint), device_param->kernel_params[28]);
+ hc_clSetKernelArg (data.ocl, kernel, 29, sizeof (cl_uint), device_param->kernel_params[29]);
+ hc_clSetKernelArg (data.ocl, kernel, 30, sizeof (cl_uint), device_param->kernel_params[30]);
+ hc_clSetKernelArg (data.ocl, kernel, 31, sizeof (cl_uint), device_param->kernel_params[31]);
if ((data.opts_type & OPTS_TYPE_PT_BITSLICE) && (data.attack_mode == ATTACK_MODE_BF))
{
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 (device_param->command_queue, kernel, 2, NULL, global_work_size, local_work_size, 0, NULL, NULL);
+ hc_clEnqueueNDRangeKernel (data.ocl, device_param->command_queue, kernel, 2, NULL, global_work_size, local_work_size, 0, NULL, NULL);
}
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 (device_param->command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);
+ hc_clEnqueueNDRangeKernel (data.ocl, device_param->command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);
}
- hc_clFlush (device_param->command_queue);
+ hc_clFlush (data.ocl, device_param->command_queue);
- hc_clFinish (device_param->command_queue);
+ hc_clFinish (data.ocl, device_param->command_queue);
}
static void run_kernel_mp (const uint kern_run, hc_device_param_t *device_param, const uint num)
switch (kern_run)
{
- case KERN_RUN_MP: hc_clSetKernelArg (kernel, 3, sizeof (cl_ulong), device_param->kernel_params_mp[3]);
- hc_clSetKernelArg (kernel, 4, sizeof (cl_uint), device_param->kernel_params_mp[4]);
- hc_clSetKernelArg (kernel, 5, sizeof (cl_uint), device_param->kernel_params_mp[5]);
- hc_clSetKernelArg (kernel, 6, sizeof (cl_uint), device_param->kernel_params_mp[6]);
- hc_clSetKernelArg (kernel, 7, sizeof (cl_uint), device_param->kernel_params_mp[7]);
- hc_clSetKernelArg (kernel, 8, sizeof (cl_uint), device_param->kernel_params_mp[8]);
+ 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]);
break;
- case KERN_RUN_MP_R: hc_clSetKernelArg (kernel, 3, sizeof (cl_ulong), device_param->kernel_params_mp_r[3]);
- hc_clSetKernelArg (kernel, 4, sizeof (cl_uint), device_param->kernel_params_mp_r[4]);
- hc_clSetKernelArg (kernel, 5, sizeof (cl_uint), device_param->kernel_params_mp_r[5]);
- hc_clSetKernelArg (kernel, 6, sizeof (cl_uint), device_param->kernel_params_mp_r[6]);
- hc_clSetKernelArg (kernel, 7, sizeof (cl_uint), device_param->kernel_params_mp_r[7]);
- hc_clSetKernelArg (kernel, 8, sizeof (cl_uint), device_param->kernel_params_mp_r[8]);
+ 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]);
break;
- case KERN_RUN_MP_L: hc_clSetKernelArg (kernel, 3, sizeof (cl_ulong), device_param->kernel_params_mp_l[3]);
- hc_clSetKernelArg (kernel, 4, sizeof (cl_uint), device_param->kernel_params_mp_l[4]);
- hc_clSetKernelArg (kernel, 5, sizeof (cl_uint), device_param->kernel_params_mp_l[5]);
- hc_clSetKernelArg (kernel, 6, sizeof (cl_uint), device_param->kernel_params_mp_l[6]);
- hc_clSetKernelArg (kernel, 7, sizeof (cl_uint), device_param->kernel_params_mp_l[7]);
- hc_clSetKernelArg (kernel, 8, sizeof (cl_uint), device_param->kernel_params_mp_l[8]);
- hc_clSetKernelArg (kernel, 9, sizeof (cl_uint), device_param->kernel_params_mp_l[9]);
+ 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]);
break;
}
const size_t global_work_size[3] = { num_elements, 1, 1 };
const size_t local_work_size[3] = { kernel_threads, 1, 1 };
- hc_clEnqueueNDRangeKernel (device_param->command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);
+ hc_clEnqueueNDRangeKernel (data.ocl, device_param->command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);
- hc_clFlush (device_param->command_queue);
+ hc_clFlush (data.ocl, device_param->command_queue);
- hc_clFinish (device_param->command_queue);
+ hc_clFinish (data.ocl, device_param->command_queue);
}
static void run_kernel_tb (hc_device_param_t *device_param, const uint num)
const size_t global_work_size[3] = { num_elements, 1, 1 };
const size_t local_work_size[3] = { kernel_threads, 1, 1 };
- hc_clEnqueueNDRangeKernel (device_param->command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);
+ hc_clEnqueueNDRangeKernel (data.ocl, device_param->command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);
- hc_clFlush (device_param->command_queue);
+ hc_clFlush (data.ocl, device_param->command_queue);
- hc_clFinish (device_param->command_queue);
+ hc_clFinish (data.ocl, device_param->command_queue);
}
static void run_kernel_tm (hc_device_param_t *device_param)
const size_t global_work_size[3] = { num_elements, 1, 1 };
const size_t local_work_size[3] = { kernel_threads, 1, 1 };
- hc_clEnqueueNDRangeKernel (device_param->command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);
+ hc_clEnqueueNDRangeKernel (data.ocl, device_param->command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);
- hc_clFlush (device_param->command_queue);
+ hc_clFlush (data.ocl, device_param->command_queue);
- hc_clFinish (device_param->command_queue);
+ hc_clFinish (data.ocl, device_param->command_queue);
}
static void run_kernel_amp (hc_device_param_t *device_param, const uint num)
cl_kernel kernel = device_param->kernel_amp;
- hc_clSetKernelArg (kernel, 5, sizeof (cl_uint), device_param->kernel_params_amp[5]);
- hc_clSetKernelArg (kernel, 6, sizeof (cl_uint), device_param->kernel_params_amp[6]);
+ 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]);
const size_t global_work_size[3] = { num_elements, 1, 1 };
const size_t local_work_size[3] = { kernel_threads, 1, 1 };
- hc_clEnqueueNDRangeKernel (device_param->command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);
+ hc_clEnqueueNDRangeKernel (data.ocl, device_param->command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);
- hc_clFlush (device_param->command_queue);
+ hc_clFlush (data.ocl, device_param->command_queue);
- hc_clFinish (device_param->command_queue);
+ hc_clFinish (data.ocl, device_param->command_queue);
}
static void run_kernel_bzero (hc_device_param_t *device_param, cl_mem buf, const uint size)
const cl_uchar zero = 0;
- hc_clEnqueueFillBuffer (device_param->command_queue, buf, &zero, sizeof (cl_uchar), 0, size, 0, NULL, NULL);
+ hc_clEnqueueFillBuffer (data.ocl, device_param->command_queue, buf, &zero, sizeof (cl_uchar), 0, size, 0, NULL, NULL);
}
else
{
const int fillsz = MIN (FILLSZ, left);
- hc_clEnqueueWriteBuffer (device_param->command_queue, buf, CL_TRUE, i, fillsz, tmp, 0, NULL, NULL);
+ hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, buf, CL_TRUE, i, fillsz, tmp, 0, NULL, NULL);
}
myfree (tmp);
{
if (data.attack_kern == ATTACK_KERN_STRAIGHT)
{
- hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_pws_buf, CL_TRUE, 0, pws_cnt * sizeof (pw_t), device_param->pws_buf, 0, NULL, NULL);
+ 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);
}
else if (data.attack_kern == ATTACK_KERN_COMBI)
{
- hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_pws_buf, CL_TRUE, 0, pws_cnt * sizeof (pw_t), device_param->pws_buf, 0, NULL, NULL);
+ 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);
}
else if (data.attack_kern == ATTACK_KERN_BF)
{
if (data.attack_mode == ATTACK_MODE_STRAIGHT)
{
- hc_clEnqueueCopyBuffer (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);
+ 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);
}
else if (data.attack_mode == ATTACK_MODE_COMBI)
{
- hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_combs_c, CL_TRUE, 0, innerloop_left * sizeof (comb_t), device_param->combs_buf, 0, NULL, NULL);
+ hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_combs_c, CL_TRUE, 0, innerloop_left * sizeof (comb_t), device_param->combs_buf, 0, NULL, NULL);
}
else if (data.attack_mode == ATTACK_MODE_BF)
{
- hc_clEnqueueCopyBuffer (device_param->command_queue, device_param->d_bfs, device_param->d_bfs_c, 0, 0, innerloop_left * sizeof (bf_t), 0, NULL, NULL);
+ 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);
}
else if (data.attack_mode == ATTACK_MODE_HYBRID1)
{
- hc_clEnqueueCopyBuffer (device_param->command_queue, device_param->d_combs, device_param->d_combs_c, 0, 0, innerloop_left * sizeof (comb_t), 0, NULL, NULL);
+ 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);
}
else if (data.attack_mode == ATTACK_MODE_HYBRID2)
{
- hc_clEnqueueCopyBuffer (device_param->command_queue, device_param->d_combs, device_param->d_combs_c, 0, 0, innerloop_left * sizeof (comb_t), 0, NULL, NULL);
+ 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 (data.attack_exec == ATTACK_EXEC_INSIDE_KERNEL)
run_kernel_tm (device_param);
- hc_clEnqueueCopyBuffer (device_param->command_queue, device_param->d_tm_c, device_param->d_bfs_c, 0, 0, size_tm, 0, NULL, NULL);
+ 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);
}
}
{
run_kernel (KERN_RUN_23, device_param, pws_cnt);
- hc_clEnqueueReadBuffer (device_param->command_queue, device_param->d_hooks, CL_TRUE, 0, device_param->size_hooks, device_param->hooks_buf, 0, NULL, NULL);
+ 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);
// do something with data
- hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_hooks, CL_TRUE, 0, device_param->size_hooks, device_param->hooks_buf, 0, NULL, NULL);
+ 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);
}
run_kernel (KERN_RUN_3, device_param, pws_cnt);
static void pw_transpose_to_hi1 (const pw_t *p1, pw_t *p2)
{
- memcpy (p2->hi1, p1->hi1, 64 * sizeof (uint));
+ memcpy (p2->h.hi1, p1->h.hi1, 64 * sizeof (uint));
}
static uint pw_add_to_hc1 (hc_device_param_t *device_param, const u8 *pw_buf, const uint pw_len)
uint cache_cnt = pw_cache->cnt;
- u8 *pw_hc1 = pw_cache->pw_buf.hc1[cache_cnt];
+ u8 *pw_hc1 = pw_cache->pw_buf.h.hc1[cache_cnt];
memcpy (pw_hc1, pw_buf, pw_len);
#ifndef OSX
char *cpu_affinity = NULL;
#endif
+ OCL_PTR *ocl = NULL;
char *opencl_devices = NULL;
char *opencl_platforms = NULL;
char *opencl_device_types = NULL;
uint workload_profile = WORKLOAD_PROFILE;
uint kernel_accel = KERNEL_ACCEL;
uint kernel_loops = KERNEL_LOOPS;
- #ifdef HAVE_HWMON
uint gpu_temp_disable = GPU_TEMP_DISABLE;
+ #ifdef HAVE_HWMON
uint gpu_temp_abort = GPU_TEMP_ABORT;
uint gpu_temp_retain = GPU_TEMP_RETAIN;
#ifdef HAVE_ADL
{"workload-profile", required_argument, 0, IDX_WORKLOAD_PROFILE},
{"kernel-accel", required_argument, 0, IDX_KERNEL_ACCEL},
{"kernel-loops", required_argument, 0, IDX_KERNEL_LOOPS},
- #ifdef HAVE_HWMON
{"gpu-temp-disable", no_argument, 0, IDX_GPU_TEMP_DISABLE},
+ #ifdef HAVE_HWMON
{"gpu-temp-abort", required_argument, 0, IDX_GPU_TEMP_ABORT},
{"gpu-temp-retain", required_argument, 0, IDX_GPU_TEMP_RETAIN},
#ifdef HAVE_ADL
kernel_accel_chgd = 1; break;
case IDX_KERNEL_LOOPS: kernel_loops = atoi (optarg);
kernel_loops_chgd = 1; break;
- #ifdef HAVE_HWMON
case IDX_GPU_TEMP_DISABLE: gpu_temp_disable = 1; break;
+ #ifdef HAVE_HWMON
case IDX_GPU_TEMP_ABORT: gpu_temp_abort = atoi (optarg);
#ifdef HAVE_ADL
gpu_temp_abort_chgd = 1;
logfile_top_uint (force);
logfile_top_uint (kernel_accel);
logfile_top_uint (kernel_loops);
+ logfile_top_uint (gpu_temp_disable);
#ifdef HAVE_HWMON
logfile_top_uint (gpu_temp_abort);
- logfile_top_uint (gpu_temp_disable);
logfile_top_uint (gpu_temp_retain);
#endif
logfile_top_uint (hash_mode);
logfile_top_string (session);
logfile_top_string (truecrypt_keyfiles);
+ /**
+ * Init OpenCL library loader
+ */
+
+ if (keyspace == 0)
+ {
+ ocl = (OCL_PTR *) mymalloc (sizeof (OCL_PTR));
+
+ ocl_init(ocl);
+
+ data.ocl = ocl;
+ }
+
/**
* OpenCL platform selection
*/
kernel_accel = 16;
break;
case 6211: kernel_loops = ROUNDS_TRUECRYPT_2K;
+ #ifndef OSX
kernel_accel = 64;
+ #endif
break;
case 6212: kernel_loops = ROUNDS_TRUECRYPT_2K;
kernel_accel = 32;
kernel_accel = 8;
break;
case 6241: kernel_loops = ROUNDS_TRUECRYPT_1K;
+ #ifndef OSX
kernel_accel = 128;
+ #endif
break;
case 6242: kernel_loops = ROUNDS_TRUECRYPT_1K;
kernel_accel = 64;
kernel_accel = 8;
break;
case 7200: kernel_loops = ROUNDS_GRUB;
+ #ifndef OSX
kernel_accel = 16;
+ #endif
break;
case 7400: kernel_loops = ROUNDS_SHA256CRYPT;
kernel_accel = 8;
cl_uint platform_devices_cnt;
- hc_clGetPlatformIDs (CL_PLATFORMS_MAX, platforms, &platforms_cnt);
-
- if (platforms_cnt == 0)
+ if (keyspace == 0)
{
- log_error ("ERROR: No OpenCL compatible platform found");
+ hc_clGetPlatformIDs (data.ocl, CL_PLATFORMS_MAX, platforms, &platforms_cnt);
- return (-1);
+ if (platforms_cnt == 0)
+ {
+ log_error ("ERROR: No OpenCL compatible platform found");
+
+ return (-1);
+ }
}
/**
char platform_vendor[INFOSZ] = { 0 };
- hc_clGetPlatformInfo (platform, CL_PLATFORM_VENDOR, sizeof (platform_vendor), platform_vendor, NULL);
+ hc_clGetPlatformInfo (data.ocl, platform, CL_PLATFORM_VENDOR, sizeof (platform_vendor), platform_vendor, NULL);
#ifdef HAVE_HWMON
#if defined(HAVE_NVML) || defined(HAVE_NVAPI)
cl_platform_id platform = platforms[platform_id];
- hc_clGetDeviceIDs (platform, CL_DEVICE_TYPE_ALL, DEVICES_MAX, platform_devices, &platform_devices_cnt);
+ hc_clGetDeviceIDs (data.ocl, platform, CL_DEVICE_TYPE_ALL, DEVICES_MAX, platform_devices, &platform_devices_cnt);
for (uint platform_devices_id = 0; platform_devices_id < platform_devices_cnt; platform_devices_id++)
{
cl_device_type device_type;
- hc_clGetDeviceInfo (device_param->device, CL_DEVICE_TYPE, sizeof (device_type), &device_type, NULL);
+ hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_TYPE, sizeof (device_type), &device_type, NULL);
device_type &= ~CL_DEVICE_TYPE_DEFAULT;
cl_uint vendor_id = 0;
- hc_clGetDeviceInfo (device_param->device, CL_DEVICE_VENDOR_ID, sizeof (vendor_id), &vendor_id, NULL);
+ hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_VENDOR_ID, sizeof (vendor_id), &vendor_id, NULL);
device_param->vendor_id = vendor_id;
char *device_name = (char *) mymalloc (INFOSZ);
- hc_clGetDeviceInfo (device_param->device, CL_DEVICE_NAME, INFOSZ, device_name, NULL);
+ hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_NAME, INFOSZ, device_name, NULL);
device_param->device_name = device_name;
char *device_version = (char *) mymalloc (INFOSZ);
- hc_clGetDeviceInfo (device_param->device, CL_DEVICE_VERSION, INFOSZ, device_version, NULL);
+ hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_VERSION, INFOSZ, device_version, NULL);
device_param->device_version = device_version;
if (opencl_vector_width == OPENCL_VECTOR_WIDTH)
{
- hc_clGetDeviceInfo (device_param->device, CL_DEVICE_NATIVE_VECTOR_WIDTH_INT, sizeof (vector_width), &vector_width, NULL);
+ hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_NATIVE_VECTOR_WIDTH_INT, sizeof (vector_width), &vector_width, NULL);
if ((vendor_id == VENDOR_ID_NV) && (strstr (device_name, " Ti") || strstr (device_name, " TI")))
{
cl_uint device_processors;
- hc_clGetDeviceInfo (device_param->device, CL_DEVICE_MAX_COMPUTE_UNITS, sizeof (device_processors), &device_processors, NULL);
+ hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_MAX_COMPUTE_UNITS, sizeof (device_processors), &device_processors, NULL);
device_param->device_processors = device_processors;
cl_ulong device_maxmem_alloc;
- hc_clGetDeviceInfo (device_param->device, CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof (device_maxmem_alloc), &device_maxmem_alloc, NULL);
+ hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof (device_maxmem_alloc), &device_maxmem_alloc, NULL);
device_param->device_maxmem_alloc = device_maxmem_alloc;
cl_ulong device_global_mem;
- hc_clGetDeviceInfo (device_param->device, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof (device_global_mem), &device_global_mem, NULL);
+ hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof (device_global_mem), &device_global_mem, NULL);
device_param->device_global_mem = device_global_mem;
cl_uint device_maxclock_frequency;
- hc_clGetDeviceInfo (device_param->device, CL_DEVICE_MAX_CLOCK_FREQUENCY, sizeof (device_maxclock_frequency), &device_maxclock_frequency, NULL);
+ hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_MAX_CLOCK_FREQUENCY, sizeof (device_maxclock_frequency), &device_maxclock_frequency, NULL);
device_param->device_maxclock_frequency = device_maxclock_frequency;
char *driver_version = (char *) mymalloc (INFOSZ);
- hc_clGetDeviceInfo (device_param->device, CL_DRIVER_VERSION, INFOSZ, driver_version, NULL);
+ hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DRIVER_VERSION, INFOSZ, driver_version, NULL);
device_param->driver_version = driver_version;
#define CL_DEVICE_WAVEFRONT_WIDTH_AMD 0x4043
- hc_clGetDeviceInfo (device_param->device, CL_DEVICE_WAVEFRONT_WIDTH_AMD, sizeof (device_processor_cores), &device_processor_cores, NULL);
+ hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_WAVEFRONT_WIDTH_AMD, sizeof (device_processor_cores), &device_processor_cores, NULL);
device_param->device_processor_cores = device_processor_cores;
}
#define CL_DEVICE_KERNEL_EXEC_TIMEOUT_NV 0x4005
- hc_clGetDeviceInfo (device_param->device, CL_DEVICE_KERNEL_EXEC_TIMEOUT_NV, sizeof (kernel_exec_timeout), &kernel_exec_timeout, NULL);
+ hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_KERNEL_EXEC_TIMEOUT_NV, sizeof (kernel_exec_timeout), &kernel_exec_timeout, NULL);
device_param->kernel_exec_timeout = kernel_exec_timeout;
#define CL_DEVICE_WARP_SIZE_NV 0x4003
- hc_clGetDeviceInfo (device_param->device, CL_DEVICE_WARP_SIZE_NV, sizeof (device_processor_cores), &device_processor_cores, NULL);
+ hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_WARP_SIZE_NV, sizeof (device_processor_cores), &device_processor_cores, NULL);
device_param->device_processor_cores = device_processor_cores;
#define CL_DEVICE_COMPUTE_CAPABILITY_MAJOR_NV 0x4000
#define CL_DEVICE_COMPUTE_CAPABILITY_MINOR_NV 0x4001
- hc_clGetDeviceInfo (device_param->device, CL_DEVICE_COMPUTE_CAPABILITY_MINOR_NV, sizeof (sm_minor), &sm_minor, NULL);
- hc_clGetDeviceInfo (device_param->device, CL_DEVICE_COMPUTE_CAPABILITY_MAJOR_NV, sizeof (sm_major), &sm_major, NULL);
+ 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);
device_param->sm_minor = sm_minor;
device_param->sm_major = sm_major;
}
}
- if (devices_active == 0)
+ if (keyspace == 0 && devices_active == 0)
{
log_error ("ERROR: No devices found/left");
#endif // HAVE_ADK
#endif // HAVE_HWMON
+ #ifdef OSX
+ /*
+ * List of OSX kernel to fix
+ */
+ if ((hash_mode == 6100) || (hash_mode == 6900) || (hash_mode == 11700) || (hash_mode == 11800) || (hash_mode == 5300) || \
+ (hash_mode == 5400) || (hash_mode == 5500) || (hash_mode == 5600) || (hash_mode == 7500) || (hash_mode == 11100) || \
+ (hash_mode == 11400) || (hash_mode == 2611) || (hash_mode == 2711) || (hash_mode == 2811) || (hash_mode == 8400) || \
+ (hash_mode == 2612) || (hash_mode == 3711) || (hash_mode == 7600) || (hash_mode == 3100) || (hash_mode == 8000) || \
+ (hash_mode == 12600) || (hash_mode == 3000) || (hash_mode == 1500) || (hash_mode == 12400) || (hash_mode == 8500) || \
+ (hash_mode == 8600) || (hash_mode == 8700) || (hash_mode == 9100) || (hash_mode == 11600) || (hash_mode == 12500) || \
+ (hash_mode == 6231) || (hash_mode == 8800) || (hash_mode == 9700) || (hash_mode == 9710) || (hash_mode == 9800) || \
+ (hash_mode == 9810) || (hash_mode == 9400) || (hash_mode == 9500) || (hash_mode == 9600) || (hash_mode == 10400) || \
+ (hash_mode == 10410) || (hash_mode == 10700) || (hash_mode == 6800) || (hash_mode == 6600) || (hash_mode == 11300) || \
+ (hash_mode == 12700))
+ {
+ if (force == 0)
+ {
+ log_info ("");
+ log_info ("Warning: Hash mode %d is not stable in OSX.", hash_mode);
+ log_info ("You can use --force to override this but do not post error reports if you do so");
+ log_info ("");
+
+ continue;
+ }
+ }
+ #endif
+
+ #ifdef DEBUG
+ if (benchmark == 1) log_info ("Hashmode: %d", data.hash_mode);
+ #endif
+
uint kernel_blocks_all = 0;
for (uint device_id = 0; device_id < devices_cnt; device_id++)
* create context for each device
*/
- device_param->context = hc_clCreateContext (NULL, 1, &device_param->device, NULL, NULL);
+ device_param->context = hc_clCreateContext (data.ocl, NULL, 1, &device_param->device, NULL, NULL);
/**
* 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 (device_param->context, device_param->device, 0);
+ device_param->command_queue = hc_clCreateCommandQueue (data.ocl, device_param->context, device_param->device, 0);
/**
* create input buffers on device
load_kernel (source_file, 1, kernel_lengths, kernel_sources);
- device_param->program = hc_clCreateProgramWithSource (device_param->context, 1, (const char **) kernel_sources, NULL);
+ device_param->program = hc_clCreateProgramWithSource (data.ocl, device_param->context, 1, (const char **) kernel_sources, NULL);
- hc_clBuildProgram (device_param->program, 1, &device_param->device, build_opts, NULL, NULL);
+ hc_clBuildProgram (data.ocl, device_param->program, 1, &device_param->device, build_opts, NULL, NULL);
size_t binary_size;
- clGetProgramInfo (device_param->program, CL_PROGRAM_BINARY_SIZES, sizeof (size_t), &binary_size, NULL);
+ hc_clGetProgramInfo (data.ocl, device_param->program, CL_PROGRAM_BINARY_SIZES, sizeof (size_t), &binary_size, NULL);
u8 *binary = (u8 *) mymalloc (binary_size);
- clGetProgramInfo (device_param->program, CL_PROGRAM_BINARIES, sizeof (binary), &binary, NULL);
+ hc_clGetProgramInfo (data.ocl, device_param->program, CL_PROGRAM_BINARIES, sizeof (binary), &binary, NULL);
writeProgramBin (cached_file, binary, binary_size);
load_kernel (cached_file, 1, kernel_lengths, kernel_sources);
- device_param->program = hc_clCreateProgramWithBinary (device_param->context, 1, &device_param->device, kernel_lengths, (const u8 **) kernel_sources, NULL);
+ device_param->program = hc_clCreateProgramWithBinary (data.ocl, device_param->context, 1, &device_param->device, kernel_lengths, (const u8 **) kernel_sources, NULL);
- hc_clBuildProgram (device_param->program, 1, &device_param->device, build_opts, NULL, NULL);
+ hc_clBuildProgram (data.ocl, device_param->program, 1, &device_param->device, build_opts, NULL, NULL);
}
}
else
load_kernel (source_file, 1, kernel_lengths, kernel_sources);
- device_param->program = hc_clCreateProgramWithSource (device_param->context, 1, (const char **) kernel_sources, NULL);
+ device_param->program = hc_clCreateProgramWithSource (data.ocl, device_param->context, 1, (const char **) kernel_sources, NULL);
if (force_jit_compilation == 1500)
{
snprintf (build_opts, sizeof (build_opts) - 1, "%s -DSCRYPT_N=%d -DSCRYPT_R=%d -DSCRYPT_P=%d -DSCRYPT_TMTO=%d", build_opts, data.salts_buf[0].scrypt_N, data.salts_buf[0].scrypt_r, data.salts_buf[0].scrypt_p, 1 << data.salts_buf[0].scrypt_tmto);
}
- hc_clBuildProgram (device_param->program, 1, &device_param->device, build_opts, NULL, NULL);
+ hc_clBuildProgram (data.ocl, device_param->program, 1, &device_param->device, build_opts, NULL, NULL);
}
local_free (kernel_lengths);
size_t ret_val_size = 0;
- clGetProgramBuildInfo (device_param->program, device_param->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size);
+ hc_clGetProgramBuildInfo (data.ocl, device_param->program, device_param->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size);
if (ret_val_size > 2)
{
char *build_log = (char *) mymalloc (ret_val_size + 1);
- clGetProgramBuildInfo (device_param->program, device_param->device, CL_PROGRAM_BUILD_LOG, ret_val_size, build_log, NULL);
+ hc_clGetProgramBuildInfo (data.ocl, device_param->program, device_param->device, CL_PROGRAM_BUILD_LOG, ret_val_size, build_log, NULL);
puts (build_log);
load_kernel (source_file, 1, kernel_lengths, kernel_sources);
- device_param->program_mp = hc_clCreateProgramWithSource (device_param->context, 1, (const char **) kernel_sources, NULL);
+ device_param->program_mp = hc_clCreateProgramWithSource (data.ocl, device_param->context, 1, (const char **) kernel_sources, NULL);
- hc_clBuildProgram (device_param->program_mp, 1, &device_param->device, build_opts, NULL, NULL);
+ hc_clBuildProgram (data.ocl, device_param->program_mp, 1, &device_param->device, build_opts, NULL, NULL);
size_t binary_size;
- clGetProgramInfo (device_param->program_mp, CL_PROGRAM_BINARY_SIZES, sizeof (size_t), &binary_size, NULL);
+ hc_clGetProgramInfo (data.ocl, device_param->program_mp, CL_PROGRAM_BINARY_SIZES, sizeof (size_t), &binary_size, NULL);
u8 *binary = (u8 *) mymalloc (binary_size);
- clGetProgramInfo (device_param->program_mp, CL_PROGRAM_BINARIES, sizeof (binary), &binary, NULL);
+ hc_clGetProgramInfo (data.ocl, device_param->program_mp, CL_PROGRAM_BINARIES, sizeof (binary), &binary, NULL);
writeProgramBin (cached_file, binary, binary_size);
load_kernel (cached_file, 1, kernel_lengths, kernel_sources);
- device_param->program_mp = hc_clCreateProgramWithBinary (device_param->context, 1, &device_param->device, kernel_lengths, (const u8 **) kernel_sources, NULL);
+ device_param->program_mp = hc_clCreateProgramWithBinary (data.ocl, device_param->context, 1, &device_param->device, kernel_lengths, (const u8 **) kernel_sources, NULL);
- hc_clBuildProgram (device_param->program_mp, 1, &device_param->device, build_opts, NULL, NULL);
+ hc_clBuildProgram (data.ocl, device_param->program_mp, 1, &device_param->device, build_opts, NULL, NULL);
}
local_free (kernel_lengths);
size_t ret_val_size = 0;
- clGetProgramBuildInfo (device_param->program_mp, device_param->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size);
+ hc_clGetProgramBuildInfo (data.ocl, device_param->program_mp, device_param->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size);
if (ret_val_size > 2)
{
char *build_log = (char *) mymalloc (ret_val_size + 1);
- clGetProgramBuildInfo (device_param->program_mp, device_param->device, CL_PROGRAM_BUILD_LOG, ret_val_size, build_log, NULL);
+ hc_clGetProgramBuildInfo (data.ocl, device_param->program_mp, device_param->device, CL_PROGRAM_BUILD_LOG, ret_val_size, build_log, NULL);
puts (build_log);
load_kernel (source_file, 1, kernel_lengths, kernel_sources);
- device_param->program_amp = hc_clCreateProgramWithSource (device_param->context, 1, (const char **) kernel_sources, NULL);
+ device_param->program_amp = hc_clCreateProgramWithSource (data.ocl, device_param->context, 1, (const char **) kernel_sources, NULL);
- hc_clBuildProgram (device_param->program_amp, 1, &device_param->device, build_opts, NULL, NULL);
+ hc_clBuildProgram (data.ocl, device_param->program_amp, 1, &device_param->device, build_opts, NULL, NULL);
size_t binary_size;
- clGetProgramInfo (device_param->program_amp, CL_PROGRAM_BINARY_SIZES, sizeof (size_t), &binary_size, NULL);
+ hc_clGetProgramInfo (data.ocl, device_param->program_amp, CL_PROGRAM_BINARY_SIZES, sizeof (size_t), &binary_size, NULL);
u8 *binary = (u8 *) mymalloc (binary_size);
- clGetProgramInfo (device_param->program_amp, CL_PROGRAM_BINARIES, sizeof (binary), &binary, NULL);
+ hc_clGetProgramInfo (data.ocl, device_param->program_amp, CL_PROGRAM_BINARIES, sizeof (binary), &binary, NULL);
writeProgramBin (cached_file, binary, binary_size);
load_kernel (cached_file, 1, kernel_lengths, kernel_sources);
- device_param->program_amp = hc_clCreateProgramWithBinary (device_param->context, 1, &device_param->device, kernel_lengths, (const u8 **) kernel_sources, NULL);
+ device_param->program_amp = hc_clCreateProgramWithBinary (data.ocl, device_param->context, 1, &device_param->device, kernel_lengths, (const u8 **) kernel_sources, NULL);
- hc_clBuildProgram (device_param->program_amp, 1, &device_param->device, build_opts, NULL, NULL);
+ hc_clBuildProgram (data.ocl, device_param->program_amp, 1, &device_param->device, build_opts, NULL, NULL);
}
local_free (kernel_lengths);
size_t ret_val_size = 0;
- clGetProgramBuildInfo (device_param->program_amp, device_param->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size);
+ hc_clGetProgramBuildInfo (data.ocl, device_param->program_amp, device_param->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size);
if (ret_val_size > 2)
{
char *build_log = (char *) mymalloc (ret_val_size + 1);
- clGetProgramBuildInfo (device_param->program_amp, device_param->device, CL_PROGRAM_BUILD_LOG, ret_val_size, build_log, NULL);
+ hc_clGetProgramBuildInfo (data.ocl, device_param->program_amp, device_param->device, CL_PROGRAM_BUILD_LOG, ret_val_size, build_log, NULL);
puts (build_log);
* global buffers
*/
- device_param->d_pws_buf = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, size_pws, NULL);
- device_param->d_pws_amp_buf = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, size_pws, NULL);
- device_param->d_tmps = hc_clCreateBuffer (device_param->context, CL_MEM_READ_WRITE, size_tmps, NULL);
- device_param->d_hooks = hc_clCreateBuffer (device_param->context, CL_MEM_READ_WRITE, size_hooks, NULL);
- device_param->d_bitmap_s1_a = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL);
- device_param->d_bitmap_s1_b = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL);
- device_param->d_bitmap_s1_c = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL);
- device_param->d_bitmap_s1_d = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL);
- device_param->d_bitmap_s2_a = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL);
- device_param->d_bitmap_s2_b = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL);
- device_param->d_bitmap_s2_c = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL);
- device_param->d_bitmap_s2_d = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL);
- device_param->d_plain_bufs = hc_clCreateBuffer (device_param->context, CL_MEM_READ_WRITE, size_plains, NULL);
- device_param->d_digests_buf = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, size_digests, NULL);
- device_param->d_digests_shown = hc_clCreateBuffer (device_param->context, CL_MEM_READ_WRITE, size_shown, NULL);
- device_param->d_salt_bufs = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, size_salts, NULL);
- device_param->d_result = hc_clCreateBuffer (device_param->context, CL_MEM_READ_WRITE, size_results, NULL);
- device_param->d_scryptV_buf = hc_clCreateBuffer (device_param->context, CL_MEM_READ_WRITE, size_scryptV, NULL);
-
- hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_bitmap_s1_a, CL_TRUE, 0, bitmap_size, bitmap_s1_a, 0, NULL, NULL);
- hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_bitmap_s1_b, CL_TRUE, 0, bitmap_size, bitmap_s1_b, 0, NULL, NULL);
- hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_bitmap_s1_c, CL_TRUE, 0, bitmap_size, bitmap_s1_c, 0, NULL, NULL);
- hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_bitmap_s1_d, CL_TRUE, 0, bitmap_size, bitmap_s1_d, 0, NULL, NULL);
- hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_bitmap_s2_a, CL_TRUE, 0, bitmap_size, bitmap_s2_a, 0, NULL, NULL);
- hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_bitmap_s2_b, CL_TRUE, 0, bitmap_size, bitmap_s2_b, 0, NULL, NULL);
- hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_bitmap_s2_c, CL_TRUE, 0, bitmap_size, bitmap_s2_c, 0, NULL, NULL);
- hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_bitmap_s2_d, CL_TRUE, 0, bitmap_size, bitmap_s2_d, 0, NULL, NULL);
- hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_digests_buf, CL_TRUE, 0, size_digests, data.digests_buf, 0, NULL, NULL);
- hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_digests_shown, CL_TRUE, 0, size_shown, data.digests_shown, 0, NULL, NULL);
- hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_salt_bufs, CL_TRUE, 0, size_salts, data.salts_buf, 0, NULL, NULL);
+ device_param->d_pws_buf = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_pws, NULL);
+ device_param->d_pws_amp_buf = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_pws, NULL);
+ device_param->d_tmps = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE, size_tmps, NULL);
+ device_param->d_hooks = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE, size_hooks, NULL);
+ device_param->d_bitmap_s1_a = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL);
+ device_param->d_bitmap_s1_b = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL);
+ device_param->d_bitmap_s1_c = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL);
+ device_param->d_bitmap_s1_d = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL);
+ device_param->d_bitmap_s2_a = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL);
+ device_param->d_bitmap_s2_b = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL);
+ device_param->d_bitmap_s2_c = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL);
+ device_param->d_bitmap_s2_d = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL);
+ device_param->d_plain_bufs = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE, size_plains, NULL);
+ device_param->d_digests_buf = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_digests, NULL);
+ device_param->d_digests_shown = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE, size_shown, NULL);
+ device_param->d_salt_bufs = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_salts, NULL);
+ device_param->d_result = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE, size_results, NULL);
+ device_param->d_scryptV_buf = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE, size_scryptV, NULL);
+
+ hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s1_a, CL_TRUE, 0, bitmap_size, bitmap_s1_a, 0, NULL, NULL);
+ hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s1_b, CL_TRUE, 0, bitmap_size, bitmap_s1_b, 0, NULL, NULL);
+ hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s1_c, CL_TRUE, 0, bitmap_size, bitmap_s1_c, 0, NULL, NULL);
+ hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s1_d, CL_TRUE, 0, bitmap_size, bitmap_s1_d, 0, NULL, NULL);
+ hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s2_a, CL_TRUE, 0, bitmap_size, bitmap_s2_a, 0, NULL, NULL);
+ hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s2_b, CL_TRUE, 0, bitmap_size, bitmap_s2_b, 0, NULL, NULL);
+ hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s2_c, CL_TRUE, 0, bitmap_size, bitmap_s2_c, 0, NULL, NULL);
+ hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s2_d, CL_TRUE, 0, bitmap_size, bitmap_s2_d, 0, NULL, NULL);
+ hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_digests_buf, CL_TRUE, 0, size_digests, data.digests_buf, 0, NULL, NULL);
+ hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_digests_shown, CL_TRUE, 0, size_shown, data.digests_shown, 0, NULL, NULL);
+ hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_salt_bufs, CL_TRUE, 0, size_salts, data.salts_buf, 0, NULL, NULL);
run_kernel_bzero (device_param, device_param->d_pws_buf, size_pws);
run_kernel_bzero (device_param, device_param->d_pws_amp_buf, size_pws);
if (attack_kern == ATTACK_KERN_STRAIGHT)
{
- device_param->d_rules = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, size_rules, NULL);
- device_param->d_rules_c = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, size_rules_c, NULL);
+ 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);
- hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_rules, CL_TRUE, 0, size_rules, kernel_rules_buf, 0, NULL, NULL);
+ hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_rules, CL_TRUE, 0, size_rules, kernel_rules_buf, 0, NULL, NULL);
run_kernel_bzero (device_param, device_param->d_rules_c, size_rules_c);
}
else if (attack_kern == ATTACK_KERN_COMBI)
{
- device_param->d_combs = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, size_combs, NULL);
- device_param->d_combs_c = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, size_combs, NULL);
- device_param->d_root_css_buf = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, size_root_css, NULL);
- device_param->d_markov_css_buf = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, size_markov_css, NULL);
+ 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);
run_kernel_bzero (device_param, device_param->d_combs, size_combs);
run_kernel_bzero (device_param, device_param->d_combs_c, size_combs);
}
else if (attack_kern == ATTACK_KERN_BF)
{
- device_param->d_bfs = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, size_bfs, NULL);
- device_param->d_bfs_c = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, size_bfs, NULL);
- device_param->d_tm_c = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, size_tm, NULL);
- device_param->d_root_css_buf = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, size_root_css, NULL);
- device_param->d_markov_css_buf = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, size_markov_css, NULL);
+ 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);
run_kernel_bzero (device_param, device_param->d_bfs, size_bfs);
run_kernel_bzero (device_param, device_param->d_bfs_c, size_bfs);
if (size_esalts)
{
- device_param->d_esalt_bufs = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, size_esalts, NULL);
+ device_param->d_esalt_bufs = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_esalts, NULL);
- hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_esalt_bufs, CL_TRUE, 0, size_esalts, data.esalts_buf, 0, NULL, NULL);
+ hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_esalt_bufs, CL_TRUE, 0, size_esalts, data.esalts_buf, 0, NULL, NULL);
}
/**
{
snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_s%02d", kern_type, 4);
- device_param->kernel1 = hc_clCreateKernel (device_param->program, kernel_name);
+ device_param->kernel1 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name);
snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_s%02d", kern_type, 8);
- device_param->kernel2 = hc_clCreateKernel (device_param->program, kernel_name);
+ device_param->kernel2 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name);
snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_s%02d", kern_type, 16);
- device_param->kernel3 = hc_clCreateKernel (device_param->program, kernel_name);
+ device_param->kernel3 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name);
}
else
{
snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_m%02d", kern_type, 4);
- device_param->kernel1 = hc_clCreateKernel (device_param->program, kernel_name);
+ device_param->kernel1 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name);
snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_m%02d", kern_type, 8);
- device_param->kernel2 = hc_clCreateKernel (device_param->program, kernel_name);
+ device_param->kernel2 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name);
snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_m%02d", kern_type, 16);
- device_param->kernel3 = hc_clCreateKernel (device_param->program, kernel_name);
+ device_param->kernel3 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name);
}
if (data.attack_mode == ATTACK_MODE_BF)
{
snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_tb", kern_type);
- device_param->kernel_tb = hc_clCreateKernel (device_param->program, kernel_name);
+ device_param->kernel_tb = hc_clCreateKernel (data.ocl, device_param->program, kernel_name);
snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_tm", kern_type);
- device_param->kernel_tm = hc_clCreateKernel (device_param->program, kernel_name);
+ device_param->kernel_tm = hc_clCreateKernel (data.ocl, device_param->program, kernel_name);
}
}
}
{
snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_init", kern_type);
- device_param->kernel1 = hc_clCreateKernel (device_param->program, kernel_name);
+ device_param->kernel1 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name);
snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_loop", kern_type);
- device_param->kernel2 = hc_clCreateKernel (device_param->program, kernel_name);
+ device_param->kernel2 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name);
snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_comp", kern_type);
- device_param->kernel3 = hc_clCreateKernel (device_param->program, kernel_name);
+ device_param->kernel3 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name);
if (opts_type & OPTS_TYPE_HOOK12)
{
snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_hook12", kern_type);
- device_param->kernel12 = hc_clCreateKernel (device_param->program, kernel_name);
+ device_param->kernel12 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name);
}
if (opts_type & OPTS_TYPE_HOOK23)
{
snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_hook23", kern_type);
- device_param->kernel23 = hc_clCreateKernel (device_param->program, kernel_name);
+ device_param->kernel23 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name);
}
}
for (uint i = 0; i <= 20; i++)
{
- hc_clSetKernelArg (device_param->kernel1, i, sizeof (cl_mem), device_param->kernel_params[i]);
- hc_clSetKernelArg (device_param->kernel2, i, sizeof (cl_mem), device_param->kernel_params[i]);
- hc_clSetKernelArg (device_param->kernel3, i, sizeof (cl_mem), device_param->kernel_params[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]);
- if (opts_type & OPTS_TYPE_HOOK12) hc_clSetKernelArg (device_param->kernel12, i, sizeof (cl_mem), device_param->kernel_params[i]);
- if (opts_type & OPTS_TYPE_HOOK23) hc_clSetKernelArg (device_param->kernel23, i, sizeof (cl_mem), device_param->kernel_params[i]);
+ 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]);
}
for (uint i = 21; i <= 31; i++)
{
- hc_clSetKernelArg (device_param->kernel1, i, sizeof (cl_uint), device_param->kernel_params[i]);
- hc_clSetKernelArg (device_param->kernel2, i, sizeof (cl_uint), device_param->kernel_params[i]);
- hc_clSetKernelArg (device_param->kernel3, i, sizeof (cl_uint), device_param->kernel_params[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]);
- if (opts_type & OPTS_TYPE_HOOK12) hc_clSetKernelArg (device_param->kernel12, i, sizeof (cl_uint), device_param->kernel_params[i]);
- if (opts_type & OPTS_TYPE_HOOK23) hc_clSetKernelArg (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 (attack_mode == ATTACK_MODE_BF)
{
- device_param->kernel_mp_l = hc_clCreateKernel (device_param->program_mp, "l_markov");
- device_param->kernel_mp_r = hc_clCreateKernel (device_param->program_mp, "r_markov");
+ 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");
if (opts_type & OPTS_TYPE_PT_BITSLICE)
{
- hc_clSetKernelArg (device_param->kernel_tb, 0, sizeof (cl_mem), device_param->kernel_params_tb[0]);
+ hc_clSetKernelArg (data.ocl, device_param->kernel_tb, 0, sizeof (cl_mem), device_param->kernel_params_tb[0]);
- hc_clSetKernelArg (device_param->kernel_tm, 0, sizeof (cl_mem), device_param->kernel_params_tm[0]);
- hc_clSetKernelArg (device_param->kernel_tm, 1, sizeof (cl_mem), device_param->kernel_params_tm[1]);
+ 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]);
}
}
else if (attack_mode == ATTACK_MODE_HYBRID1)
{
- device_param->kernel_mp = hc_clCreateKernel (device_param->program_mp, "C_markov");
+ device_param->kernel_mp = hc_clCreateKernel (data.ocl, device_param->program_mp, "C_markov");
}
else if (attack_mode == ATTACK_MODE_HYBRID2)
{
- device_param->kernel_mp = hc_clCreateKernel (device_param->program_mp, "C_markov");
+ device_param->kernel_mp = hc_clCreateKernel (data.ocl, device_param->program_mp, "C_markov");
}
if (attack_exec == ATTACK_EXEC_INSIDE_KERNEL)
}
else
{
- device_param->kernel_amp = hc_clCreateKernel (device_param->program_amp, "amp");
+ device_param->kernel_amp = hc_clCreateKernel (data.ocl, device_param->program_amp, "amp");
}
if (attack_exec == ATTACK_EXEC_INSIDE_KERNEL)
{
for (uint i = 0; i < 5; i++)
{
- hc_clSetKernelArg (device_param->kernel_amp, i, sizeof (cl_mem), device_param->kernel_params_amp[i]);
+ hc_clSetKernelArg (data.ocl, device_param->kernel_amp, i, sizeof (cl_mem), device_param->kernel_params_amp[i]);
}
for (uint i = 5; i < 7; i++)
{
- hc_clSetKernelArg (device_param->kernel_amp, i, sizeof (cl_uint), device_param->kernel_params_amp[i]);
+ hc_clSetKernelArg (data.ocl, device_param->kernel_amp, i, sizeof (cl_uint), device_param->kernel_params_amp[i]);
}
}
device_param->kernel_params_mp_buf32[7] = 0;
}
- for (uint i = 0; i < 3; i++) hc_clSetKernelArg (device_param->kernel_mp, i, sizeof (cl_mem), (void *) device_param->kernel_params_mp[i]);
- for (uint i = 3; i < 4; i++) hc_clSetKernelArg (device_param->kernel_mp, i, sizeof (cl_ulong), (void *) device_param->kernel_params_mp[i]);
- for (uint i = 4; i < 8; i++) hc_clSetKernelArg (device_param->kernel_mp, i, sizeof (cl_uint), (void *) device_param->kernel_params_mp[i]);
+ 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]);
- hc_clEnqueueWriteBuffer (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 (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);
}
}
else if (attack_mode == ATTACK_MODE_BF)
// some more resets:
- memset (device_param->pw_caches, 0, 64 * sizeof (pw_cache_t));
+ if (device_param->pw_caches) memset (device_param->pw_caches, 0, 64 * sizeof (pw_cache_t));
- memset (device_param->pws_buf, 0, device_param->size_pws);
+ if (device_param->pws_buf) memset (device_param->pws_buf, 0, device_param->size_pws);
device_param->pw_cnt = 0;
device_param->pws_cnt = 0;
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 (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 (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 (device_param->kernel_mp_l, i, sizeof (cl_uint), (void *) device_param->kernel_params_mp_l[i]);
+ 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]);
- for (uint i = 0; i < 3; i++) hc_clSetKernelArg (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 (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 (device_param->kernel_mp_r, i, sizeof (cl_uint), (void *) device_param->kernel_params_mp_r[i]);
+ 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]);
- hc_clEnqueueWriteBuffer (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 (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);
}
}
local_free (device_param->driver_version);
if (device_param->pws_buf) myfree (device_param->pws_buf);
- if (device_param->d_pws_buf) hc_clReleaseMemObject (device_param->d_pws_buf);
- if (device_param->d_pws_amp_buf) hc_clReleaseMemObject (device_param->d_pws_amp_buf);
- if (device_param->d_rules) hc_clReleaseMemObject (device_param->d_rules);
- if (device_param->d_rules_c) hc_clReleaseMemObject (device_param->d_rules_c);
- if (device_param->d_combs) hc_clReleaseMemObject (device_param->d_combs);
- if (device_param->d_combs_c) hc_clReleaseMemObject (device_param->d_combs_c);
- if (device_param->d_bfs) hc_clReleaseMemObject (device_param->d_bfs);
- if (device_param->d_bfs_c) hc_clReleaseMemObject (device_param->d_bfs_c);
- if (device_param->d_bitmap_s1_a) hc_clReleaseMemObject (device_param->d_bitmap_s1_a);
- if (device_param->d_bitmap_s1_b) hc_clReleaseMemObject (device_param->d_bitmap_s1_b);
- if (device_param->d_bitmap_s1_c) hc_clReleaseMemObject (device_param->d_bitmap_s1_c);
- if (device_param->d_bitmap_s1_d) hc_clReleaseMemObject (device_param->d_bitmap_s1_d);
- if (device_param->d_bitmap_s2_a) hc_clReleaseMemObject (device_param->d_bitmap_s2_a);
- if (device_param->d_bitmap_s2_b) hc_clReleaseMemObject (device_param->d_bitmap_s2_b);
- if (device_param->d_bitmap_s2_c) hc_clReleaseMemObject (device_param->d_bitmap_s2_c);
- if (device_param->d_bitmap_s2_d) hc_clReleaseMemObject (device_param->d_bitmap_s2_d);
- if (device_param->d_plain_bufs) hc_clReleaseMemObject (device_param->d_plain_bufs);
- if (device_param->d_digests_buf) hc_clReleaseMemObject (device_param->d_digests_buf);
- if (device_param->d_digests_shown) hc_clReleaseMemObject (device_param->d_digests_shown);
- if (device_param->d_salt_bufs) hc_clReleaseMemObject (device_param->d_salt_bufs);
- if (device_param->d_esalt_bufs) hc_clReleaseMemObject (device_param->d_esalt_bufs);
- if (device_param->d_tmps) hc_clReleaseMemObject (device_param->d_tmps);
- if (device_param->d_hooks) hc_clReleaseMemObject (device_param->d_hooks);
- if (device_param->d_result) hc_clReleaseMemObject (device_param->d_result);
- if (device_param->d_scryptV_buf) hc_clReleaseMemObject (device_param->d_scryptV_buf);
- if (device_param->d_root_css_buf) hc_clReleaseMemObject (device_param->d_root_css_buf);
- if (device_param->d_markov_css_buf) hc_clReleaseMemObject (device_param->d_markov_css_buf);
- if (device_param->d_tm_c) hc_clReleaseMemObject (device_param->d_tm_c);
-
- if (device_param->kernel1) hc_clReleaseKernel (device_param->kernel1);
- if (device_param->kernel12) hc_clReleaseKernel (device_param->kernel12);
- if (device_param->kernel2) hc_clReleaseKernel (device_param->kernel2);
- if (device_param->kernel23) hc_clReleaseKernel (device_param->kernel23);
- if (device_param->kernel3) hc_clReleaseKernel (device_param->kernel3);
- if (device_param->kernel_mp) hc_clReleaseKernel (device_param->kernel_mp);
- if (device_param->kernel_mp_l) hc_clReleaseKernel (device_param->kernel_mp_l);
- if (device_param->kernel_mp_r) hc_clReleaseKernel (device_param->kernel_mp_r);
- if (device_param->kernel_tb) hc_clReleaseKernel (device_param->kernel_tb);
- if (device_param->kernel_tm) hc_clReleaseKernel (device_param->kernel_tm);
- if (device_param->kernel_amp) hc_clReleaseKernel (device_param->kernel_amp);
-
- if (device_param->program) hc_clReleaseProgram (device_param->program);
- if (device_param->program_mp) hc_clReleaseProgram (device_param->program_mp);
- if (device_param->program_amp) hc_clReleaseProgram (device_param->program_amp);
-
- if (device_param->command_queue) hc_clReleaseCommandQueue (device_param->command_queue);
- if (device_param->context) hc_clReleaseContext (device_param->context);
+ if (device_param->d_pws_buf) hc_clReleaseMemObject (data.ocl, device_param->d_pws_buf);
+ if (device_param->d_pws_amp_buf) hc_clReleaseMemObject (data.ocl, device_param->d_pws_amp_buf);
+ if (device_param->d_rules) hc_clReleaseMemObject (data.ocl, device_param->d_rules);
+ if (device_param->d_rules_c) hc_clReleaseMemObject (data.ocl, device_param->d_rules_c);
+ if (device_param->d_combs) hc_clReleaseMemObject (data.ocl, device_param->d_combs);
+ if (device_param->d_combs_c) hc_clReleaseMemObject (data.ocl, device_param->d_combs_c);
+ if (device_param->d_bfs) hc_clReleaseMemObject (data.ocl, device_param->d_bfs);
+ if (device_param->d_bfs_c) hc_clReleaseMemObject (data.ocl, device_param->d_bfs_c);
+ if (device_param->d_bitmap_s1_a) hc_clReleaseMemObject (data.ocl, device_param->d_bitmap_s1_a);
+ if (device_param->d_bitmap_s1_b) hc_clReleaseMemObject (data.ocl, device_param->d_bitmap_s1_b);
+ if (device_param->d_bitmap_s1_c) hc_clReleaseMemObject (data.ocl, device_param->d_bitmap_s1_c);
+ if (device_param->d_bitmap_s1_d) hc_clReleaseMemObject (data.ocl, device_param->d_bitmap_s1_d);
+ if (device_param->d_bitmap_s2_a) hc_clReleaseMemObject (data.ocl, device_param->d_bitmap_s2_a);
+ if (device_param->d_bitmap_s2_b) hc_clReleaseMemObject (data.ocl, device_param->d_bitmap_s2_b);
+ if (device_param->d_bitmap_s2_c) hc_clReleaseMemObject (data.ocl, device_param->d_bitmap_s2_c);
+ if (device_param->d_bitmap_s2_d) hc_clReleaseMemObject (data.ocl, device_param->d_bitmap_s2_d);
+ if (device_param->d_plain_bufs) hc_clReleaseMemObject (data.ocl, device_param->d_plain_bufs);
+ if (device_param->d_digests_buf) hc_clReleaseMemObject (data.ocl, device_param->d_digests_buf);
+ if (device_param->d_digests_shown) hc_clReleaseMemObject (data.ocl, device_param->d_digests_shown);
+ if (device_param->d_salt_bufs) hc_clReleaseMemObject (data.ocl, device_param->d_salt_bufs);
+ if (device_param->d_esalt_bufs) hc_clReleaseMemObject (data.ocl, device_param->d_esalt_bufs);
+ if (device_param->d_tmps) hc_clReleaseMemObject (data.ocl, device_param->d_tmps);
+ if (device_param->d_hooks) hc_clReleaseMemObject (data.ocl, device_param->d_hooks);
+ if (device_param->d_result) hc_clReleaseMemObject (data.ocl, device_param->d_result);
+ if (device_param->d_scryptV_buf) hc_clReleaseMemObject (data.ocl, device_param->d_scryptV_buf);
+ if (device_param->d_root_css_buf) hc_clReleaseMemObject (data.ocl, device_param->d_root_css_buf);
+ if (device_param->d_markov_css_buf) hc_clReleaseMemObject (data.ocl, device_param->d_markov_css_buf);
+ if (device_param->d_tm_c) hc_clReleaseMemObject (data.ocl, device_param->d_tm_c);
+
+ if (device_param->kernel1) hc_clReleaseKernel (data.ocl, device_param->kernel1);
+ if (device_param->kernel12) hc_clReleaseKernel (data.ocl, device_param->kernel12);
+ if (device_param->kernel2) hc_clReleaseKernel (data.ocl, device_param->kernel2);
+ if (device_param->kernel23) hc_clReleaseKernel (data.ocl, device_param->kernel23);
+ if (device_param->kernel3) hc_clReleaseKernel (data.ocl, device_param->kernel3);
+ if (device_param->kernel_mp) hc_clReleaseKernel (data.ocl, device_param->kernel_mp);
+ if (device_param->kernel_mp_l) hc_clReleaseKernel (data.ocl, device_param->kernel_mp_l);
+ if (device_param->kernel_mp_r) hc_clReleaseKernel (data.ocl, device_param->kernel_mp_r);
+ if (device_param->kernel_tb) hc_clReleaseKernel (data.ocl, device_param->kernel_tb);
+ 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->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);
}
// reset default fan speed
if (quiet == 0) log_info_nn ("Started: %s", ctime (&proc_start));
if (quiet == 0) log_info_nn ("Stopped: %s", ctime (&proc_stop));
+ if (data.ocl) ocl_close (data.ocl);
+
if (data.devices_status == STATUS_ABORTED) return 2;
if (data.devices_status == STATUS_QUIT) return 2;
if (data.devices_status == STATUS_STOP_AT_CHECKPOINT) return 2;
projects / hashcat.git / commitdiff