diff --git a/include/opencl.h b/include/opencl.h index 9091aecf5..cd5728dfd 100644 --- a/include/opencl.h +++ b/include/opencl.h @@ -63,7 +63,10 @@ int run_cracker (opencl_ctx_t *opencl_ctx, hc_device_param_t *device_param, hash int opencl_ctx_init (opencl_ctx_t *opencl_ctx, const char *opencl_platforms, const char *opencl_devices, const char *opencl_device_types, const uint opencl_vector_width, const uint opencl_vector_width_chgd, const uint nvidia_spin_damp, const uint nvidia_spin_damp_chgd, const uint workload_profile, const uint kernel_accel, const uint kernel_accel_chgd, const uint kernel_loops, const uint kernel_loops_chgd, const uint keyspace, const uint stdout_flag); void opencl_ctx_destroy (opencl_ctx_t *opencl_ctx); -int opencl_ctx_devices_init (opencl_ctx_t *opencl_ctx, const hashconfig_t *hashconfig, const tuning_db_t *tuning_db, const uint attack_mode, const uint quiet, const uint force, const uint benchmark, const uint machine_readable, const uint algorithm_pos); +int opencl_ctx_devices_init (opencl_ctx_t *opencl_ctx, const hashconfig_t *hashconfig, const tuning_db_t *tuning_db, const uint attack_mode, const bool quiet, const bool force, const bool benchmark, const bool machine_readable, const uint algorithm_pos); void opencl_ctx_devices_destroy (opencl_ctx_t *opencl_ctx); +int opencl_session_begin (opencl_ctx_t *opencl_ctx, const hashconfig_t *hashconfig, const hashes_t *hashes, const session_ctx_t *session_ctx); +int opencl_session_destroy (opencl_ctx_t *opencl_ctx); + #endif // _OPENCL_H diff --git a/include/session.h b/include/session.h new file mode 100644 index 000000000..56be15af8 --- /dev/null +++ b/include/session.h @@ -0,0 +1,13 @@ +/** + * Author......: See docs/credits.txt + * License.....: MIT + */ + +#ifndef _SESSION_H +#define _SESSION_H + +void session_ctx_init (session_ctx_t *session_ctx, const bool quiet, const bool force, const bool benchmark, const u32 scrypt_tmto, char *cwd, char *install_dir, char *profile_dir, char *session_dir, char *shared_dir, char *cpath_real, const u32 wordlist_mode, char *rule_buf_l, char *rule_buf_r, const int rule_len_l, const int rule_len_r, const u32 kernel_rules_cnt, kernel_rule_t *kernel_rules_buf, const u32 attack_mode, const u32 attack_kern, const u32 bitmap_size, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, u32 *bitmap_s1_a, u32 *bitmap_s1_b, u32 *bitmap_s1_c, u32 *bitmap_s1_d, u32 *bitmap_s2_a, u32 *bitmap_s2_b, u32 *bitmap_s2_c, u32 *bitmap_s2_d); + +void session_ctx_destroy (session_ctx_t *session_ctx); + +#endif // _SESSION_H diff --git a/include/types.h b/include/types.h index c1c52c987..d420103a9 100644 --- a/include/types.h +++ b/include/types.h @@ -482,10 +482,10 @@ typedef struct void *ocl; - cl_uint platforms_cnt; + cl_uint platforms_cnt; cl_platform_id *platforms; - cl_uint platform_devices_cnt; + cl_uint platform_devices_cnt; cl_device_id *platform_devices; u32 devices_cnt; @@ -518,6 +518,8 @@ typedef struct int need_nvapi; int need_xnvctrl; + int force_jit_compilation; + } opencl_ctx_t; #if defined (__APPLE__) @@ -731,6 +733,50 @@ typedef struct } wl_data_t; +typedef struct +{ + bool quiet; + bool force; + bool benchmark; + + u32 scrypt_tmto; + + char *cwd; + char *install_dir; + char *profile_dir; + char *session_dir; + char *shared_dir; + char *cpath_real; + + u32 wordlist_mode; + + char *rule_buf_l; + char *rule_buf_r; + int rule_len_l; + int rule_len_r; + + u32 kernel_rules_cnt; + kernel_rule_t *kernel_rules_buf; + + u32 attack_mode; + u32 attack_kern; + + u32 bitmap_size; + u32 bitmap_mask; + u32 bitmap_shift1; + u32 bitmap_shift2; + + u32 *bitmap_s1_a; + u32 *bitmap_s1_b; + u32 *bitmap_s1_c; + u32 *bitmap_s1_d; + u32 *bitmap_s2_a; + u32 *bitmap_s2_b; + u32 *bitmap_s2_c; + u32 *bitmap_s2_d; + +} session_ctx_t; + typedef struct { /** @@ -794,13 +840,6 @@ typedef struct hm_attrs_t hm_device[DEVICES_MAX]; #endif - /** - * hashes - */ - - u32 scrypt_tmp_size; - u32 scrypt_tmto_final; - /** * logging */ @@ -830,7 +869,6 @@ typedef struct u32 maskpos; char *session; - char *homedir; char *install_dir; char *profile_dir; char *session_dir; @@ -880,6 +918,7 @@ typedef struct potfile_ctx_t *potfile_ctx; loopback_ctx_t *loopback_ctx; debugfile_ctx_t *debugfile_ctx; + session_ctx_t *session_ctx; #if defined (HAVE_HWMON) u32 gpu_temp_disable; diff --git a/src/Makefile b/src/Makefile index 25131a4d3..3ccf99151 100644 --- a/src/Makefile +++ b/src/Makefile @@ -163,7 +163,7 @@ LFLAGS_CROSS_WIN := -lpsapi ## Objects ## -OBJS_ALL := affinity attack_mode autotune benchmark bitmap bitops common convert cpu_aes cpu_crc32 cpu_des cpu_md5 cpu_sha1 cpu_sha256 data debugfile dictstat dispatch dynloader ext_ADL ext_nvapi ext_nvml ext_OpenCL ext_xnvctrl filehandling filenames folder hash_management hlfmt hwmon induct interface locking logfile logging loopback memory monitor mpsp opencl outfile_check outfile potfile powertune remove restore rp_cpu rp_kernel_on_cpu runtime shared status stdout terminal thread timer tuningdb usage version weak_hash wordlist +OBJS_ALL := affinity attack_mode autotune benchmark bitmap bitops common convert cpu_aes cpu_crc32 cpu_des cpu_md5 cpu_sha1 cpu_sha256 data debugfile dictstat dispatch dynloader ext_ADL ext_nvapi ext_nvml ext_OpenCL ext_xnvctrl filehandling filenames folder hash_management hlfmt hwmon induct interface locking logfile logging loopback memory monitor mpsp opencl outfile_check outfile potfile powertune remove restore rp_cpu rp_kernel_on_cpu runtime session shared status stdout terminal thread timer tuningdb usage version weak_hash wordlist NATIVE_OBJS := $(foreach OBJ,$(OBJS_ALL),obj/$(OBJ).NATIVE.o) diff --git a/src/hashcat.c b/src/hashcat.c index 232364790..ff08d3d8d 100644 --- a/src/hashcat.c +++ b/src/hashcat.c @@ -84,6 +84,7 @@ #include "induct.h" #include "dispatch.h" #include "monitor.h" +#include "session.h" extern hc_global_data_t data; @@ -2069,19 +2070,19 @@ int main (int argc, char **argv) * Some algorithm, like descrypt, can benefit from JIT compilation */ - int force_jit_compilation = -1; + opencl_ctx->force_jit_compilation = -1; if (hashconfig->hash_mode == 8900) { - force_jit_compilation = 8900; + opencl_ctx->force_jit_compilation = 8900; } else if (hashconfig->hash_mode == 9300) { - force_jit_compilation = 8900; + opencl_ctx->force_jit_compilation = 8900; } else if (hashconfig->hash_mode == 1500 && attack_mode == ATTACK_MODE_BF && hashes->salts_cnt == 1) { - force_jit_compilation = 1500; + opencl_ctx->force_jit_compilation = 1500; } /** @@ -2901,2018 +2902,84 @@ int main (int argc, char **argv) if (data.quiet == 0) log_info_nn ("Initializing device kernels and memory..."); + session_ctx_t *session_ctx = (session_ctx_t *) mymalloc (sizeof (session_ctx_t)); + + data.session_ctx = session_ctx; + + session_ctx_init (session_ctx, quiet, force, benchmark, scrypt_tmto, cwd, install_dir, profile_dir, session_dir, shared_dir, cpath_real, wordlist_mode, rule_buf_l, rule_buf_r, rule_len_l, rule_len_r, kernel_rules_cnt, kernel_rules_buf, attack_mode, attack_kern, bitmap_size, bitmap_mask, bitmap_shift1, bitmap_shift2, bitmap_s1_a, bitmap_s1_b, bitmap_s1_c, bitmap_s1_d, bitmap_s2_a, bitmap_s2_b, bitmap_s2_c, bitmap_s2_d); + + opencl_session_begin (opencl_ctx, hashconfig, hashes, session_ctx); + + if (data.quiet == 0) log_info_nn (""); + + /** + * Store initial fanspeed if gpu_temp_retain is enabled + */ + + #if defined(HAVE_HWMON) + for (uint device_id = 0; device_id < opencl_ctx->devices_cnt; device_id++) { - cl_int CL_err = CL_SUCCESS; - - /** - * host buffer - */ - hc_device_param_t *device_param = &opencl_ctx->devices_param[device_id]; if (device_param->skipped) continue; - /** - * device properties - */ + if (gpu_temp_disable == 1) continue; - const char *device_name_chksum = device_param->device_name_chksum; - const u32 device_processors = device_param->device_processors; + if (gpu_temp_retain == 0) continue; - /** - * create context for each device - */ + hc_thread_mutex_lock (mux_hwmon); - 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 (opencl_ctx->ocl, properties, 1, &device_param->device, NULL, NULL, &device_param->context); - - if (CL_err != CL_SUCCESS) + if (data.hm_device[device_id].fan_get_supported == 1) { - log_error ("ERROR: clCreateContext(): %s\n", val2cstr_cl (CL_err)); + const int fanspeed = hm_get_fanspeed_with_device_id (opencl_ctx, device_id); + const int fanpolicy = hm_get_fanpolicy_with_device_id (opencl_ctx, device_id); - return -1; - } + // we also set it to tell the OS we take control over the fan and it's automatic controller + // if it was set to automatic. we do not control user-defined fanspeeds. - /** - * create command-queue - */ - - // not supported with NV - // device_param->command_queue = hc_clCreateCommandQueueWithProperties (device_param->context, device_param->device, NULL); - - CL_err = hc_clCreateCommandQueue (opencl_ctx->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 - * because of shared memory usage or bitslice - * there needs to be some upper limit, otherwise there's too much overhead - */ - - uint kernel_threads = MIN (KERNEL_THREADS_MAX, device_param->device_maxworkgroup_size); - - if (hashconfig->hash_mode == 8900) kernel_threads = 64; // Scrypt - if (hashconfig->hash_mode == 9300) kernel_threads = 64; // Scrypt - - if (device_param->device_type & CL_DEVICE_TYPE_CPU) - { - kernel_threads = KERNEL_THREADS_MAX_CPU; - } - - if (hashconfig->hash_mode == 1500) kernel_threads = 64; // DES - if (hashconfig->hash_mode == 3000) kernel_threads = 64; // DES - if (hashconfig->hash_mode == 3100) kernel_threads = 64; // DES - if (hashconfig->hash_mode == 3200) kernel_threads = 8; // Blowfish - if (hashconfig->hash_mode == 7500) kernel_threads = 64; // RC4 - if (hashconfig->hash_mode == 8500) kernel_threads = 64; // DES - if (hashconfig->hash_mode == 9000) kernel_threads = 8; // Blowfish - if (hashconfig->hash_mode == 9700) kernel_threads = 64; // RC4 - if (hashconfig->hash_mode == 9710) kernel_threads = 64; // RC4 - if (hashconfig->hash_mode == 9800) kernel_threads = 64; // RC4 - if (hashconfig->hash_mode == 9810) kernel_threads = 64; // RC4 - if (hashconfig->hash_mode == 10400) kernel_threads = 64; // RC4 - if (hashconfig->hash_mode == 10410) kernel_threads = 64; // RC4 - if (hashconfig->hash_mode == 10500) kernel_threads = 64; // RC4 - if (hashconfig->hash_mode == 13100) kernel_threads = 64; // RC4 - if (hashconfig->hash_mode == 14000) kernel_threads = 64; // DES - if (hashconfig->hash_mode == 14100) kernel_threads = 64; // DES - - device_param->kernel_threads = kernel_threads; - - device_param->hardware_power = device_processors * kernel_threads; - - /** - * create input buffers on device : calculate size of fixed memory buffers - */ - - size_t size_root_css = SP_PW_MAX * sizeof (cs_t); - size_t size_markov_css = SP_PW_MAX * CHARSIZ * sizeof (cs_t); - - device_param->size_root_css = size_root_css; - device_param->size_markov_css = size_markov_css; - - size_t size_results = sizeof (uint); - - device_param->size_results = size_results; - - size_t size_rules = kernel_rules_cnt * sizeof (kernel_rule_t); - size_t size_rules_c = KERNEL_RULES * sizeof (kernel_rule_t); - - size_t size_plains = hashes->digests_cnt * sizeof (plain_t); - size_t size_salts = hashes->salts_cnt * sizeof (salt_t); - size_t size_esalts = hashes->salts_cnt * hashconfig->esalt_size; - size_t size_shown = hashes->digests_cnt * sizeof (uint); - size_t size_digests = hashes->digests_cnt * hashconfig->dgst_size; - - device_param->size_plains = size_plains; - device_param->size_digests = size_digests; - device_param->size_shown = size_shown; - device_param->size_salts = size_salts; - - size_t size_combs = KERNEL_COMBS * sizeof (comb_t); - size_t size_bfs = KERNEL_BFS * sizeof (bf_t); - size_t size_tm = 32 * sizeof (bs_word_t); - - // scryptV stuff - - size_t size_scrypt = 4; - - if ((hashconfig->hash_mode == 8900) || (hashconfig->hash_mode == 9300)) - { - // we need to check that all hashes have the same scrypt settings - - const u32 scrypt_N = hashes->salts_buf[0].scrypt_N; - const u32 scrypt_r = hashes->salts_buf[0].scrypt_r; - const u32 scrypt_p = hashes->salts_buf[0].scrypt_p; - - for (uint i = 1; i < hashes->salts_cnt; i++) + if (fanpolicy == 1) { - if ((hashes->salts_buf[i].scrypt_N != scrypt_N) - || (hashes->salts_buf[i].scrypt_r != scrypt_r) - || (hashes->salts_buf[i].scrypt_p != scrypt_p)) - { - log_error ("ERROR: Mixed scrypt settings not supported"); + data.hm_device[device_id].fan_set_supported = 1; - return -1; + int rc = -1; + + if (device_param->device_vendor_id == VENDOR_ID_AMD) + { + rc = hm_set_fanspeed_with_device_id_adl (device_id, fanspeed, 1); } - } - - uint tmto_start = 0; - uint tmto_stop = 10; - - if (scrypt_tmto) - { - tmto_start = scrypt_tmto; - } - else - { - // in case the user did not specify the tmto manually - // use some values known to run best (tested on 290x for AMD and GTX1080 for NV) - - if (hashconfig->hash_mode == 8900) + else if (device_param->device_vendor_id == VENDOR_ID_NV) { - if (device_param->device_vendor_id == VENDOR_ID_AMD) - { - tmto_start = 3; - } - else if (device_param->device_vendor_id == VENDOR_ID_NV) - { - tmto_start = 2; - } - } - else if (hashconfig->hash_mode == 9300) - { - if (device_param->device_vendor_id == VENDOR_ID_AMD) - { - tmto_start = 2; - } - else if (device_param->device_vendor_id == VENDOR_ID_NV) - { - tmto_start = 4; - } - } - } - - data.scrypt_tmp_size = (128 * scrypt_r * scrypt_p); - - device_param->kernel_accel_min = 1; - device_param->kernel_accel_max = 8; - - uint tmto; - - for (tmto = tmto_start; tmto < tmto_stop; tmto++) - { - size_scrypt = (128 * scrypt_r) * scrypt_N; - - size_scrypt /= 1u << tmto; - - size_scrypt *= device_param->device_processors * device_param->kernel_threads * device_param->kernel_accel_max; - - if ((size_scrypt / 4) > device_param->device_maxmem_alloc) - { - if (quiet == 0) log_info ("WARNING: Not enough single-block device memory allocatable to use --scrypt-tmto %d, increasing...", tmto); - - continue; - } - - if (size_scrypt > device_param->device_global_mem) - { - if (quiet == 0) log_info ("WARNING: Not enough total device memory allocatable to use --scrypt-tmto %d, increasing...", tmto); - - continue; - } - - for (uint salts_pos = 0; salts_pos < hashes->salts_cnt; salts_pos++) - { - data.scrypt_tmto_final = tmto; - } - - break; - } - - if (tmto == tmto_stop) - { - log_error ("ERROR: Can't allocate enough device memory"); - - return -1; - } - - if (quiet == 0) log_info ("SCRYPT tmto optimizer value set to: %u, mem: %" PRIu64 "\n", data.scrypt_tmto_final, size_scrypt); - } - - size_t size_scrypt4 = size_scrypt / 4; - - /** - * some algorithms need a fixed kernel-loops count - */ - - if (hashconfig->hash_mode == 1500 && attack_mode == ATTACK_MODE_BF) - { - const u32 kernel_loops_fixed = 1024; - - device_param->kernel_loops_min = kernel_loops_fixed; - device_param->kernel_loops_max = kernel_loops_fixed; - } - - if (hashconfig->hash_mode == 3000 && attack_mode == ATTACK_MODE_BF) - { - const u32 kernel_loops_fixed = 1024; - - device_param->kernel_loops_min = kernel_loops_fixed; - device_param->kernel_loops_max = kernel_loops_fixed; - } - - if (hashconfig->hash_mode == 8900) - { - const u32 kernel_loops_fixed = 1; - - device_param->kernel_loops_min = kernel_loops_fixed; - device_param->kernel_loops_max = kernel_loops_fixed; - } - - if (hashconfig->hash_mode == 9300) - { - const u32 kernel_loops_fixed = 1; - - device_param->kernel_loops_min = kernel_loops_fixed; - device_param->kernel_loops_max = kernel_loops_fixed; - } - - if (hashconfig->hash_mode == 12500) - { - const u32 kernel_loops_fixed = ROUNDS_RAR3 / 16; - - device_param->kernel_loops_min = kernel_loops_fixed; - device_param->kernel_loops_max = kernel_loops_fixed; - } - - if (hashconfig->hash_mode == 14000 && attack_mode == ATTACK_MODE_BF) - { - const u32 kernel_loops_fixed = 1024; - - device_param->kernel_loops_min = kernel_loops_fixed; - device_param->kernel_loops_max = kernel_loops_fixed; - } - - if (hashconfig->hash_mode == 14100 && attack_mode == ATTACK_MODE_BF) - { - const u32 kernel_loops_fixed = 1024; - - device_param->kernel_loops_min = kernel_loops_fixed; - device_param->kernel_loops_max = kernel_loops_fixed; - } - - /** - * some algorithms have a maximum kernel-loops count - */ - - if (device_param->kernel_loops_min < device_param->kernel_loops_max) - { - u32 innerloop_cnt = 0; - - if (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL) - { - if (data.attack_kern == ATTACK_KERN_STRAIGHT) innerloop_cnt = data.kernel_rules_cnt; - else if (data.attack_kern == ATTACK_KERN_COMBI) innerloop_cnt = data.combs_cnt; - else if (data.attack_kern == ATTACK_KERN_BF) innerloop_cnt = data.bfs_cnt; - } - else - { - innerloop_cnt = hashes->salts_buf[0].salt_iter; - } - - if ((innerloop_cnt >= device_param->kernel_loops_min) && - (innerloop_cnt <= device_param->kernel_loops_max)) - { - device_param->kernel_loops_max = innerloop_cnt; - } - } - - u32 kernel_accel_min = device_param->kernel_accel_min; - u32 kernel_accel_max = device_param->kernel_accel_max; - - // find out if we would request too much memory on memory blocks which are based on kernel_accel - - size_t size_pws = 4; - size_t size_tmps = 4; - size_t size_hooks = 4; - - while (kernel_accel_max >= kernel_accel_min) - { - const u32 kernel_power_max = device_processors * kernel_threads * kernel_accel_max; - - // size_pws - - size_pws = kernel_power_max * sizeof (pw_t); - - // size_tmps - - switch (hashconfig->hash_mode) - { - case 400: size_tmps = kernel_power_max * sizeof (phpass_tmp_t); break; - case 500: size_tmps = kernel_power_max * sizeof (md5crypt_tmp_t); break; - case 501: size_tmps = kernel_power_max * sizeof (md5crypt_tmp_t); break; - case 1600: size_tmps = kernel_power_max * sizeof (md5crypt_tmp_t); break; - case 1800: size_tmps = kernel_power_max * sizeof (sha512crypt_tmp_t); break; - case 2100: size_tmps = kernel_power_max * sizeof (dcc2_tmp_t); break; - case 2500: size_tmps = kernel_power_max * sizeof (wpa_tmp_t); break; - case 3200: size_tmps = kernel_power_max * sizeof (bcrypt_tmp_t); break; - case 5200: size_tmps = kernel_power_max * sizeof (pwsafe3_tmp_t); break; - case 5800: size_tmps = kernel_power_max * sizeof (androidpin_tmp_t); break; - case 6211: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; - case 6212: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; - case 6213: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; - case 6221: size_tmps = kernel_power_max * sizeof (tc64_tmp_t); break; - case 6222: size_tmps = kernel_power_max * sizeof (tc64_tmp_t); break; - case 6223: size_tmps = kernel_power_max * sizeof (tc64_tmp_t); break; - case 6231: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; - case 6232: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; - case 6233: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; - case 6241: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; - case 6242: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; - case 6243: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; - case 6300: size_tmps = kernel_power_max * sizeof (md5crypt_tmp_t); break; - case 6400: size_tmps = kernel_power_max * sizeof (sha256aix_tmp_t); break; - case 6500: size_tmps = kernel_power_max * sizeof (sha512aix_tmp_t); break; - case 6600: size_tmps = kernel_power_max * sizeof (agilekey_tmp_t); break; - case 6700: size_tmps = kernel_power_max * sizeof (sha1aix_tmp_t); break; - case 6800: size_tmps = kernel_power_max * sizeof (lastpass_tmp_t); break; - case 7100: size_tmps = kernel_power_max * sizeof (pbkdf2_sha512_tmp_t); break; - case 7200: size_tmps = kernel_power_max * sizeof (pbkdf2_sha512_tmp_t); break; - case 7400: size_tmps = kernel_power_max * sizeof (sha256crypt_tmp_t); break; - case 7900: size_tmps = kernel_power_max * sizeof (drupal7_tmp_t); break; - case 8200: size_tmps = kernel_power_max * sizeof (pbkdf2_sha512_tmp_t); break; - case 8800: size_tmps = kernel_power_max * sizeof (androidfde_tmp_t); break; - case 8900: size_tmps = kernel_power_max * data.scrypt_tmp_size; break; - case 9000: size_tmps = kernel_power_max * sizeof (pwsafe2_tmp_t); break; - case 9100: size_tmps = kernel_power_max * sizeof (lotus8_tmp_t); break; - case 9200: size_tmps = kernel_power_max * sizeof (pbkdf2_sha256_tmp_t); break; - case 9300: size_tmps = kernel_power_max * data.scrypt_tmp_size; break; - case 9400: size_tmps = kernel_power_max * sizeof (office2007_tmp_t); break; - case 9500: size_tmps = kernel_power_max * sizeof (office2010_tmp_t); break; - case 9600: size_tmps = kernel_power_max * sizeof (office2013_tmp_t); break; - case 10000: size_tmps = kernel_power_max * sizeof (pbkdf2_sha256_tmp_t); break; - case 10200: size_tmps = kernel_power_max * sizeof (cram_md5_t); break; - case 10300: size_tmps = kernel_power_max * sizeof (saph_sha1_tmp_t); break; - case 10500: size_tmps = kernel_power_max * sizeof (pdf14_tmp_t); break; - case 10700: size_tmps = kernel_power_max * sizeof (pdf17l8_tmp_t); break; - case 10900: size_tmps = kernel_power_max * sizeof (pbkdf2_sha256_tmp_t); break; - case 11300: size_tmps = kernel_power_max * sizeof (bitcoin_wallet_tmp_t); break; - case 11600: size_tmps = kernel_power_max * sizeof (seven_zip_tmp_t); break; - case 11900: size_tmps = kernel_power_max * sizeof (pbkdf2_md5_tmp_t); break; - case 12000: size_tmps = kernel_power_max * sizeof (pbkdf2_sha1_tmp_t); break; - case 12100: size_tmps = kernel_power_max * sizeof (pbkdf2_sha512_tmp_t); break; - case 12200: size_tmps = kernel_power_max * sizeof (ecryptfs_tmp_t); break; - case 12300: size_tmps = kernel_power_max * sizeof (oraclet_tmp_t); break; - case 12400: size_tmps = kernel_power_max * sizeof (bsdicrypt_tmp_t); break; - case 12500: size_tmps = kernel_power_max * sizeof (rar3_tmp_t); break; - case 12700: size_tmps = kernel_power_max * sizeof (mywallet_tmp_t); break; - case 12800: size_tmps = kernel_power_max * sizeof (pbkdf2_sha256_tmp_t); break; - case 12900: size_tmps = kernel_power_max * sizeof (pbkdf2_sha256_tmp_t); break; - case 13000: size_tmps = kernel_power_max * sizeof (pbkdf2_sha256_tmp_t); break; - case 13200: size_tmps = kernel_power_max * sizeof (axcrypt_tmp_t); break; - case 13400: size_tmps = kernel_power_max * sizeof (keepass_tmp_t); break; - case 13600: size_tmps = kernel_power_max * sizeof (pbkdf2_sha1_tmp_t); break; - case 13711: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; - case 13712: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; - case 13713: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; - case 13721: size_tmps = kernel_power_max * sizeof (tc64_tmp_t); break; - case 13722: size_tmps = kernel_power_max * sizeof (tc64_tmp_t); break; - case 13723: size_tmps = kernel_power_max * sizeof (tc64_tmp_t); break; - case 13731: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; - case 13732: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; - case 13733: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; - case 13741: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; - case 13742: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; - case 13743: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; - case 13751: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; - case 13752: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; - case 13753: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; - case 13761: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; - case 13762: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; - case 13763: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; - }; - - // size_hooks - - if ((hashconfig->opts_type & OPTS_TYPE_HOOK12) || (hashconfig->opts_type & OPTS_TYPE_HOOK23)) - { - switch (hashconfig->hash_mode) - { - } - } - - // now check if all device-memory sizes which depend on the kernel_accel_max amplifier are within its boundaries - // if not, decrease amplifier and try again - - int memory_limit_hit = 0; - - if (size_pws > device_param->device_maxmem_alloc) memory_limit_hit = 1; - if (size_tmps > device_param->device_maxmem_alloc) memory_limit_hit = 1; - if (size_hooks > device_param->device_maxmem_alloc) memory_limit_hit = 1; - - const u64 size_total - = bitmap_size - + bitmap_size - + bitmap_size - + bitmap_size - + bitmap_size - + bitmap_size - + bitmap_size - + bitmap_size - + size_bfs - + size_combs - + size_digests - + size_esalts - + size_hooks - + size_markov_css - + size_plains - + size_pws - + size_pws // not a bug - + size_results - + size_root_css - + size_rules - + size_rules_c - + size_salts - + size_scrypt4 - + size_scrypt4 - + size_scrypt4 - + size_scrypt4 - + size_shown - + size_tm - + size_tmps; - - if (size_total > device_param->device_global_mem) memory_limit_hit = 1; - - if (memory_limit_hit == 1) - { - kernel_accel_max--; - - continue; - } - - break; - } - - if (kernel_accel_max < kernel_accel_min) - { - log_error ("- Device #%u: Device does not provide enough allocatable device-memory to handle this attack", device_id + 1); - - return -1; - } - - device_param->kernel_accel_min = kernel_accel_min; - device_param->kernel_accel_max = kernel_accel_max; - - /* - if (kernel_accel_max < kernel_accel) - { - if (quiet == 0) log_info ("- Device #%u: Reduced maximum kernel-accel to %u", device_id + 1, kernel_accel_max); - - device_param->kernel_accel = kernel_accel_max; - } - */ - - device_param->size_bfs = size_bfs; - device_param->size_combs = size_combs; - device_param->size_rules = size_rules; - device_param->size_rules_c = size_rules_c; - device_param->size_pws = size_pws; - device_param->size_tmps = size_tmps; - device_param->size_hooks = size_hooks; - - /** - * default building options - */ - - if (chdir (cpath_real) == -1) - { - log_error ("ERROR: %s: %s", cpath_real, strerror (errno)); - - return -1; - } - - char build_opts[1024] = { 0 }; - - #if defined (_WIN) - snprintf (build_opts, sizeof (build_opts) - 1, "-I \"%s\"", cpath_real); - #else - snprintf (build_opts, sizeof (build_opts) - 1, "-I %s", cpath_real); - #endif - - // include check - // this test needs to be done manually because of osx opencl runtime - // if there's a problem with permission, its not reporting back and erroring out silently - - #define files_cnt 15 - - const char *files_names[files_cnt] = - { - "inc_cipher_aes256.cl", - "inc_cipher_serpent256.cl", - "inc_cipher_twofish256.cl", - "inc_common.cl", - "inc_comp_multi_bs.cl", - "inc_comp_multi.cl", - "inc_comp_single_bs.cl", - "inc_comp_single.cl", - "inc_hash_constants.h", - "inc_hash_functions.cl", - "inc_rp.cl", - "inc_rp.h", - "inc_simd.cl", - "inc_types.cl", - "inc_vendor.cl", - }; - - for (int i = 0; i < files_cnt; i++) - { - FILE *fd = fopen (files_names[i], "r"); - - if (fd == NULL) - { - log_error ("ERROR: %s: fopen(): %s", files_names[i], strerror (errno)); - - return -1; - } - - char buf[1]; - - size_t n = fread (buf, 1, 1, fd); - - if (n != 1) - { - log_error ("ERROR: %s: fread(): %s", files_names[i], strerror (errno)); - - return -1; - } - - fclose (fd); - } - - // we don't have sm_* on vendors not NV but it doesn't matter - - char build_opts_new[1024] = { 0 }; - - #if defined (DEBUG) - snprintf (build_opts_new, sizeof (build_opts_new) - 1, "%s -D VENDOR_ID=%u -D CUDA_ARCH=%d -D VECT_SIZE=%u -D DEVICE_TYPE=%u -D DGST_R0=%u -D DGST_R1=%u -D DGST_R2=%u -D DGST_R3=%u -D DGST_ELEM=%u -D KERN_TYPE=%u -D _unroll -cl-std=CL1.1", build_opts, device_param->device_vendor_id, (device_param->sm_major * 100) + device_param->sm_minor, device_param->vector_width, (u32) device_param->device_type, hashconfig->dgst_pos0, hashconfig->dgst_pos1, hashconfig->dgst_pos2, hashconfig->dgst_pos3, hashconfig->dgst_size / 4, hashconfig->kern_type); - #else - snprintf (build_opts_new, sizeof (build_opts_new) - 1, "%s -D VENDOR_ID=%u -D CUDA_ARCH=%d -D VECT_SIZE=%u -D DEVICE_TYPE=%u -D DGST_R0=%u -D DGST_R1=%u -D DGST_R2=%u -D DGST_R3=%u -D DGST_ELEM=%u -D KERN_TYPE=%u -D _unroll -cl-std=CL1.1 -w", build_opts, device_param->device_vendor_id, (device_param->sm_major * 100) + device_param->sm_minor, device_param->vector_width, (u32) device_param->device_type, hashconfig->dgst_pos0, hashconfig->dgst_pos1, hashconfig->dgst_pos2, hashconfig->dgst_pos3, hashconfig->dgst_size / 4, hashconfig->kern_type); - #endif - - strncpy (build_opts, build_opts_new, sizeof (build_opts)); - - #if defined (DEBUG) - log_info ("- Device #%u: build_opts '%s'\n", device_id + 1, build_opts); - #endif - - /** - * main kernel - */ - - { - /** - * kernel source filename - */ - - char source_file[256] = { 0 }; - - generate_source_kernel_filename (hashconfig->attack_exec, attack_kern, hashconfig->kern_type, shared_dir, source_file); - - struct stat sst; - - if (stat (source_file, &sst) == -1) - { - log_error ("ERROR: %s: %s", source_file, strerror (errno)); - - return -1; - } - - /** - * kernel cached filename - */ - - char cached_file[256] = { 0 }; - - generate_cached_kernel_filename (hashconfig->attack_exec, attack_kern, hashconfig->kern_type, profile_dir, device_name_chksum, cached_file); - - int cached = 1; - - struct stat cst; - - if ((stat (cached_file, &cst) == -1) || cst.st_size == 0) - { - cached = 0; - } - - /** - * kernel compile or load - */ - - size_t *kernel_lengths = (size_t *) mymalloc (sizeof (size_t)); - - const u8 **kernel_sources = (const u8 **) mymalloc (sizeof (u8 *)); - - if (force_jit_compilation == -1) - { - if (cached == 0) - { - if (quiet == 0) log_info ("- Device #%u: Kernel %s not found in cache! Building may take a while...", device_id + 1, filename_from_filepath (cached_file)); - - load_kernel (source_file, 1, kernel_lengths, kernel_sources); - - CL_err = hc_clCreateProgramWithSource (opencl_ctx->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 (opencl_ctx->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; - } - - size_t build_log_size = 0; - - /* - CL_err = hc_clGetProgramBuildInfo (opencl_ctx->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; - } - */ - - hc_clGetProgramBuildInfo (opencl_ctx->ocl, device_param->program, device_param->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &build_log_size); - - #if defined (DEBUG) - if ((build_log_size != 0) || (CL_err != CL_SUCCESS)) - #else - if (CL_err != CL_SUCCESS) + #if defined (__linux__) + rc = set_fan_control (data.hm_xnvctrl, data.hm_device[device_id].xnvctrl, NV_CTRL_GPU_COOLER_MANUAL_CONTROL_TRUE); #endif - { - char *build_log = (char *) mymalloc (build_log_size + 1); - CL_err = hc_clGetProgramBuildInfo (opencl_ctx->ocl, device_param->program, device_param->device, CL_PROGRAM_BUILD_LOG, build_log_size, build_log, NULL); + #if defined (_WIN) + rc = hm_set_fanspeed_with_device_id_nvapi (device_id, fanspeed, 1); + #endif + } - if (CL_err != CL_SUCCESS) - { - log_error ("ERROR: clGetProgramBuildInfo(): %s\n", val2cstr_cl (CL_err)); - - return -1; - } - - puts (build_log); - - myfree (build_log); - } - - if (CL_err != CL_SUCCESS) - { - device_param->skipped = true; - - log_info ("- Device #%u: Kernel %s build failure. Proceeding without this device.", device_id + 1, source_file); - - continue; - } - - size_t binary_size; - - CL_err = hc_clGetProgramInfo (opencl_ctx->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); - - CL_err = hc_clGetProgramInfo (opencl_ctx->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); - - local_free (binary); + if (rc == 0) + { + data.hm_device[device_id].fan_set_supported = 1; } else { - #if defined (DEBUG) - log_info ("- Device #%u: Kernel %s (%ld bytes)", device_id + 1, cached_file, cst.st_size); - #endif + log_info ("WARNING: Failed to set initial fan speed for device #%u", device_id + 1); - load_kernel (cached_file, 1, kernel_lengths, kernel_sources); - - CL_err = hc_clCreateProgramWithBinary (opencl_ctx->ocl, device_param->context, 1, &device_param->device, kernel_lengths, (const u8 **) kernel_sources, NULL, &device_param->program); - - if (CL_err != CL_SUCCESS) - { - log_error ("ERROR: clCreateProgramWithBinary(): %s\n", val2cstr_cl (CL_err)); - - return -1; - } - - CL_err = hc_clBuildProgram (opencl_ctx->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; - } + data.hm_device[device_id].fan_set_supported = 0; } } else { - #if defined (DEBUG) - log_info ("- Device #%u: Kernel %s (%ld bytes)", device_id + 1, source_file, sst.st_size); - #endif - - load_kernel (source_file, 1, kernel_lengths, kernel_sources); - - CL_err = hc_clCreateProgramWithSource (opencl_ctx->ocl, device_param->context, 1, (const char **) kernel_sources, NULL, &device_param->program); - - if (CL_err != CL_SUCCESS) - { - log_error ("ERROR: clCreateProgramWithSource(): %s\n", val2cstr_cl (CL_err)); - - return -1; - } - - char build_opts_update[1024] = { 0 }; - - if (force_jit_compilation == 1500) - { - snprintf (build_opts_update, sizeof (build_opts_update) - 1, "%s -DDESCRYPT_SALT=%u", build_opts, hashes->salts_buf[0].salt_buf[0]); - } - else if (force_jit_compilation == 8900) - { - snprintf (build_opts_update, sizeof (build_opts_update) - 1, "%s -DSCRYPT_N=%u -DSCRYPT_R=%u -DSCRYPT_P=%u -DSCRYPT_TMTO=%u -DSCRYPT_TMP_ELEM=%u", build_opts, hashes->salts_buf[0].scrypt_N, hashes->salts_buf[0].scrypt_r, hashes->salts_buf[0].scrypt_p, 1 << data.scrypt_tmto_final, data.scrypt_tmp_size / 16); - } - else - { - snprintf (build_opts_update, sizeof (build_opts_update) - 1, "%s", build_opts); - } - - CL_err = hc_clBuildProgram (opencl_ctx->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; - } - - size_t build_log_size = 0; - - /* - CL_err = hc_clGetProgramBuildInfo (opencl_ctx->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; - } - */ - - hc_clGetProgramBuildInfo (opencl_ctx->ocl, device_param->program, device_param->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &build_log_size); - - #if defined (DEBUG) - if ((build_log_size != 0) || (CL_err != CL_SUCCESS)) - #else - if (CL_err != CL_SUCCESS) - #endif - { - char *build_log = (char *) mymalloc (build_log_size + 1); - - CL_err = hc_clGetProgramBuildInfo (opencl_ctx->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); - - myfree (build_log); - } - - if (CL_err != CL_SUCCESS) - { - device_param->skipped = true; - - log_info ("- Device #%u: Kernel %s build failure. Proceeding without this device.", device_id + 1, source_file); - } - } - - local_free (kernel_lengths); - local_free (kernel_sources[0]); - local_free (kernel_sources); - } - - /** - * word generator kernel - */ - - if (attack_mode != ATTACK_MODE_STRAIGHT) - { - /** - * kernel mp source filename - */ - - char source_file[256] = { 0 }; - - generate_source_kernel_mp_filename (hashconfig->opti_type, hashconfig->opts_type, shared_dir, source_file); - - struct stat sst; - - if (stat (source_file, &sst) == -1) - { - log_error ("ERROR: %s: %s", source_file, strerror (errno)); - - return -1; - } - - /** - * kernel mp cached filename - */ - - char cached_file[256] = { 0 }; - - generate_cached_kernel_mp_filename (hashconfig->opti_type, hashconfig->opts_type, profile_dir, device_name_chksum, cached_file); - - int cached = 1; - - struct stat cst; - - if (stat (cached_file, &cst) == -1) - { - cached = 0; - } - - /** - * kernel compile or load - */ - - size_t *kernel_lengths = (size_t *) mymalloc (sizeof (size_t)); - - const u8 **kernel_sources = (const u8 **) mymalloc (sizeof (u8 *)); - - if (cached == 0) - { - if (quiet == 0) log_info ("- Device #%u: Kernel %s not found in cache! Building may take a while...", device_id + 1, filename_from_filepath (cached_file)); - if (quiet == 0) log_info (""); - - load_kernel (source_file, 1, kernel_lengths, kernel_sources); - - CL_err = hc_clCreateProgramWithSource (opencl_ctx->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; - } - - CL_err = hc_clBuildProgram (opencl_ctx->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; - } - - size_t build_log_size = 0; - - /* - CL_err = hc_clGetProgramBuildInfo (opencl_ctx->ocl, device_param->program_mp, 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; - } - */ - - hc_clGetProgramBuildInfo (opencl_ctx->ocl, device_param->program_mp, device_param->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &build_log_size); - - #if defined (DEBUG) - if ((build_log_size != 0) || (CL_err != CL_SUCCESS)) - #else - if (CL_err != CL_SUCCESS) - #endif - { - char *build_log = (char *) mymalloc (build_log_size + 1); - - CL_err = hc_clGetProgramBuildInfo (opencl_ctx->ocl, device_param->program_mp, 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); - - myfree (build_log); - } - - if (CL_err != CL_SUCCESS) - { - device_param->skipped = true; - - log_info ("- Device #%u: Kernel %s build failure. Proceeding without this device.", device_id + 1, source_file); - - continue; - } - - size_t binary_size; - - CL_err = hc_clGetProgramInfo (opencl_ctx->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); - - CL_err = hc_clGetProgramInfo (opencl_ctx->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); - - local_free (binary); - } - else - { - #if defined (DEBUG) - log_info ("- Device #%u: Kernel %s (%ld bytes)", device_id + 1, cached_file, cst.st_size); - #endif - - load_kernel (cached_file, 1, kernel_lengths, kernel_sources); - - CL_err = hc_clCreateProgramWithBinary (opencl_ctx->ocl, device_param->context, 1, &device_param->device, kernel_lengths, (const u8 **) kernel_sources, NULL, &device_param->program_mp); - - if (CL_err != CL_SUCCESS) - { - log_error ("ERROR: clCreateProgramWithBinary(): %s\n", val2cstr_cl (CL_err)); - - return -1; - } - - CL_err = hc_clBuildProgram (opencl_ctx->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); - local_free (kernel_sources[0]); - local_free (kernel_sources); - } - - /** - * amplifier kernel - */ - - if (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL) - { - - } - else - { - /** - * kernel amp source filename - */ - - char source_file[256] = { 0 }; - - generate_source_kernel_amp_filename (attack_kern, shared_dir, source_file); - - struct stat sst; - - if (stat (source_file, &sst) == -1) - { - log_error ("ERROR: %s: %s", source_file, strerror (errno)); - - return -1; - } - - /** - * kernel amp cached filename - */ - - char cached_file[256] = { 0 }; - - generate_cached_kernel_amp_filename (attack_kern, profile_dir, device_name_chksum, cached_file); - - int cached = 1; - - struct stat cst; - - if (stat (cached_file, &cst) == -1) - { - cached = 0; - } - - /** - * kernel compile or load - */ - - size_t *kernel_lengths = (size_t *) mymalloc (sizeof (size_t)); - - const u8 **kernel_sources = (const u8 **) mymalloc (sizeof (u8 *)); - - if (cached == 0) - { - if (quiet == 0) log_info ("- Device #%u: Kernel %s not found in cache! Building may take a while...", device_id + 1, filename_from_filepath (cached_file)); - if (quiet == 0) log_info (""); - - load_kernel (source_file, 1, kernel_lengths, kernel_sources); - - CL_err = hc_clCreateProgramWithSource (opencl_ctx->ocl, device_param->context, 1, (const char **) kernel_sources, NULL, &device_param->program_amp); - - if (CL_err != CL_SUCCESS) - { - log_error ("ERROR: clCreateProgramWithSource(): %s\n", val2cstr_cl (CL_err)); - - return -1; - } - - CL_err = hc_clBuildProgram (opencl_ctx->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; - } - - size_t build_log_size = 0; - - /* - CL_err = hc_clGetProgramBuildInfo (opencl_ctx->ocl, device_param->program_amp, 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; - } - */ - - hc_clGetProgramBuildInfo (opencl_ctx->ocl, device_param->program_amp, device_param->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &build_log_size); - - #if defined (DEBUG) - if ((build_log_size != 0) || (CL_err != CL_SUCCESS)) - #else - if (CL_err != CL_SUCCESS) - #endif - { - char *build_log = (char *) mymalloc (build_log_size + 1); - - CL_err = hc_clGetProgramBuildInfo (opencl_ctx->ocl, device_param->program_amp, 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); - - myfree (build_log); - } - - if (CL_err != CL_SUCCESS) - { - device_param->skipped = true; - - log_info ("- Device #%u: Kernel %s build failure. Proceed without this device.", device_id + 1, source_file); - - continue; - } - - size_t binary_size; - - CL_err = hc_clGetProgramInfo (opencl_ctx->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); - - CL_err = hc_clGetProgramInfo (opencl_ctx->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); - - local_free (binary); - } - else - { - #if defined (DEBUG) - if (quiet == 0) log_info ("- Device #%u: Kernel %s (%ld bytes)", device_id + 1, cached_file, cst.st_size); - #endif - - load_kernel (cached_file, 1, kernel_lengths, kernel_sources); - - CL_err = hc_clCreateProgramWithBinary (opencl_ctx->ocl, device_param->context, 1, &device_param->device, kernel_lengths, (const u8 **) kernel_sources, NULL, &device_param->program_amp); - - if (CL_err != CL_SUCCESS) - { - log_error ("ERROR: clCreateProgramWithBinary(): %s\n", val2cstr_cl (CL_err)); - - return -1; - } - - CL_err = hc_clBuildProgram (opencl_ctx->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[0]); - 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) - { - ((uint *) hashes->digests_buf)[0] = -1u; - ((uint *) hashes->digests_buf)[1] = -1u; - ((uint *) hashes->digests_buf)[2] = -1u; - ((uint *) hashes->digests_buf)[3] = -1u; - } - - /** - * global buffers - */ - - CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_ONLY, size_pws, NULL, &device_param->d_pws_buf); - CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_ONLY, size_pws, NULL, &device_param->d_pws_amp_buf); - CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_WRITE, size_tmps, NULL, &device_param->d_tmps); - CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_WRITE, size_hooks, NULL, &device_param->d_hooks); - CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL, &device_param->d_bitmap_s1_a); - CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL, &device_param->d_bitmap_s1_b); - CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL, &device_param->d_bitmap_s1_c); - CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL, &device_param->d_bitmap_s1_d); - CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL, &device_param->d_bitmap_s2_a); - CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL, &device_param->d_bitmap_s2_b); - CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL, &device_param->d_bitmap_s2_c); - CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL, &device_param->d_bitmap_s2_d); - CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_WRITE, size_plains, NULL, &device_param->d_plain_bufs); - CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_ONLY, size_digests, NULL, &device_param->d_digests_buf); - CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_WRITE, size_shown, NULL, &device_param->d_digests_shown); - CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_ONLY, size_salts, NULL, &device_param->d_salt_bufs); - CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_WRITE, size_results, NULL, &device_param->d_result); - CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_WRITE, size_scrypt4, NULL, &device_param->d_scryptV0_buf); - CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_WRITE, size_scrypt4, NULL, &device_param->d_scryptV1_buf); - CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_WRITE, size_scrypt4, NULL, &device_param->d_scryptV2_buf); - CL_err |= hc_clCreateBuffer (opencl_ctx->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 (opencl_ctx->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 (opencl_ctx->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 (opencl_ctx->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 (opencl_ctx->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 (opencl_ctx->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 (opencl_ctx->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 (opencl_ctx->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 (opencl_ctx->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 (opencl_ctx->ocl, device_param->command_queue, device_param->d_digests_buf, CL_TRUE, 0, size_digests, hashes->digests_buf, 0, NULL, NULL); - CL_err |= hc_clEnqueueWriteBuffer (opencl_ctx->ocl, device_param->command_queue, device_param->d_digests_shown, CL_TRUE, 0, size_shown, hashes->digests_shown, 0, NULL, NULL); - CL_err |= hc_clEnqueueWriteBuffer (opencl_ctx->ocl, device_param->command_queue, device_param->d_salt_bufs, CL_TRUE, 0, size_salts, hashes->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) - { - CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_ONLY, size_rules, NULL, &device_param->d_rules); - CL_err |= hc_clCreateBuffer (opencl_ctx->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; - } - - CL_err = hc_clEnqueueWriteBuffer (opencl_ctx->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) - { - CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_ONLY, size_combs, NULL, &device_param->d_combs); - CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_ONLY, size_combs, NULL, &device_param->d_combs_c); - CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_ONLY, size_root_css, NULL, &device_param->d_root_css_buf); - CL_err |= hc_clCreateBuffer (opencl_ctx->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) - { - CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_ONLY, size_bfs, NULL, &device_param->d_bfs); - CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_ONLY, size_bfs, NULL, &device_param->d_bfs_c); - CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_ONLY, size_tm, NULL, &device_param->d_tm_c); - CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_ONLY, size_root_css, NULL, &device_param->d_root_css_buf); - CL_err |= hc_clCreateBuffer (opencl_ctx->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; + data.hm_device[device_id].fan_set_supported = 0; } } - if (size_esalts) - { - CL_err = hc_clCreateBuffer (opencl_ctx->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 (opencl_ctx->ocl, device_param->command_queue, device_param->d_esalt_bufs, CL_TRUE, 0, size_esalts, hashes->esalts_buf, 0, NULL, NULL); - - if (CL_err != CL_SUCCESS) - { - log_error ("ERROR: clEnqueueWriteBuffer(): %s\n", val2cstr_cl (CL_err)); - - return -1; - } - } - - /** - * main host data - */ - - pw_t *pws_buf = (pw_t *) mymalloc (size_pws); - - device_param->pws_buf = pws_buf; - - comb_t *combs_buf = (comb_t *) mycalloc (KERNEL_COMBS, sizeof (comb_t)); - - device_param->combs_buf = combs_buf; - - void *hooks_buf = mymalloc (size_hooks); - - device_param->hooks_buf = hooks_buf; - - /** - * kernel args - */ - - device_param->kernel_params_buf32[24] = bitmap_mask; - device_param->kernel_params_buf32[25] = bitmap_shift1; - device_param->kernel_params_buf32[26] = bitmap_shift2; - device_param->kernel_params_buf32[27] = 0; // salt_pos - device_param->kernel_params_buf32[28] = 0; // loop_pos - device_param->kernel_params_buf32[29] = 0; // loop_cnt - device_param->kernel_params_buf32[30] = 0; // kernel_rules_cnt - device_param->kernel_params_buf32[31] = 0; // digests_cnt - device_param->kernel_params_buf32[32] = 0; // digests_offset - device_param->kernel_params_buf32[33] = 0; // combs_mode - device_param->kernel_params_buf32[34] = 0; // gid_max - - device_param->kernel_params[ 0] = (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL) - ? &device_param->d_pws_buf - : &device_param->d_pws_amp_buf; - device_param->kernel_params[ 1] = &device_param->d_rules_c; - device_param->kernel_params[ 2] = &device_param->d_combs_c; - device_param->kernel_params[ 3] = &device_param->d_bfs_c; - device_param->kernel_params[ 4] = &device_param->d_tmps; - device_param->kernel_params[ 5] = &device_param->d_hooks; - device_param->kernel_params[ 6] = &device_param->d_bitmap_s1_a; - device_param->kernel_params[ 7] = &device_param->d_bitmap_s1_b; - device_param->kernel_params[ 8] = &device_param->d_bitmap_s1_c; - device_param->kernel_params[ 9] = &device_param->d_bitmap_s1_d; - device_param->kernel_params[10] = &device_param->d_bitmap_s2_a; - device_param->kernel_params[11] = &device_param->d_bitmap_s2_b; - device_param->kernel_params[12] = &device_param->d_bitmap_s2_c; - device_param->kernel_params[13] = &device_param->d_bitmap_s2_d; - device_param->kernel_params[14] = &device_param->d_plain_bufs; - device_param->kernel_params[15] = &device_param->d_digests_buf; - device_param->kernel_params[16] = &device_param->d_digests_shown; - device_param->kernel_params[17] = &device_param->d_salt_bufs; - device_param->kernel_params[18] = &device_param->d_esalt_bufs; - device_param->kernel_params[19] = &device_param->d_result; - device_param->kernel_params[20] = &device_param->d_scryptV0_buf; - device_param->kernel_params[21] = &device_param->d_scryptV1_buf; - device_param->kernel_params[22] = &device_param->d_scryptV2_buf; - device_param->kernel_params[23] = &device_param->d_scryptV3_buf; - device_param->kernel_params[24] = &device_param->kernel_params_buf32[24]; - device_param->kernel_params[25] = &device_param->kernel_params_buf32[25]; - device_param->kernel_params[26] = &device_param->kernel_params_buf32[26]; - device_param->kernel_params[27] = &device_param->kernel_params_buf32[27]; - device_param->kernel_params[28] = &device_param->kernel_params_buf32[28]; - device_param->kernel_params[29] = &device_param->kernel_params_buf32[29]; - device_param->kernel_params[30] = &device_param->kernel_params_buf32[30]; - device_param->kernel_params[31] = &device_param->kernel_params_buf32[31]; - device_param->kernel_params[32] = &device_param->kernel_params_buf32[32]; - device_param->kernel_params[33] = &device_param->kernel_params_buf32[33]; - device_param->kernel_params[34] = &device_param->kernel_params_buf32[34]; - - device_param->kernel_params_mp_buf64[3] = 0; - device_param->kernel_params_mp_buf32[4] = 0; - device_param->kernel_params_mp_buf32[5] = 0; - device_param->kernel_params_mp_buf32[6] = 0; - device_param->kernel_params_mp_buf32[7] = 0; - device_param->kernel_params_mp_buf32[8] = 0; - - device_param->kernel_params_mp[0] = NULL; - device_param->kernel_params_mp[1] = NULL; - device_param->kernel_params_mp[2] = NULL; - device_param->kernel_params_mp[3] = &device_param->kernel_params_mp_buf64[3]; - device_param->kernel_params_mp[4] = &device_param->kernel_params_mp_buf32[4]; - device_param->kernel_params_mp[5] = &device_param->kernel_params_mp_buf32[5]; - device_param->kernel_params_mp[6] = &device_param->kernel_params_mp_buf32[6]; - device_param->kernel_params_mp[7] = &device_param->kernel_params_mp_buf32[7]; - device_param->kernel_params_mp[8] = &device_param->kernel_params_mp_buf32[8]; - - device_param->kernel_params_mp_l_buf64[3] = 0; - device_param->kernel_params_mp_l_buf32[4] = 0; - device_param->kernel_params_mp_l_buf32[5] = 0; - device_param->kernel_params_mp_l_buf32[6] = 0; - device_param->kernel_params_mp_l_buf32[7] = 0; - device_param->kernel_params_mp_l_buf32[8] = 0; - device_param->kernel_params_mp_l_buf32[9] = 0; - - device_param->kernel_params_mp_l[0] = NULL; - device_param->kernel_params_mp_l[1] = NULL; - device_param->kernel_params_mp_l[2] = NULL; - device_param->kernel_params_mp_l[3] = &device_param->kernel_params_mp_l_buf64[3]; - device_param->kernel_params_mp_l[4] = &device_param->kernel_params_mp_l_buf32[4]; - device_param->kernel_params_mp_l[5] = &device_param->kernel_params_mp_l_buf32[5]; - device_param->kernel_params_mp_l[6] = &device_param->kernel_params_mp_l_buf32[6]; - device_param->kernel_params_mp_l[7] = &device_param->kernel_params_mp_l_buf32[7]; - device_param->kernel_params_mp_l[8] = &device_param->kernel_params_mp_l_buf32[8]; - device_param->kernel_params_mp_l[9] = &device_param->kernel_params_mp_l_buf32[9]; - - device_param->kernel_params_mp_r_buf64[3] = 0; - device_param->kernel_params_mp_r_buf32[4] = 0; - device_param->kernel_params_mp_r_buf32[5] = 0; - device_param->kernel_params_mp_r_buf32[6] = 0; - device_param->kernel_params_mp_r_buf32[7] = 0; - device_param->kernel_params_mp_r_buf32[8] = 0; - - device_param->kernel_params_mp_r[0] = NULL; - device_param->kernel_params_mp_r[1] = NULL; - device_param->kernel_params_mp_r[2] = NULL; - device_param->kernel_params_mp_r[3] = &device_param->kernel_params_mp_r_buf64[3]; - device_param->kernel_params_mp_r[4] = &device_param->kernel_params_mp_r_buf32[4]; - device_param->kernel_params_mp_r[5] = &device_param->kernel_params_mp_r_buf32[5]; - device_param->kernel_params_mp_r[6] = &device_param->kernel_params_mp_r_buf32[6]; - device_param->kernel_params_mp_r[7] = &device_param->kernel_params_mp_r_buf32[7]; - device_param->kernel_params_mp_r[8] = &device_param->kernel_params_mp_r_buf32[8]; - - device_param->kernel_params_amp_buf32[5] = 0; // combs_mode - device_param->kernel_params_amp_buf32[6] = 0; // gid_max - - device_param->kernel_params_amp[0] = &device_param->d_pws_buf; - device_param->kernel_params_amp[1] = &device_param->d_pws_amp_buf; - device_param->kernel_params_amp[2] = &device_param->d_rules_c; - device_param->kernel_params_amp[3] = &device_param->d_combs_c; - device_param->kernel_params_amp[4] = &device_param->d_bfs_c; - device_param->kernel_params_amp[5] = &device_param->kernel_params_amp_buf32[5]; - device_param->kernel_params_amp[6] = &device_param->kernel_params_amp_buf32[6]; - - device_param->kernel_params_tm[0] = &device_param->d_bfs_c; - device_param->kernel_params_tm[1] = &device_param->d_tm_c; - - device_param->kernel_params_memset_buf32[1] = 0; // value - device_param->kernel_params_memset_buf32[2] = 0; // gid_max - - device_param->kernel_params_memset[0] = NULL; - device_param->kernel_params_memset[1] = &device_param->kernel_params_memset_buf32[1]; - device_param->kernel_params_memset[2] = &device_param->kernel_params_memset_buf32[2]; - - /** - * kernel name - */ - - size_t kernel_wgs_tmp; - - char kernel_name[64] = { 0 }; - - if (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL) - { - if (hashconfig->opti_type & OPTI_TYPE_SINGLE_HASH) - { - snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_s%02d", hashconfig->kern_type, 4); - - CL_err = hc_clCreateKernel (opencl_ctx->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", hashconfig->kern_type, 8); - - CL_err = hc_clCreateKernel (opencl_ctx->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", hashconfig->kern_type, 16); - - CL_err = hc_clCreateKernel (opencl_ctx->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", hashconfig->kern_type, 4); - - CL_err = hc_clCreateKernel (opencl_ctx->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", hashconfig->kern_type, 8); - - CL_err = hc_clCreateKernel (opencl_ctx->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", hashconfig->kern_type, 16); - - CL_err = hc_clCreateKernel (opencl_ctx->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 (attack_mode == ATTACK_MODE_BF) - { - if (hashconfig->opts_type & OPTS_TYPE_PT_BITSLICE) - { - snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_tm", hashconfig->kern_type); - - CL_err = hc_clCreateKernel (opencl_ctx->ocl, device_param->program, kernel_name, &device_param->kernel_tm); - - if (CL_err != CL_SUCCESS) - { - log_error ("ERROR: clCreateKernel(): %s\n", val2cstr_cl (CL_err)); - - return -1; - } - - CL_err = hc_clGetKernelWorkGroupInfo (opencl_ctx->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; - } - } - } - } - else - { - snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_init", hashconfig->kern_type); - - CL_err = hc_clCreateKernel (opencl_ctx->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", hashconfig->kern_type); - - CL_err = hc_clCreateKernel (opencl_ctx->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", hashconfig->kern_type); - - CL_err = hc_clCreateKernel (opencl_ctx->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 (hashconfig->opts_type & OPTS_TYPE_HOOK12) - { - snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_hook12", hashconfig->kern_type); - - CL_err = hc_clCreateKernel (opencl_ctx->ocl, device_param->program, kernel_name, &device_param->kernel12); - - if (CL_err != CL_SUCCESS) - { - log_error ("ERROR: clCreateKernel(): %s\n", val2cstr_cl (CL_err)); - - return -1; - } - - CL_err = hc_clGetKernelWorkGroupInfo (opencl_ctx->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 (hashconfig->opts_type & OPTS_TYPE_HOOK23) - { - snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_hook23", hashconfig->kern_type); - - CL_err = hc_clCreateKernel (opencl_ctx->ocl, device_param->program, kernel_name, &device_param->kernel23); - - if (CL_err != CL_SUCCESS) - { - log_error ("ERROR: clCreateKernel(): %s\n", val2cstr_cl (CL_err)); - - return -1; - } - - CL_err = hc_clGetKernelWorkGroupInfo (opencl_ctx->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; - } - } - } - - CL_err |= hc_clGetKernelWorkGroupInfo (opencl_ctx->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 (opencl_ctx->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 (opencl_ctx->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++) - { - CL_err |= hc_clSetKernelArg (opencl_ctx->ocl, device_param->kernel1, i, sizeof (cl_mem), device_param->kernel_params[i]); - CL_err |= hc_clSetKernelArg (opencl_ctx->ocl, device_param->kernel2, i, sizeof (cl_mem), device_param->kernel_params[i]); - CL_err |= hc_clSetKernelArg (opencl_ctx->ocl, device_param->kernel3, i, sizeof (cl_mem), device_param->kernel_params[i]); - - if (hashconfig->opts_type & OPTS_TYPE_HOOK12) CL_err |= hc_clSetKernelArg (opencl_ctx->ocl, device_param->kernel12, i, sizeof (cl_mem), device_param->kernel_params[i]); - if (hashconfig->opts_type & OPTS_TYPE_HOOK23) CL_err |= hc_clSetKernelArg (opencl_ctx->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)); - - return -1; - } - } - - for (uint i = 24; i <= 34; i++) - { - CL_err |= hc_clSetKernelArg (opencl_ctx->ocl, device_param->kernel1, i, sizeof (cl_uint), device_param->kernel_params[i]); - CL_err |= hc_clSetKernelArg (opencl_ctx->ocl, device_param->kernel2, i, sizeof (cl_uint), device_param->kernel_params[i]); - CL_err |= hc_clSetKernelArg (opencl_ctx->ocl, device_param->kernel3, i, sizeof (cl_uint), device_param->kernel_params[i]); - - if (hashconfig->opts_type & OPTS_TYPE_HOOK12) CL_err |= hc_clSetKernelArg (opencl_ctx->ocl, device_param->kernel12, i, sizeof (cl_uint), device_param->kernel_params[i]); - if (hashconfig->opts_type & OPTS_TYPE_HOOK23) CL_err |= hc_clSetKernelArg (opencl_ctx->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 - - CL_err = hc_clCreateKernel (opencl_ctx->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 (opencl_ctx->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 (opencl_ctx->ocl, device_param->kernel_memset, 0, sizeof (cl_mem), device_param->kernel_params_memset[0]); - CL_err |= hc_clSetKernelArg (opencl_ctx->ocl, device_param->kernel_memset, 1, sizeof (cl_uint), device_param->kernel_params_memset[1]); - CL_err |= hc_clSetKernelArg (opencl_ctx->ocl, device_param->kernel_memset, 2, sizeof (cl_uint), device_param->kernel_params_memset[2]); - - if (CL_err != CL_SUCCESS) - { - log_error ("ERROR: clSetKernelArg(): %s\n", val2cstr_cl (CL_err)); - - return -1; - } - - // MP start - - if (attack_mode == ATTACK_MODE_BF) - { - CL_err |= hc_clCreateKernel (opencl_ctx->ocl, device_param->program_mp, "l_markov", &device_param->kernel_mp_l); - CL_err |= hc_clCreateKernel (opencl_ctx->ocl, device_param->program_mp, "r_markov", &device_param->kernel_mp_r); - - if (CL_err != CL_SUCCESS) - { - log_error ("ERROR: clCreateKernel(): %s\n", val2cstr_cl (CL_err)); - - return -1; - } - - CL_err |= hc_clGetKernelWorkGroupInfo (opencl_ctx->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 (opencl_ctx->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 (hashconfig->opts_type & OPTS_TYPE_PT_BITSLICE) - { - CL_err |= hc_clSetKernelArg (opencl_ctx->ocl, device_param->kernel_tm, 0, sizeof (cl_mem), device_param->kernel_params_tm[0]); - CL_err |= hc_clSetKernelArg (opencl_ctx->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) - { - CL_err = hc_clCreateKernel (opencl_ctx->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 (opencl_ctx->ocl, device_param->kernel_mp, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp); - - if (CL_err != CL_SUCCESS) - { - log_error ("ERROR: clGetKernelWorkGroupInfo(): %s\n", val2cstr_cl (CL_err)); - - return -1; - } - } - else if (attack_mode == ATTACK_MODE_HYBRID2) - { - CL_err = hc_clCreateKernel (opencl_ctx->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 (opencl_ctx->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 (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL) - { - // nothing to do - } - else - { - CL_err = hc_clCreateKernel (opencl_ctx->ocl, device_param->program_amp, "amp", &device_param->kernel_amp); - - if (CL_err != CL_SUCCESS) - { - log_error ("ERROR: clCreateKernel(): %s\n", val2cstr_cl (CL_err)); - - return -1; - } - - CL_err = hc_clGetKernelWorkGroupInfo (opencl_ctx->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 (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL) - { - // nothing to do - } - else - { - for (uint i = 0; i < 5; i++) - { - CL_err = hc_clSetKernelArg (opencl_ctx->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++) - { - CL_err = hc_clSetKernelArg (opencl_ctx->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; - } - } - } - - // maybe this has been updated by clGetKernelWorkGroupInfo() - // value can only be decreased, so we don't need to reallocate buffers - - device_param->kernel_threads = kernel_threads; - - // zero some data buffers - - run_kernel_bzero (opencl_ctx, device_param, device_param->d_pws_buf, size_pws); - run_kernel_bzero (opencl_ctx, device_param, device_param->d_pws_amp_buf, size_pws); - run_kernel_bzero (opencl_ctx, device_param, device_param->d_tmps, size_tmps); - run_kernel_bzero (opencl_ctx, device_param, device_param->d_hooks, size_hooks); - run_kernel_bzero (opencl_ctx, device_param, device_param->d_plain_bufs, size_plains); - run_kernel_bzero (opencl_ctx, device_param, device_param->d_result, size_results); - - /** - * special buffers - */ - - if (attack_kern == ATTACK_KERN_STRAIGHT) - { - run_kernel_bzero (opencl_ctx, device_param, device_param->d_rules_c, size_rules_c); - } - else if (attack_kern == ATTACK_KERN_COMBI) - { - run_kernel_bzero (opencl_ctx, device_param, device_param->d_combs, size_combs); - run_kernel_bzero (opencl_ctx, device_param, device_param->d_combs_c, size_combs); - run_kernel_bzero (opencl_ctx, device_param, device_param->d_root_css_buf, size_root_css); - run_kernel_bzero (opencl_ctx, device_param, device_param->d_markov_css_buf, size_markov_css); - } - else if (attack_kern == ATTACK_KERN_BF) - { - run_kernel_bzero (opencl_ctx, device_param, device_param->d_bfs, size_bfs); - run_kernel_bzero (opencl_ctx, device_param, device_param->d_bfs_c, size_bfs); - run_kernel_bzero (opencl_ctx, device_param, device_param->d_tm_c, size_tm); - run_kernel_bzero (opencl_ctx, device_param, device_param->d_root_css_buf, size_root_css); - run_kernel_bzero (opencl_ctx, device_param, device_param->d_markov_css_buf, size_markov_css); - } - - #if defined(HAVE_HWMON) - - /** - * Store initial fanspeed if gpu_temp_retain is enabled - */ - - if (gpu_temp_disable == 0) - { - if (gpu_temp_retain != 0) - { - hc_thread_mutex_lock (mux_hwmon); - - if (data.hm_device[device_id].fan_get_supported == 1) - { - const int fanspeed = hm_get_fanspeed_with_device_id (opencl_ctx, device_id); - const int fanpolicy = hm_get_fanpolicy_with_device_id (opencl_ctx, device_id); - - // we also set it to tell the OS we take control over the fan and it's automatic controller - // if it was set to automatic. we do not control user-defined fanspeeds. - - if (fanpolicy == 1) - { - data.hm_device[device_id].fan_set_supported = 1; - - int rc = -1; - - if (device_param->device_vendor_id == VENDOR_ID_AMD) - { - rc = hm_set_fanspeed_with_device_id_adl (device_id, fanspeed, 1); - } - else if (device_param->device_vendor_id == VENDOR_ID_NV) - { - #if defined (__linux__) - rc = set_fan_control (data.hm_xnvctrl, data.hm_device[device_id].xnvctrl, NV_CTRL_GPU_COOLER_MANUAL_CONTROL_TRUE); - #endif - - #if defined (_WIN) - rc = hm_set_fanspeed_with_device_id_nvapi (device_id, fanspeed, 1); - #endif - } - - if (rc == 0) - { - data.hm_device[device_id].fan_set_supported = 1; - } - else - { - log_info ("WARNING: Failed to set initial fan speed for device #%u", device_id + 1); - - data.hm_device[device_id].fan_set_supported = 0; - } - } - else - { - data.hm_device[device_id].fan_set_supported = 0; - } - } - - hc_thread_mutex_unlock (mux_hwmon); - } - } - - #endif // HAVE_HWMON + hc_thread_mutex_unlock (mux_hwmon); } - if (data.quiet == 0) log_info_nn (""); + #endif // HAVE_HWMON /** * In benchmark-mode, inform user which algorithm is checked @@ -6685,6 +4752,39 @@ int main (int argc, char **argv) loopback_write_open (loopback_ctx, induction_directory); } + /** + * some algorithms have a maximum kernel-loops count + */ + + for (uint device_id = 0; device_id < opencl_ctx->devices_cnt; device_id++) + { + hc_device_param_t *device_param = &opencl_ctx->devices_param[device_id]; + + if (device_param->skipped) continue; + + if (device_param->kernel_loops_min < device_param->kernel_loops_max) + { + u32 innerloop_cnt = 0; + + if (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL) + { + if (session_ctx->attack_kern == ATTACK_KERN_STRAIGHT) innerloop_cnt = data.kernel_rules_cnt; + else if (session_ctx->attack_kern == ATTACK_KERN_COMBI) innerloop_cnt = data.combs_cnt; + else if (session_ctx->attack_kern == ATTACK_KERN_BF) innerloop_cnt = data.bfs_cnt; + } + else + { + innerloop_cnt = hashes->salts_buf[0].salt_iter; + } + + if ((innerloop_cnt >= device_param->kernel_loops_min) && + (innerloop_cnt <= device_param->kernel_loops_max)) + { + device_param->kernel_loops_max = innerloop_cnt; + } + } + } + /** * create autotune threads */ @@ -7003,111 +5103,6 @@ int main (int argc, char **argv) * Clean up */ - for (uint device_id = 0; device_id < opencl_ctx->devices_cnt; device_id++) - { - hc_device_param_t *device_param = &opencl_ctx->devices_param[device_id]; - - 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_name_chksum); - 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) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_pws_buf); - if (device_param->d_pws_amp_buf) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_pws_amp_buf); - if (device_param->d_rules) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_rules); - if (device_param->d_rules_c) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_rules_c); - if (device_param->d_combs) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_combs); - if (device_param->d_combs_c) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_combs_c); - if (device_param->d_bfs) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_bfs); - if (device_param->d_bfs_c) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_bfs_c); - if (device_param->d_bitmap_s1_a) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_bitmap_s1_a); - if (device_param->d_bitmap_s1_b) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_bitmap_s1_b); - if (device_param->d_bitmap_s1_c) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_bitmap_s1_c); - if (device_param->d_bitmap_s1_d) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_bitmap_s1_d); - if (device_param->d_bitmap_s2_a) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_bitmap_s2_a); - if (device_param->d_bitmap_s2_b) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_bitmap_s2_b); - if (device_param->d_bitmap_s2_c) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_bitmap_s2_c); - if (device_param->d_bitmap_s2_d) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_bitmap_s2_d); - if (device_param->d_plain_bufs) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_plain_bufs); - if (device_param->d_digests_buf) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_digests_buf); - if (device_param->d_digests_shown) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_digests_shown); - if (device_param->d_salt_bufs) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_salt_bufs); - if (device_param->d_esalt_bufs) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_esalt_bufs); - if (device_param->d_tmps) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_tmps); - if (device_param->d_hooks) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_hooks); - if (device_param->d_result) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_result); - if (device_param->d_scryptV0_buf) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_scryptV0_buf); - if (device_param->d_scryptV1_buf) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_scryptV1_buf); - if (device_param->d_scryptV2_buf) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_scryptV2_buf); - if (device_param->d_scryptV3_buf) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_scryptV3_buf); - if (device_param->d_root_css_buf) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_root_css_buf); - if (device_param->d_markov_css_buf) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_markov_css_buf); - if (device_param->d_tm_c) CL_err |= hc_clReleaseMemObject (opencl_ctx->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 (opencl_ctx->ocl, device_param->kernel1); - if (device_param->kernel12) CL_err |= hc_clReleaseKernel (opencl_ctx->ocl, device_param->kernel12); - if (device_param->kernel2) CL_err |= hc_clReleaseKernel (opencl_ctx->ocl, device_param->kernel2); - if (device_param->kernel23) CL_err |= hc_clReleaseKernel (opencl_ctx->ocl, device_param->kernel23); - if (device_param->kernel3) CL_err |= hc_clReleaseKernel (opencl_ctx->ocl, device_param->kernel3); - if (device_param->kernel_mp) CL_err |= hc_clReleaseKernel (opencl_ctx->ocl, device_param->kernel_mp); - if (device_param->kernel_mp_l) CL_err |= hc_clReleaseKernel (opencl_ctx->ocl, device_param->kernel_mp_l); - if (device_param->kernel_mp_r) CL_err |= hc_clReleaseKernel (opencl_ctx->ocl, device_param->kernel_mp_r); - if (device_param->kernel_tm) CL_err |= hc_clReleaseKernel (opencl_ctx->ocl, device_param->kernel_tm); - if (device_param->kernel_amp) CL_err |= hc_clReleaseKernel (opencl_ctx->ocl, device_param->kernel_amp); - if (device_param->kernel_memset) CL_err |= hc_clReleaseKernel (opencl_ctx->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 (opencl_ctx->ocl, device_param->program); - if (device_param->program_mp) CL_err |= hc_clReleaseProgram (opencl_ctx->ocl, device_param->program_mp); - if (device_param->program_amp) CL_err |= hc_clReleaseProgram (opencl_ctx->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 (opencl_ctx->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 (opencl_ctx->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 #if defined (HAVE_HWMON) @@ -7267,6 +5262,10 @@ int main (int argc, char **argv) // free memory + opencl_session_destroy (opencl_ctx); + + opencl_ctx_devices_destroy (opencl_ctx); + local_free (masks); debugfile_destroy (debugfile_ctx); diff --git a/src/opencl.c b/src/opencl.c index 00f86d874..57b44d435 100644 --- a/src/opencl.c +++ b/src/opencl.c @@ -30,6 +30,7 @@ #include "potfile.h" #include "debugfile.h" #include "loopback.h" +#include "filenames.h" #include "data.h" #include "shared.h" #include "filehandling.h" @@ -40,6 +41,7 @@ extern hc_global_data_t data; extern hc_thread_mutex_t mux_counter; +extern hc_thread_mutex_t mux_hwmon; extern const int comptime; @@ -1495,7 +1497,7 @@ void opencl_ctx_destroy (opencl_ctx_t *opencl_ctx) myfree (opencl_ctx); } -int opencl_ctx_devices_init (opencl_ctx_t *opencl_ctx, const hashconfig_t *hashconfig, const tuning_db_t *tuning_db, const uint attack_mode, const uint quiet, const uint force, const uint benchmark, const uint machine_readable, const uint algorithm_pos) +int opencl_ctx_devices_init (opencl_ctx_t *opencl_ctx, const hashconfig_t *hashconfig, const tuning_db_t *tuning_db, const uint attack_mode, const bool quiet, const bool force, const bool benchmark, const bool machine_readable, const uint algorithm_pos) { if (opencl_ctx->disable == 1) return 0; @@ -2441,6 +2443,18 @@ int opencl_ctx_devices_init (opencl_ctx_t *opencl_ctx, const hashconfig_t *hashc void opencl_ctx_devices_destroy (opencl_ctx_t *opencl_ctx) { + for (uint device_id = 0; device_id < opencl_ctx->devices_cnt; device_id++) + { + hc_device_param_t *device_param = &opencl_ctx->devices_param[device_id]; + + if (device_param->skipped) continue; + + myfree (device_param->device_name); + myfree (device_param->device_name_chksum); + myfree (device_param->device_version); + myfree (device_param->driver_version); + } + opencl_ctx->devices_cnt = 0; opencl_ctx->devices_active = 0; @@ -2449,3 +2463,2037 @@ void opencl_ctx_devices_destroy (opencl_ctx_t *opencl_ctx) opencl_ctx->need_nvapi = 0; opencl_ctx->need_xnvctrl = 0; } + +int opencl_session_begin (opencl_ctx_t *opencl_ctx, const hashconfig_t *hashconfig, const hashes_t *hashes, const session_ctx_t *session_ctx) +{ + for (uint device_id = 0; device_id < opencl_ctx->devices_cnt; device_id++) + { + cl_int CL_err = CL_SUCCESS; + + /** + * host buffer + */ + + hc_device_param_t *device_param = &opencl_ctx->devices_param[device_id]; + + if (device_param->skipped) continue; + + /** + * device properties + */ + + const char *device_name_chksum = device_param->device_name_chksum; + const u32 device_processors = device_param->device_processors; + + /** + * create context for each device + */ + + 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 (opencl_ctx->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); + + CL_err = hc_clCreateCommandQueue (opencl_ctx->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 + * because of shared memory usage or bitslice + * there needs to be some upper limit, otherwise there's too much overhead + */ + + uint kernel_threads = MIN (KERNEL_THREADS_MAX, device_param->device_maxworkgroup_size); + + if (hashconfig->hash_mode == 8900) kernel_threads = 64; // Scrypt + if (hashconfig->hash_mode == 9300) kernel_threads = 64; // Scrypt + + if (device_param->device_type & CL_DEVICE_TYPE_CPU) + { + kernel_threads = KERNEL_THREADS_MAX_CPU; + } + + if (hashconfig->hash_mode == 1500) kernel_threads = 64; // DES + if (hashconfig->hash_mode == 3000) kernel_threads = 64; // DES + if (hashconfig->hash_mode == 3100) kernel_threads = 64; // DES + if (hashconfig->hash_mode == 3200) kernel_threads = 8; // Blowfish + if (hashconfig->hash_mode == 7500) kernel_threads = 64; // RC4 + if (hashconfig->hash_mode == 8500) kernel_threads = 64; // DES + if (hashconfig->hash_mode == 9000) kernel_threads = 8; // Blowfish + if (hashconfig->hash_mode == 9700) kernel_threads = 64; // RC4 + if (hashconfig->hash_mode == 9710) kernel_threads = 64; // RC4 + if (hashconfig->hash_mode == 9800) kernel_threads = 64; // RC4 + if (hashconfig->hash_mode == 9810) kernel_threads = 64; // RC4 + if (hashconfig->hash_mode == 10400) kernel_threads = 64; // RC4 + if (hashconfig->hash_mode == 10410) kernel_threads = 64; // RC4 + if (hashconfig->hash_mode == 10500) kernel_threads = 64; // RC4 + if (hashconfig->hash_mode == 13100) kernel_threads = 64; // RC4 + if (hashconfig->hash_mode == 14000) kernel_threads = 64; // DES + if (hashconfig->hash_mode == 14100) kernel_threads = 64; // DES + + device_param->kernel_threads = kernel_threads; + + device_param->hardware_power = device_processors * kernel_threads; + + /** + * create input buffers on device : calculate size of fixed memory buffers + */ + + size_t size_root_css = SP_PW_MAX * sizeof (cs_t); + size_t size_markov_css = SP_PW_MAX * CHARSIZ * sizeof (cs_t); + + device_param->size_root_css = size_root_css; + device_param->size_markov_css = size_markov_css; + + size_t size_results = sizeof (uint); + + device_param->size_results = size_results; + + size_t size_rules = session_ctx->kernel_rules_cnt * sizeof (kernel_rule_t); + size_t size_rules_c = KERNEL_RULES * sizeof (kernel_rule_t); + + size_t size_plains = hashes->digests_cnt * sizeof (plain_t); + size_t size_salts = hashes->salts_cnt * sizeof (salt_t); + size_t size_esalts = hashes->salts_cnt * hashconfig->esalt_size; + size_t size_shown = hashes->digests_cnt * sizeof (uint); + size_t size_digests = hashes->digests_cnt * hashconfig->dgst_size; + + device_param->size_plains = size_plains; + device_param->size_digests = size_digests; + device_param->size_shown = size_shown; + device_param->size_salts = size_salts; + + size_t size_combs = KERNEL_COMBS * sizeof (comb_t); + size_t size_bfs = KERNEL_BFS * sizeof (bf_t); + size_t size_tm = 32 * sizeof (bs_word_t); + + // scryptV stuff + + u32 scrypt_tmp_size = 0; + u32 scrypt_tmto_final = 0; + + size_t size_scrypt = 4; + + if ((hashconfig->hash_mode == 8900) || (hashconfig->hash_mode == 9300)) + { + // we need to check that all hashes have the same scrypt settings + + const u32 scrypt_N = hashes->salts_buf[0].scrypt_N; + const u32 scrypt_r = hashes->salts_buf[0].scrypt_r; + const u32 scrypt_p = hashes->salts_buf[0].scrypt_p; + + for (uint i = 1; i < hashes->salts_cnt; i++) + { + if ((hashes->salts_buf[i].scrypt_N != scrypt_N) + || (hashes->salts_buf[i].scrypt_r != scrypt_r) + || (hashes->salts_buf[i].scrypt_p != scrypt_p)) + { + log_error ("ERROR: Mixed scrypt settings not supported"); + + return -1; + } + } + + scrypt_tmp_size = (128 * scrypt_r * scrypt_p); + + uint tmto_start = 0; + uint tmto_stop = 10; + + if (session_ctx->scrypt_tmto) + { + tmto_start = session_ctx->scrypt_tmto; + } + else + { + // in case the user did not specify the tmto manually + // use some values known to run best (tested on 290x for AMD and GTX1080 for NV) + + if (hashconfig->hash_mode == 8900) + { + if (device_param->device_vendor_id == VENDOR_ID_AMD) + { + tmto_start = 3; + } + else if (device_param->device_vendor_id == VENDOR_ID_NV) + { + tmto_start = 2; + } + } + else if (hashconfig->hash_mode == 9300) + { + if (device_param->device_vendor_id == VENDOR_ID_AMD) + { + tmto_start = 2; + } + else if (device_param->device_vendor_id == VENDOR_ID_NV) + { + tmto_start = 4; + } + } + } + + device_param->kernel_accel_min = 1; + device_param->kernel_accel_max = 8; + + uint tmto; + + for (tmto = tmto_start; tmto < tmto_stop; tmto++) + { + size_scrypt = (128 * scrypt_r) * scrypt_N; + + size_scrypt /= 1u << tmto; + + size_scrypt *= device_param->device_processors * device_param->kernel_threads * device_param->kernel_accel_max; + + if ((size_scrypt / 4) > device_param->device_maxmem_alloc) + { + if (session_ctx->quiet == 0) log_info ("WARNING: Not enough single-block device memory allocatable to use --scrypt-tmto %d, increasing...", tmto); + + continue; + } + + if (size_scrypt > device_param->device_global_mem) + { + if (session_ctx->quiet == 0) log_info ("WARNING: Not enough total device memory allocatable to use --scrypt-tmto %d, increasing...", tmto); + + continue; + } + + for (uint salts_pos = 0; salts_pos < hashes->salts_cnt; salts_pos++) + { + scrypt_tmto_final = tmto; + } + + break; + } + + if (tmto == tmto_stop) + { + log_error ("ERROR: Can't allocate enough device memory"); + + return -1; + } + + if (session_ctx->quiet == 0) log_info ("SCRYPT tmto optimizer value set to: %u, mem: %" PRIu64 "\n", scrypt_tmto_final, size_scrypt); + } + + size_t size_scrypt4 = size_scrypt / 4; + + /** + * some algorithms need a fixed kernel-loops count + */ + + if (hashconfig->hash_mode == 1500 && session_ctx->attack_mode == ATTACK_MODE_BF) + { + const u32 kernel_loops_fixed = 1024; + + device_param->kernel_loops_min = kernel_loops_fixed; + device_param->kernel_loops_max = kernel_loops_fixed; + } + + if (hashconfig->hash_mode == 3000 && session_ctx->attack_mode == ATTACK_MODE_BF) + { + const u32 kernel_loops_fixed = 1024; + + device_param->kernel_loops_min = kernel_loops_fixed; + device_param->kernel_loops_max = kernel_loops_fixed; + } + + if (hashconfig->hash_mode == 8900) + { + const u32 kernel_loops_fixed = 1; + + device_param->kernel_loops_min = kernel_loops_fixed; + device_param->kernel_loops_max = kernel_loops_fixed; + } + + if (hashconfig->hash_mode == 9300) + { + const u32 kernel_loops_fixed = 1; + + device_param->kernel_loops_min = kernel_loops_fixed; + device_param->kernel_loops_max = kernel_loops_fixed; + } + + if (hashconfig->hash_mode == 12500) + { + const u32 kernel_loops_fixed = ROUNDS_RAR3 / 16; + + device_param->kernel_loops_min = kernel_loops_fixed; + device_param->kernel_loops_max = kernel_loops_fixed; + } + + if (hashconfig->hash_mode == 14000 && session_ctx->attack_mode == ATTACK_MODE_BF) + { + const u32 kernel_loops_fixed = 1024; + + device_param->kernel_loops_min = kernel_loops_fixed; + device_param->kernel_loops_max = kernel_loops_fixed; + } + + if (hashconfig->hash_mode == 14100 && session_ctx->attack_mode == ATTACK_MODE_BF) + { + const u32 kernel_loops_fixed = 1024; + + device_param->kernel_loops_min = kernel_loops_fixed; + device_param->kernel_loops_max = kernel_loops_fixed; + } + + u32 kernel_accel_min = device_param->kernel_accel_min; + u32 kernel_accel_max = device_param->kernel_accel_max; + + // find out if we would request too much memory on memory blocks which are based on kernel_accel + + size_t size_pws = 4; + size_t size_tmps = 4; + size_t size_hooks = 4; + + while (kernel_accel_max >= kernel_accel_min) + { + const u32 kernel_power_max = device_processors * kernel_threads * kernel_accel_max; + + // size_pws + + size_pws = kernel_power_max * sizeof (pw_t); + + // size_tmps + + switch (hashconfig->hash_mode) + { + case 400: size_tmps = kernel_power_max * sizeof (phpass_tmp_t); break; + case 500: size_tmps = kernel_power_max * sizeof (md5crypt_tmp_t); break; + case 501: size_tmps = kernel_power_max * sizeof (md5crypt_tmp_t); break; + case 1600: size_tmps = kernel_power_max * sizeof (md5crypt_tmp_t); break; + case 1800: size_tmps = kernel_power_max * sizeof (sha512crypt_tmp_t); break; + case 2100: size_tmps = kernel_power_max * sizeof (dcc2_tmp_t); break; + case 2500: size_tmps = kernel_power_max * sizeof (wpa_tmp_t); break; + case 3200: size_tmps = kernel_power_max * sizeof (bcrypt_tmp_t); break; + case 5200: size_tmps = kernel_power_max * sizeof (pwsafe3_tmp_t); break; + case 5800: size_tmps = kernel_power_max * sizeof (androidpin_tmp_t); break; + case 6211: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; + case 6212: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; + case 6213: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; + case 6221: size_tmps = kernel_power_max * sizeof (tc64_tmp_t); break; + case 6222: size_tmps = kernel_power_max * sizeof (tc64_tmp_t); break; + case 6223: size_tmps = kernel_power_max * sizeof (tc64_tmp_t); break; + case 6231: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; + case 6232: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; + case 6233: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; + case 6241: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; + case 6242: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; + case 6243: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; + case 6300: size_tmps = kernel_power_max * sizeof (md5crypt_tmp_t); break; + case 6400: size_tmps = kernel_power_max * sizeof (sha256aix_tmp_t); break; + case 6500: size_tmps = kernel_power_max * sizeof (sha512aix_tmp_t); break; + case 6600: size_tmps = kernel_power_max * sizeof (agilekey_tmp_t); break; + case 6700: size_tmps = kernel_power_max * sizeof (sha1aix_tmp_t); break; + case 6800: size_tmps = kernel_power_max * sizeof (lastpass_tmp_t); break; + case 7100: size_tmps = kernel_power_max * sizeof (pbkdf2_sha512_tmp_t); break; + case 7200: size_tmps = kernel_power_max * sizeof (pbkdf2_sha512_tmp_t); break; + case 7400: size_tmps = kernel_power_max * sizeof (sha256crypt_tmp_t); break; + case 7900: size_tmps = kernel_power_max * sizeof (drupal7_tmp_t); break; + case 8200: size_tmps = kernel_power_max * sizeof (pbkdf2_sha512_tmp_t); break; + case 8800: size_tmps = kernel_power_max * sizeof (androidfde_tmp_t); break; + case 8900: size_tmps = kernel_power_max * scrypt_tmp_size; break; + case 9000: size_tmps = kernel_power_max * sizeof (pwsafe2_tmp_t); break; + case 9100: size_tmps = kernel_power_max * sizeof (lotus8_tmp_t); break; + case 9200: size_tmps = kernel_power_max * sizeof (pbkdf2_sha256_tmp_t); break; + case 9300: size_tmps = kernel_power_max * scrypt_tmp_size; break; + case 9400: size_tmps = kernel_power_max * sizeof (office2007_tmp_t); break; + case 9500: size_tmps = kernel_power_max * sizeof (office2010_tmp_t); break; + case 9600: size_tmps = kernel_power_max * sizeof (office2013_tmp_t); break; + case 10000: size_tmps = kernel_power_max * sizeof (pbkdf2_sha256_tmp_t); break; + case 10200: size_tmps = kernel_power_max * sizeof (cram_md5_t); break; + case 10300: size_tmps = kernel_power_max * sizeof (saph_sha1_tmp_t); break; + case 10500: size_tmps = kernel_power_max * sizeof (pdf14_tmp_t); break; + case 10700: size_tmps = kernel_power_max * sizeof (pdf17l8_tmp_t); break; + case 10900: size_tmps = kernel_power_max * sizeof (pbkdf2_sha256_tmp_t); break; + case 11300: size_tmps = kernel_power_max * sizeof (bitcoin_wallet_tmp_t); break; + case 11600: size_tmps = kernel_power_max * sizeof (seven_zip_tmp_t); break; + case 11900: size_tmps = kernel_power_max * sizeof (pbkdf2_md5_tmp_t); break; + case 12000: size_tmps = kernel_power_max * sizeof (pbkdf2_sha1_tmp_t); break; + case 12100: size_tmps = kernel_power_max * sizeof (pbkdf2_sha512_tmp_t); break; + case 12200: size_tmps = kernel_power_max * sizeof (ecryptfs_tmp_t); break; + case 12300: size_tmps = kernel_power_max * sizeof (oraclet_tmp_t); break; + case 12400: size_tmps = kernel_power_max * sizeof (bsdicrypt_tmp_t); break; + case 12500: size_tmps = kernel_power_max * sizeof (rar3_tmp_t); break; + case 12700: size_tmps = kernel_power_max * sizeof (mywallet_tmp_t); break; + case 12800: size_tmps = kernel_power_max * sizeof (pbkdf2_sha256_tmp_t); break; + case 12900: size_tmps = kernel_power_max * sizeof (pbkdf2_sha256_tmp_t); break; + case 13000: size_tmps = kernel_power_max * sizeof (pbkdf2_sha256_tmp_t); break; + case 13200: size_tmps = kernel_power_max * sizeof (axcrypt_tmp_t); break; + case 13400: size_tmps = kernel_power_max * sizeof (keepass_tmp_t); break; + case 13600: size_tmps = kernel_power_max * sizeof (pbkdf2_sha1_tmp_t); break; + case 13711: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; + case 13712: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; + case 13713: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; + case 13721: size_tmps = kernel_power_max * sizeof (tc64_tmp_t); break; + case 13722: size_tmps = kernel_power_max * sizeof (tc64_tmp_t); break; + case 13723: size_tmps = kernel_power_max * sizeof (tc64_tmp_t); break; + case 13731: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; + case 13732: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; + case 13733: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; + case 13741: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; + case 13742: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; + case 13743: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; + case 13751: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; + case 13752: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; + case 13753: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; + case 13761: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; + case 13762: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; + case 13763: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break; + }; + + // size_hooks + + if ((hashconfig->opts_type & OPTS_TYPE_HOOK12) || (hashconfig->opts_type & OPTS_TYPE_HOOK23)) + { + switch (hashconfig->hash_mode) + { + } + } + + // now check if all device-memory sizes which depend on the kernel_accel_max amplifier are within its boundaries + // if not, decrease amplifier and try again + + int memory_limit_hit = 0; + + if (size_pws > device_param->device_maxmem_alloc) memory_limit_hit = 1; + if (size_tmps > device_param->device_maxmem_alloc) memory_limit_hit = 1; + if (size_hooks > device_param->device_maxmem_alloc) memory_limit_hit = 1; + + const u64 size_total + = session_ctx->bitmap_size + + session_ctx->bitmap_size + + session_ctx->bitmap_size + + session_ctx->bitmap_size + + session_ctx->bitmap_size + + session_ctx->bitmap_size + + session_ctx->bitmap_size + + session_ctx->bitmap_size + + size_bfs + + size_combs + + size_digests + + size_esalts + + size_hooks + + size_markov_css + + size_plains + + size_pws + + size_pws // not a bug + + size_results + + size_root_css + + size_rules + + size_rules_c + + size_salts + + size_scrypt4 + + size_scrypt4 + + size_scrypt4 + + size_scrypt4 + + size_shown + + size_tm + + size_tmps; + + if (size_total > device_param->device_global_mem) memory_limit_hit = 1; + + if (memory_limit_hit == 1) + { + kernel_accel_max--; + + continue; + } + + break; + } + + if (kernel_accel_max < kernel_accel_min) + { + log_error ("- Device #%u: Device does not provide enough allocatable device-memory to handle this attack", device_id + 1); + + return -1; + } + + device_param->kernel_accel_min = kernel_accel_min; + device_param->kernel_accel_max = kernel_accel_max; + + /* + if (kernel_accel_max < kernel_accel) + { + if (session_ctx->quiet == 0) log_info ("- Device #%u: Reduced maximum kernel-accel to %u", device_id + 1, kernel_accel_max); + + device_param->kernel_accel = kernel_accel_max; + } + */ + + device_param->size_bfs = size_bfs; + device_param->size_combs = size_combs; + device_param->size_rules = size_rules; + device_param->size_rules_c = size_rules_c; + device_param->size_pws = size_pws; + device_param->size_tmps = size_tmps; + device_param->size_hooks = size_hooks; + + /** + * default building options + */ + + if (chdir (session_ctx->cpath_real) == -1) + { + log_error ("ERROR: %s: %s", session_ctx->cpath_real, strerror (errno)); + + return -1; + } + + char build_opts[1024] = { 0 }; + + #if defined (_WIN) + snprintf (build_opts, sizeof (build_opts) - 1, "-I \"%s\"", session_ctx->cpath_real); + #else + snprintf (build_opts, sizeof (build_opts) - 1, "-I %s", session_ctx->cpath_real); + #endif + + // include check + // this test needs to be done manually because of osx opencl runtime + // if there's a problem with permission, its not reporting back and erroring out silently + + #define files_cnt 15 + + const char *files_names[files_cnt] = + { + "inc_cipher_aes256.cl", + "inc_cipher_serpent256.cl", + "inc_cipher_twofish256.cl", + "inc_common.cl", + "inc_comp_multi_bs.cl", + "inc_comp_multi.cl", + "inc_comp_single_bs.cl", + "inc_comp_single.cl", + "inc_hash_constants.h", + "inc_hash_functions.cl", + "inc_rp.cl", + "inc_rp.h", + "inc_simd.cl", + "inc_types.cl", + "inc_vendor.cl", + }; + + for (int i = 0; i < files_cnt; i++) + { + FILE *fd = fopen (files_names[i], "r"); + + if (fd == NULL) + { + log_error ("ERROR: %s: fopen(): %s", files_names[i], strerror (errno)); + + return -1; + } + + char buf[1]; + + size_t n = fread (buf, 1, 1, fd); + + if (n != 1) + { + log_error ("ERROR: %s: fread(): %s", files_names[i], strerror (errno)); + + return -1; + } + + fclose (fd); + } + + // we don't have sm_* on vendors not NV but it doesn't matter + + char build_opts_new[1024] = { 0 }; + + #if defined (DEBUG) + snprintf (build_opts_new, sizeof (build_opts_new) - 1, "%s -D VENDOR_ID=%u -D CUDA_ARCH=%d -D VECT_SIZE=%u -D DEVICE_TYPE=%u -D DGST_R0=%u -D DGST_R1=%u -D DGST_R2=%u -D DGST_R3=%u -D DGST_ELEM=%u -D KERN_TYPE=%u -D _unroll -cl-std=CL1.1", build_opts, device_param->device_vendor_id, (device_param->sm_major * 100) + device_param->sm_minor, device_param->vector_width, (u32) device_param->device_type, hashconfig->dgst_pos0, hashconfig->dgst_pos1, hashconfig->dgst_pos2, hashconfig->dgst_pos3, hashconfig->dgst_size / 4, hashconfig->kern_type); + #else + snprintf (build_opts_new, sizeof (build_opts_new) - 1, "%s -D VENDOR_ID=%u -D CUDA_ARCH=%d -D VECT_SIZE=%u -D DEVICE_TYPE=%u -D DGST_R0=%u -D DGST_R1=%u -D DGST_R2=%u -D DGST_R3=%u -D DGST_ELEM=%u -D KERN_TYPE=%u -D _unroll -cl-std=CL1.1 -w", build_opts, device_param->device_vendor_id, (device_param->sm_major * 100) + device_param->sm_minor, device_param->vector_width, (u32) device_param->device_type, hashconfig->dgst_pos0, hashconfig->dgst_pos1, hashconfig->dgst_pos2, hashconfig->dgst_pos3, hashconfig->dgst_size / 4, hashconfig->kern_type); + #endif + + strncpy (build_opts, build_opts_new, sizeof (build_opts)); + + #if defined (DEBUG) + log_info ("- Device #%u: build_opts '%s'\n", device_id + 1, build_opts); + #endif + + /** + * main kernel + */ + + { + /** + * kernel source filename + */ + + char source_file[256] = { 0 }; + + generate_source_kernel_filename (hashconfig->attack_exec, session_ctx->attack_kern, hashconfig->kern_type, session_ctx->shared_dir, source_file); + + struct stat sst; + + if (stat (source_file, &sst) == -1) + { + log_error ("ERROR: %s: %s", source_file, strerror (errno)); + + return -1; + } + + /** + * kernel cached filename + */ + + char cached_file[256] = { 0 }; + + generate_cached_kernel_filename (hashconfig->attack_exec, session_ctx->attack_kern, hashconfig->kern_type, session_ctx->profile_dir, device_name_chksum, cached_file); + + int cached = 1; + + struct stat cst; + + if ((stat (cached_file, &cst) == -1) || cst.st_size == 0) + { + cached = 0; + } + + /** + * kernel compile or load + */ + + size_t *kernel_lengths = (size_t *) mymalloc (sizeof (size_t)); + + const u8 **kernel_sources = (const u8 **) mymalloc (sizeof (u8 *)); + + if (opencl_ctx->force_jit_compilation == -1) + { + if (cached == 0) + { + if (session_ctx->quiet == 0) log_info ("- Device #%u: Kernel %s not found in cache! Building may take a while...", device_id + 1, filename_from_filepath (cached_file)); + + load_kernel (source_file, 1, kernel_lengths, kernel_sources); + + CL_err = hc_clCreateProgramWithSource (opencl_ctx->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 (opencl_ctx->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; + } + + size_t build_log_size = 0; + + /* + CL_err = hc_clGetProgramBuildInfo (opencl_ctx->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; + } + */ + + hc_clGetProgramBuildInfo (opencl_ctx->ocl, device_param->program, device_param->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &build_log_size); + + #if defined (DEBUG) + if ((build_log_size != 0) || (CL_err != CL_SUCCESS)) + #else + if (CL_err != CL_SUCCESS) + #endif + { + char *build_log = (char *) mymalloc (build_log_size + 1); + + CL_err = hc_clGetProgramBuildInfo (opencl_ctx->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); + + myfree (build_log); + } + + if (CL_err != CL_SUCCESS) + { + device_param->skipped = true; + + log_info ("- Device #%u: Kernel %s build failure. Proceeding without this device.", device_id + 1, source_file); + + continue; + } + + size_t binary_size; + + CL_err = hc_clGetProgramInfo (opencl_ctx->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); + + CL_err = hc_clGetProgramInfo (opencl_ctx->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); + + local_free (binary); + } + else + { + #if defined (DEBUG) + log_info ("- Device #%u: Kernel %s (%ld bytes)", device_id + 1, cached_file, cst.st_size); + #endif + + load_kernel (cached_file, 1, kernel_lengths, kernel_sources); + + CL_err = hc_clCreateProgramWithBinary (opencl_ctx->ocl, device_param->context, 1, &device_param->device, kernel_lengths, (const u8 **) kernel_sources, NULL, &device_param->program); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clCreateProgramWithBinary(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } + + CL_err = hc_clBuildProgram (opencl_ctx->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 + { + #if defined (DEBUG) + log_info ("- Device #%u: Kernel %s (%ld bytes)", device_id + 1, source_file, sst.st_size); + #endif + + load_kernel (source_file, 1, kernel_lengths, kernel_sources); + + CL_err = hc_clCreateProgramWithSource (opencl_ctx->ocl, device_param->context, 1, (const char **) kernel_sources, NULL, &device_param->program); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clCreateProgramWithSource(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } + + char build_opts_update[1024] = { 0 }; + + if (opencl_ctx->force_jit_compilation == 1500) + { + snprintf (build_opts_update, sizeof (build_opts_update) - 1, "%s -DDESCRYPT_SALT=%u", build_opts, hashes->salts_buf[0].salt_buf[0]); + } + else if (opencl_ctx->force_jit_compilation == 8900) + { + snprintf (build_opts_update, sizeof (build_opts_update) - 1, "%s -DSCRYPT_N=%u -DSCRYPT_R=%u -DSCRYPT_P=%u -DSCRYPT_TMTO=%u -DSCRYPT_TMP_ELEM=%u", build_opts, hashes->salts_buf[0].scrypt_N, hashes->salts_buf[0].scrypt_r, hashes->salts_buf[0].scrypt_p, 1 << scrypt_tmto_final, scrypt_tmp_size / 16); + } + else + { + snprintf (build_opts_update, sizeof (build_opts_update) - 1, "%s", build_opts); + } + + CL_err = hc_clBuildProgram (opencl_ctx->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; + } + + size_t build_log_size = 0; + + /* + CL_err = hc_clGetProgramBuildInfo (opencl_ctx->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; + } + */ + + hc_clGetProgramBuildInfo (opencl_ctx->ocl, device_param->program, device_param->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &build_log_size); + + #if defined (DEBUG) + if ((build_log_size != 0) || (CL_err != CL_SUCCESS)) + #else + if (CL_err != CL_SUCCESS) + #endif + { + char *build_log = (char *) mymalloc (build_log_size + 1); + + CL_err = hc_clGetProgramBuildInfo (opencl_ctx->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); + + myfree (build_log); + } + + if (CL_err != CL_SUCCESS) + { + device_param->skipped = true; + + log_info ("- Device #%u: Kernel %s build failure. Proceeding without this device.", device_id + 1, source_file); + } + } + + local_free (kernel_lengths); + local_free (kernel_sources[0]); + local_free (kernel_sources); + } + + /** + * word generator kernel + */ + + if (session_ctx->attack_mode != ATTACK_MODE_STRAIGHT) + { + /** + * kernel mp source filename + */ + + char source_file[256] = { 0 }; + + generate_source_kernel_mp_filename (hashconfig->opti_type, hashconfig->opts_type, session_ctx->shared_dir, source_file); + + struct stat sst; + + if (stat (source_file, &sst) == -1) + { + log_error ("ERROR: %s: %s", source_file, strerror (errno)); + + return -1; + } + + /** + * kernel mp cached filename + */ + + char cached_file[256] = { 0 }; + + generate_cached_kernel_mp_filename (hashconfig->opti_type, hashconfig->opts_type, session_ctx->profile_dir, device_name_chksum, cached_file); + + int cached = 1; + + struct stat cst; + + if (stat (cached_file, &cst) == -1) + { + cached = 0; + } + + /** + * kernel compile or load + */ + + size_t *kernel_lengths = (size_t *) mymalloc (sizeof (size_t)); + + const u8 **kernel_sources = (const u8 **) mymalloc (sizeof (u8 *)); + + if (cached == 0) + { + if (session_ctx->quiet == 0) log_info ("- Device #%u: Kernel %s not found in cache! Building may take a while...", device_id + 1, filename_from_filepath (cached_file)); + if (session_ctx->quiet == 0) log_info (""); + + load_kernel (source_file, 1, kernel_lengths, kernel_sources); + + CL_err = hc_clCreateProgramWithSource (opencl_ctx->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; + } + + CL_err = hc_clBuildProgram (opencl_ctx->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; + } + + size_t build_log_size = 0; + + /* + CL_err = hc_clGetProgramBuildInfo (opencl_ctx->ocl, device_param->program_mp, 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; + } + */ + + hc_clGetProgramBuildInfo (opencl_ctx->ocl, device_param->program_mp, device_param->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &build_log_size); + + #if defined (DEBUG) + if ((build_log_size != 0) || (CL_err != CL_SUCCESS)) + #else + if (CL_err != CL_SUCCESS) + #endif + { + char *build_log = (char *) mymalloc (build_log_size + 1); + + CL_err = hc_clGetProgramBuildInfo (opencl_ctx->ocl, device_param->program_mp, 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); + + myfree (build_log); + } + + if (CL_err != CL_SUCCESS) + { + device_param->skipped = true; + + log_info ("- Device #%u: Kernel %s build failure. Proceeding without this device.", device_id + 1, source_file); + + continue; + } + + size_t binary_size; + + CL_err = hc_clGetProgramInfo (opencl_ctx->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); + + CL_err = hc_clGetProgramInfo (opencl_ctx->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); + + local_free (binary); + } + else + { + #if defined (DEBUG) + log_info ("- Device #%u: Kernel %s (%ld bytes)", device_id + 1, cached_file, cst.st_size); + #endif + + load_kernel (cached_file, 1, kernel_lengths, kernel_sources); + + CL_err = hc_clCreateProgramWithBinary (opencl_ctx->ocl, device_param->context, 1, &device_param->device, kernel_lengths, (const u8 **) kernel_sources, NULL, &device_param->program_mp); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clCreateProgramWithBinary(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } + + CL_err = hc_clBuildProgram (opencl_ctx->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); + local_free (kernel_sources[0]); + local_free (kernel_sources); + } + + /** + * amplifier kernel + */ + + if (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL) + { + + } + else + { + /** + * kernel amp source filename + */ + + char source_file[256] = { 0 }; + + generate_source_kernel_amp_filename (session_ctx->attack_kern, session_ctx->shared_dir, source_file); + + struct stat sst; + + if (stat (source_file, &sst) == -1) + { + log_error ("ERROR: %s: %s", source_file, strerror (errno)); + + return -1; + } + + /** + * kernel amp cached filename + */ + + char cached_file[256] = { 0 }; + + generate_cached_kernel_amp_filename (session_ctx->attack_kern, session_ctx->profile_dir, device_name_chksum, cached_file); + + int cached = 1; + + struct stat cst; + + if (stat (cached_file, &cst) == -1) + { + cached = 0; + } + + /** + * kernel compile or load + */ + + size_t *kernel_lengths = (size_t *) mymalloc (sizeof (size_t)); + + const u8 **kernel_sources = (const u8 **) mymalloc (sizeof (u8 *)); + + if (cached == 0) + { + if (session_ctx->quiet == 0) log_info ("- Device #%u: Kernel %s not found in cache! Building may take a while...", device_id + 1, filename_from_filepath (cached_file)); + if (session_ctx->quiet == 0) log_info (""); + + load_kernel (source_file, 1, kernel_lengths, kernel_sources); + + CL_err = hc_clCreateProgramWithSource (opencl_ctx->ocl, device_param->context, 1, (const char **) kernel_sources, NULL, &device_param->program_amp); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clCreateProgramWithSource(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } + + CL_err = hc_clBuildProgram (opencl_ctx->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; + } + + size_t build_log_size = 0; + + /* + CL_err = hc_clGetProgramBuildInfo (opencl_ctx->ocl, device_param->program_amp, 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; + } + */ + + hc_clGetProgramBuildInfo (opencl_ctx->ocl, device_param->program_amp, device_param->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &build_log_size); + + #if defined (DEBUG) + if ((build_log_size != 0) || (CL_err != CL_SUCCESS)) + #else + if (CL_err != CL_SUCCESS) + #endif + { + char *build_log = (char *) mymalloc (build_log_size + 1); + + CL_err = hc_clGetProgramBuildInfo (opencl_ctx->ocl, device_param->program_amp, 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); + + myfree (build_log); + } + + if (CL_err != CL_SUCCESS) + { + device_param->skipped = true; + + log_info ("- Device #%u: Kernel %s build failure. Proceed without this device.", device_id + 1, source_file); + + continue; + } + + size_t binary_size; + + CL_err = hc_clGetProgramInfo (opencl_ctx->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); + + CL_err = hc_clGetProgramInfo (opencl_ctx->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); + + local_free (binary); + } + else + { + #if defined (DEBUG) + if (session_ctx->quiet == 0) log_info ("- Device #%u: Kernel %s (%ld bytes)", device_id + 1, cached_file, cst.st_size); + #endif + + load_kernel (cached_file, 1, kernel_lengths, kernel_sources); + + CL_err = hc_clCreateProgramWithBinary (opencl_ctx->ocl, device_param->context, 1, &device_param->device, kernel_lengths, (const u8 **) kernel_sources, NULL, &device_param->program_amp); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clCreateProgramWithBinary(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } + + CL_err = hc_clBuildProgram (opencl_ctx->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[0]); + local_free (kernel_sources); + } + + // return back to the folder we came from initially (workaround) + + if (chdir (session_ctx->cwd) == -1) + { + log_error ("ERROR: %s: %s", session_ctx->cwd, strerror (errno)); + + return -1; + } + + // some algorithm collide too fast, make that impossible + + if (session_ctx->benchmark == 1) + { + ((uint *) hashes->digests_buf)[0] = -1u; + ((uint *) hashes->digests_buf)[1] = -1u; + ((uint *) hashes->digests_buf)[2] = -1u; + ((uint *) hashes->digests_buf)[3] = -1u; + } + + /** + * global buffers + */ + + CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_ONLY, size_pws, NULL, &device_param->d_pws_buf); + CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_ONLY, size_pws, NULL, &device_param->d_pws_amp_buf); + CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_WRITE, size_tmps, NULL, &device_param->d_tmps); + CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_WRITE, size_hooks, NULL, &device_param->d_hooks); + CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_ONLY, session_ctx->bitmap_size, NULL, &device_param->d_bitmap_s1_a); + CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_ONLY, session_ctx->bitmap_size, NULL, &device_param->d_bitmap_s1_b); + CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_ONLY, session_ctx->bitmap_size, NULL, &device_param->d_bitmap_s1_c); + CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_ONLY, session_ctx->bitmap_size, NULL, &device_param->d_bitmap_s1_d); + CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_ONLY, session_ctx->bitmap_size, NULL, &device_param->d_bitmap_s2_a); + CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_ONLY, session_ctx->bitmap_size, NULL, &device_param->d_bitmap_s2_b); + CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_ONLY, session_ctx->bitmap_size, NULL, &device_param->d_bitmap_s2_c); + CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_ONLY, session_ctx->bitmap_size, NULL, &device_param->d_bitmap_s2_d); + CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_WRITE, size_plains, NULL, &device_param->d_plain_bufs); + CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_ONLY, size_digests, NULL, &device_param->d_digests_buf); + CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_WRITE, size_shown, NULL, &device_param->d_digests_shown); + CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_ONLY, size_salts, NULL, &device_param->d_salt_bufs); + CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_WRITE, size_results, NULL, &device_param->d_result); + CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_WRITE, size_scrypt4, NULL, &device_param->d_scryptV0_buf); + CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_WRITE, size_scrypt4, NULL, &device_param->d_scryptV1_buf); + CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_WRITE, size_scrypt4, NULL, &device_param->d_scryptV2_buf); + CL_err |= hc_clCreateBuffer (opencl_ctx->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 (opencl_ctx->ocl, device_param->command_queue, device_param->d_bitmap_s1_a, CL_TRUE, 0, session_ctx->bitmap_size, session_ctx->bitmap_s1_a, 0, NULL, NULL); + CL_err |= hc_clEnqueueWriteBuffer (opencl_ctx->ocl, device_param->command_queue, device_param->d_bitmap_s1_b, CL_TRUE, 0, session_ctx->bitmap_size, session_ctx->bitmap_s1_b, 0, NULL, NULL); + CL_err |= hc_clEnqueueWriteBuffer (opencl_ctx->ocl, device_param->command_queue, device_param->d_bitmap_s1_c, CL_TRUE, 0, session_ctx->bitmap_size, session_ctx->bitmap_s1_c, 0, NULL, NULL); + CL_err |= hc_clEnqueueWriteBuffer (opencl_ctx->ocl, device_param->command_queue, device_param->d_bitmap_s1_d, CL_TRUE, 0, session_ctx->bitmap_size, session_ctx->bitmap_s1_d, 0, NULL, NULL); + CL_err |= hc_clEnqueueWriteBuffer (opencl_ctx->ocl, device_param->command_queue, device_param->d_bitmap_s2_a, CL_TRUE, 0, session_ctx->bitmap_size, session_ctx->bitmap_s2_a, 0, NULL, NULL); + CL_err |= hc_clEnqueueWriteBuffer (opencl_ctx->ocl, device_param->command_queue, device_param->d_bitmap_s2_b, CL_TRUE, 0, session_ctx->bitmap_size, session_ctx->bitmap_s2_b, 0, NULL, NULL); + CL_err |= hc_clEnqueueWriteBuffer (opencl_ctx->ocl, device_param->command_queue, device_param->d_bitmap_s2_c, CL_TRUE, 0, session_ctx->bitmap_size, session_ctx->bitmap_s2_c, 0, NULL, NULL); + CL_err |= hc_clEnqueueWriteBuffer (opencl_ctx->ocl, device_param->command_queue, device_param->d_bitmap_s2_d, CL_TRUE, 0, session_ctx->bitmap_size, session_ctx->bitmap_s2_d, 0, NULL, NULL); + CL_err |= hc_clEnqueueWriteBuffer (opencl_ctx->ocl, device_param->command_queue, device_param->d_digests_buf, CL_TRUE, 0, size_digests, hashes->digests_buf, 0, NULL, NULL); + CL_err |= hc_clEnqueueWriteBuffer (opencl_ctx->ocl, device_param->command_queue, device_param->d_digests_shown, CL_TRUE, 0, size_shown, hashes->digests_shown, 0, NULL, NULL); + CL_err |= hc_clEnqueueWriteBuffer (opencl_ctx->ocl, device_param->command_queue, device_param->d_salt_bufs, CL_TRUE, 0, size_salts, hashes->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 (session_ctx->attack_kern == ATTACK_KERN_STRAIGHT) + { + CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_ONLY, size_rules, NULL, &device_param->d_rules); + CL_err |= hc_clCreateBuffer (opencl_ctx->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; + } + + CL_err = hc_clEnqueueWriteBuffer (opencl_ctx->ocl, device_param->command_queue, device_param->d_rules, CL_TRUE, 0, size_rules, session_ctx->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 (session_ctx->attack_kern == ATTACK_KERN_COMBI) + { + CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_ONLY, size_combs, NULL, &device_param->d_combs); + CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_ONLY, size_combs, NULL, &device_param->d_combs_c); + CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_ONLY, size_root_css, NULL, &device_param->d_root_css_buf); + CL_err |= hc_clCreateBuffer (opencl_ctx->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 (session_ctx->attack_kern == ATTACK_KERN_BF) + { + CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_ONLY, size_bfs, NULL, &device_param->d_bfs); + CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_ONLY, size_bfs, NULL, &device_param->d_bfs_c); + CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_ONLY, size_tm, NULL, &device_param->d_tm_c); + CL_err |= hc_clCreateBuffer (opencl_ctx->ocl, device_param->context, CL_MEM_READ_ONLY, size_root_css, NULL, &device_param->d_root_css_buf); + CL_err |= hc_clCreateBuffer (opencl_ctx->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) + { + CL_err = hc_clCreateBuffer (opencl_ctx->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 (opencl_ctx->ocl, device_param->command_queue, device_param->d_esalt_bufs, CL_TRUE, 0, size_esalts, hashes->esalts_buf, 0, NULL, NULL); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clEnqueueWriteBuffer(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } + } + + /** + * main host data + */ + + pw_t *pws_buf = (pw_t *) mymalloc (size_pws); + + device_param->pws_buf = pws_buf; + + comb_t *combs_buf = (comb_t *) mycalloc (KERNEL_COMBS, sizeof (comb_t)); + + device_param->combs_buf = combs_buf; + + void *hooks_buf = mymalloc (size_hooks); + + device_param->hooks_buf = hooks_buf; + + /** + * kernel args + */ + + device_param->kernel_params_buf32[24] = session_ctx->bitmap_mask; + device_param->kernel_params_buf32[25] = session_ctx->bitmap_shift1; + device_param->kernel_params_buf32[26] = session_ctx->bitmap_shift2; + device_param->kernel_params_buf32[27] = 0; // salt_pos + device_param->kernel_params_buf32[28] = 0; // loop_pos + device_param->kernel_params_buf32[29] = 0; // loop_cnt + device_param->kernel_params_buf32[30] = 0; // kernel_rules_cnt + device_param->kernel_params_buf32[31] = 0; // digests_cnt + device_param->kernel_params_buf32[32] = 0; // digests_offset + device_param->kernel_params_buf32[33] = 0; // combs_mode + device_param->kernel_params_buf32[34] = 0; // gid_max + + device_param->kernel_params[ 0] = (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL) + ? &device_param->d_pws_buf + : &device_param->d_pws_amp_buf; + device_param->kernel_params[ 1] = &device_param->d_rules_c; + device_param->kernel_params[ 2] = &device_param->d_combs_c; + device_param->kernel_params[ 3] = &device_param->d_bfs_c; + device_param->kernel_params[ 4] = &device_param->d_tmps; + device_param->kernel_params[ 5] = &device_param->d_hooks; + device_param->kernel_params[ 6] = &device_param->d_bitmap_s1_a; + device_param->kernel_params[ 7] = &device_param->d_bitmap_s1_b; + device_param->kernel_params[ 8] = &device_param->d_bitmap_s1_c; + device_param->kernel_params[ 9] = &device_param->d_bitmap_s1_d; + device_param->kernel_params[10] = &device_param->d_bitmap_s2_a; + device_param->kernel_params[11] = &device_param->d_bitmap_s2_b; + device_param->kernel_params[12] = &device_param->d_bitmap_s2_c; + device_param->kernel_params[13] = &device_param->d_bitmap_s2_d; + device_param->kernel_params[14] = &device_param->d_plain_bufs; + device_param->kernel_params[15] = &device_param->d_digests_buf; + device_param->kernel_params[16] = &device_param->d_digests_shown; + device_param->kernel_params[17] = &device_param->d_salt_bufs; + device_param->kernel_params[18] = &device_param->d_esalt_bufs; + device_param->kernel_params[19] = &device_param->d_result; + device_param->kernel_params[20] = &device_param->d_scryptV0_buf; + device_param->kernel_params[21] = &device_param->d_scryptV1_buf; + device_param->kernel_params[22] = &device_param->d_scryptV2_buf; + device_param->kernel_params[23] = &device_param->d_scryptV3_buf; + device_param->kernel_params[24] = &device_param->kernel_params_buf32[24]; + device_param->kernel_params[25] = &device_param->kernel_params_buf32[25]; + device_param->kernel_params[26] = &device_param->kernel_params_buf32[26]; + device_param->kernel_params[27] = &device_param->kernel_params_buf32[27]; + device_param->kernel_params[28] = &device_param->kernel_params_buf32[28]; + device_param->kernel_params[29] = &device_param->kernel_params_buf32[29]; + device_param->kernel_params[30] = &device_param->kernel_params_buf32[30]; + device_param->kernel_params[31] = &device_param->kernel_params_buf32[31]; + device_param->kernel_params[32] = &device_param->kernel_params_buf32[32]; + device_param->kernel_params[33] = &device_param->kernel_params_buf32[33]; + device_param->kernel_params[34] = &device_param->kernel_params_buf32[34]; + + device_param->kernel_params_mp_buf64[3] = 0; + device_param->kernel_params_mp_buf32[4] = 0; + device_param->kernel_params_mp_buf32[5] = 0; + device_param->kernel_params_mp_buf32[6] = 0; + device_param->kernel_params_mp_buf32[7] = 0; + device_param->kernel_params_mp_buf32[8] = 0; + + device_param->kernel_params_mp[0] = NULL; + device_param->kernel_params_mp[1] = NULL; + device_param->kernel_params_mp[2] = NULL; + device_param->kernel_params_mp[3] = &device_param->kernel_params_mp_buf64[3]; + device_param->kernel_params_mp[4] = &device_param->kernel_params_mp_buf32[4]; + device_param->kernel_params_mp[5] = &device_param->kernel_params_mp_buf32[5]; + device_param->kernel_params_mp[6] = &device_param->kernel_params_mp_buf32[6]; + device_param->kernel_params_mp[7] = &device_param->kernel_params_mp_buf32[7]; + device_param->kernel_params_mp[8] = &device_param->kernel_params_mp_buf32[8]; + + device_param->kernel_params_mp_l_buf64[3] = 0; + device_param->kernel_params_mp_l_buf32[4] = 0; + device_param->kernel_params_mp_l_buf32[5] = 0; + device_param->kernel_params_mp_l_buf32[6] = 0; + device_param->kernel_params_mp_l_buf32[7] = 0; + device_param->kernel_params_mp_l_buf32[8] = 0; + device_param->kernel_params_mp_l_buf32[9] = 0; + + device_param->kernel_params_mp_l[0] = NULL; + device_param->kernel_params_mp_l[1] = NULL; + device_param->kernel_params_mp_l[2] = NULL; + device_param->kernel_params_mp_l[3] = &device_param->kernel_params_mp_l_buf64[3]; + device_param->kernel_params_mp_l[4] = &device_param->kernel_params_mp_l_buf32[4]; + device_param->kernel_params_mp_l[5] = &device_param->kernel_params_mp_l_buf32[5]; + device_param->kernel_params_mp_l[6] = &device_param->kernel_params_mp_l_buf32[6]; + device_param->kernel_params_mp_l[7] = &device_param->kernel_params_mp_l_buf32[7]; + device_param->kernel_params_mp_l[8] = &device_param->kernel_params_mp_l_buf32[8]; + device_param->kernel_params_mp_l[9] = &device_param->kernel_params_mp_l_buf32[9]; + + device_param->kernel_params_mp_r_buf64[3] = 0; + device_param->kernel_params_mp_r_buf32[4] = 0; + device_param->kernel_params_mp_r_buf32[5] = 0; + device_param->kernel_params_mp_r_buf32[6] = 0; + device_param->kernel_params_mp_r_buf32[7] = 0; + device_param->kernel_params_mp_r_buf32[8] = 0; + + device_param->kernel_params_mp_r[0] = NULL; + device_param->kernel_params_mp_r[1] = NULL; + device_param->kernel_params_mp_r[2] = NULL; + device_param->kernel_params_mp_r[3] = &device_param->kernel_params_mp_r_buf64[3]; + device_param->kernel_params_mp_r[4] = &device_param->kernel_params_mp_r_buf32[4]; + device_param->kernel_params_mp_r[5] = &device_param->kernel_params_mp_r_buf32[5]; + device_param->kernel_params_mp_r[6] = &device_param->kernel_params_mp_r_buf32[6]; + device_param->kernel_params_mp_r[7] = &device_param->kernel_params_mp_r_buf32[7]; + device_param->kernel_params_mp_r[8] = &device_param->kernel_params_mp_r_buf32[8]; + + device_param->kernel_params_amp_buf32[5] = 0; // combs_mode + device_param->kernel_params_amp_buf32[6] = 0; // gid_max + + device_param->kernel_params_amp[0] = &device_param->d_pws_buf; + device_param->kernel_params_amp[1] = &device_param->d_pws_amp_buf; + device_param->kernel_params_amp[2] = &device_param->d_rules_c; + device_param->kernel_params_amp[3] = &device_param->d_combs_c; + device_param->kernel_params_amp[4] = &device_param->d_bfs_c; + device_param->kernel_params_amp[5] = &device_param->kernel_params_amp_buf32[5]; + device_param->kernel_params_amp[6] = &device_param->kernel_params_amp_buf32[6]; + + device_param->kernel_params_tm[0] = &device_param->d_bfs_c; + device_param->kernel_params_tm[1] = &device_param->d_tm_c; + + device_param->kernel_params_memset_buf32[1] = 0; // value + device_param->kernel_params_memset_buf32[2] = 0; // gid_max + + device_param->kernel_params_memset[0] = NULL; + device_param->kernel_params_memset[1] = &device_param->kernel_params_memset_buf32[1]; + device_param->kernel_params_memset[2] = &device_param->kernel_params_memset_buf32[2]; + + /** + * kernel name + */ + + size_t kernel_wgs_tmp; + + char kernel_name[64] = { 0 }; + + if (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL) + { + if (hashconfig->opti_type & OPTI_TYPE_SINGLE_HASH) + { + snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_s%02d", hashconfig->kern_type, 4); + + CL_err = hc_clCreateKernel (opencl_ctx->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", hashconfig->kern_type, 8); + + CL_err = hc_clCreateKernel (opencl_ctx->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", hashconfig->kern_type, 16); + + CL_err = hc_clCreateKernel (opencl_ctx->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", hashconfig->kern_type, 4); + + CL_err = hc_clCreateKernel (opencl_ctx->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", hashconfig->kern_type, 8); + + CL_err = hc_clCreateKernel (opencl_ctx->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", hashconfig->kern_type, 16); + + CL_err = hc_clCreateKernel (opencl_ctx->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 (session_ctx->attack_mode == ATTACK_MODE_BF) + { + if (hashconfig->opts_type & OPTS_TYPE_PT_BITSLICE) + { + snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_tm", hashconfig->kern_type); + + CL_err = hc_clCreateKernel (opencl_ctx->ocl, device_param->program, kernel_name, &device_param->kernel_tm); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clCreateKernel(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } + + CL_err = hc_clGetKernelWorkGroupInfo (opencl_ctx->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; + } + } + } + } + else + { + snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_init", hashconfig->kern_type); + + CL_err = hc_clCreateKernel (opencl_ctx->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", hashconfig->kern_type); + + CL_err = hc_clCreateKernel (opencl_ctx->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", hashconfig->kern_type); + + CL_err = hc_clCreateKernel (opencl_ctx->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 (hashconfig->opts_type & OPTS_TYPE_HOOK12) + { + snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_hook12", hashconfig->kern_type); + + CL_err = hc_clCreateKernel (opencl_ctx->ocl, device_param->program, kernel_name, &device_param->kernel12); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clCreateKernel(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } + + CL_err = hc_clGetKernelWorkGroupInfo (opencl_ctx->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 (hashconfig->opts_type & OPTS_TYPE_HOOK23) + { + snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_hook23", hashconfig->kern_type); + + CL_err = hc_clCreateKernel (opencl_ctx->ocl, device_param->program, kernel_name, &device_param->kernel23); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clCreateKernel(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } + + CL_err = hc_clGetKernelWorkGroupInfo (opencl_ctx->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; + } + } + } + + CL_err |= hc_clGetKernelWorkGroupInfo (opencl_ctx->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 (opencl_ctx->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 (opencl_ctx->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++) + { + CL_err |= hc_clSetKernelArg (opencl_ctx->ocl, device_param->kernel1, i, sizeof (cl_mem), device_param->kernel_params[i]); + CL_err |= hc_clSetKernelArg (opencl_ctx->ocl, device_param->kernel2, i, sizeof (cl_mem), device_param->kernel_params[i]); + CL_err |= hc_clSetKernelArg (opencl_ctx->ocl, device_param->kernel3, i, sizeof (cl_mem), device_param->kernel_params[i]); + + if (hashconfig->opts_type & OPTS_TYPE_HOOK12) CL_err |= hc_clSetKernelArg (opencl_ctx->ocl, device_param->kernel12, i, sizeof (cl_mem), device_param->kernel_params[i]); + if (hashconfig->opts_type & OPTS_TYPE_HOOK23) CL_err |= hc_clSetKernelArg (opencl_ctx->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)); + + return -1; + } + } + + for (uint i = 24; i <= 34; i++) + { + CL_err |= hc_clSetKernelArg (opencl_ctx->ocl, device_param->kernel1, i, sizeof (cl_uint), device_param->kernel_params[i]); + CL_err |= hc_clSetKernelArg (opencl_ctx->ocl, device_param->kernel2, i, sizeof (cl_uint), device_param->kernel_params[i]); + CL_err |= hc_clSetKernelArg (opencl_ctx->ocl, device_param->kernel3, i, sizeof (cl_uint), device_param->kernel_params[i]); + + if (hashconfig->opts_type & OPTS_TYPE_HOOK12) CL_err |= hc_clSetKernelArg (opencl_ctx->ocl, device_param->kernel12, i, sizeof (cl_uint), device_param->kernel_params[i]); + if (hashconfig->opts_type & OPTS_TYPE_HOOK23) CL_err |= hc_clSetKernelArg (opencl_ctx->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 + + CL_err = hc_clCreateKernel (opencl_ctx->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 (opencl_ctx->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 (opencl_ctx->ocl, device_param->kernel_memset, 0, sizeof (cl_mem), device_param->kernel_params_memset[0]); + CL_err |= hc_clSetKernelArg (opencl_ctx->ocl, device_param->kernel_memset, 1, sizeof (cl_uint), device_param->kernel_params_memset[1]); + CL_err |= hc_clSetKernelArg (opencl_ctx->ocl, device_param->kernel_memset, 2, sizeof (cl_uint), device_param->kernel_params_memset[2]); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clSetKernelArg(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } + + // MP start + + if (session_ctx->attack_mode == ATTACK_MODE_BF) + { + CL_err |= hc_clCreateKernel (opencl_ctx->ocl, device_param->program_mp, "l_markov", &device_param->kernel_mp_l); + CL_err |= hc_clCreateKernel (opencl_ctx->ocl, device_param->program_mp, "r_markov", &device_param->kernel_mp_r); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clCreateKernel(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } + + CL_err |= hc_clGetKernelWorkGroupInfo (opencl_ctx->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 (opencl_ctx->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 (hashconfig->opts_type & OPTS_TYPE_PT_BITSLICE) + { + CL_err |= hc_clSetKernelArg (opencl_ctx->ocl, device_param->kernel_tm, 0, sizeof (cl_mem), device_param->kernel_params_tm[0]); + CL_err |= hc_clSetKernelArg (opencl_ctx->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 (session_ctx->attack_mode == ATTACK_MODE_HYBRID1) + { + CL_err = hc_clCreateKernel (opencl_ctx->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 (opencl_ctx->ocl, device_param->kernel_mp, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clGetKernelWorkGroupInfo(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } + } + else if (session_ctx->attack_mode == ATTACK_MODE_HYBRID2) + { + CL_err = hc_clCreateKernel (opencl_ctx->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 (opencl_ctx->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 (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL) + { + // nothing to do + } + else + { + CL_err = hc_clCreateKernel (opencl_ctx->ocl, device_param->program_amp, "amp", &device_param->kernel_amp); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: clCreateKernel(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } + + CL_err = hc_clGetKernelWorkGroupInfo (opencl_ctx->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 (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL) + { + // nothing to do + } + else + { + for (uint i = 0; i < 5; i++) + { + CL_err = hc_clSetKernelArg (opencl_ctx->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++) + { + CL_err = hc_clSetKernelArg (opencl_ctx->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; + } + } + } + + // maybe this has been updated by clGetKernelWorkGroupInfo() + // value can only be decreased, so we don't need to reallocate buffers + + device_param->kernel_threads = kernel_threads; + + // zero some data buffers + + run_kernel_bzero (opencl_ctx, device_param, device_param->d_pws_buf, size_pws); + run_kernel_bzero (opencl_ctx, device_param, device_param->d_pws_amp_buf, size_pws); + run_kernel_bzero (opencl_ctx, device_param, device_param->d_tmps, size_tmps); + run_kernel_bzero (opencl_ctx, device_param, device_param->d_hooks, size_hooks); + run_kernel_bzero (opencl_ctx, device_param, device_param->d_plain_bufs, size_plains); + run_kernel_bzero (opencl_ctx, device_param, device_param->d_result, size_results); + + /** + * special buffers + */ + + if (session_ctx->attack_kern == ATTACK_KERN_STRAIGHT) + { + run_kernel_bzero (opencl_ctx, device_param, device_param->d_rules_c, size_rules_c); + } + else if (session_ctx->attack_kern == ATTACK_KERN_COMBI) + { + run_kernel_bzero (opencl_ctx, device_param, device_param->d_combs, size_combs); + run_kernel_bzero (opencl_ctx, device_param, device_param->d_combs_c, size_combs); + run_kernel_bzero (opencl_ctx, device_param, device_param->d_root_css_buf, size_root_css); + run_kernel_bzero (opencl_ctx, device_param, device_param->d_markov_css_buf, size_markov_css); + } + else if (session_ctx->attack_kern == ATTACK_KERN_BF) + { + run_kernel_bzero (opencl_ctx, device_param, device_param->d_bfs, size_bfs); + run_kernel_bzero (opencl_ctx, device_param, device_param->d_bfs_c, size_bfs); + run_kernel_bzero (opencl_ctx, device_param, device_param->d_tm_c, size_tm); + run_kernel_bzero (opencl_ctx, device_param, device_param->d_root_css_buf, size_root_css); + run_kernel_bzero (opencl_ctx, device_param, device_param->d_markov_css_buf, size_markov_css); + } + } + + return 0; +} + +int opencl_session_destroy (opencl_ctx_t *opencl_ctx) +{ + for (uint device_id = 0; device_id < opencl_ctx->devices_cnt; device_id++) + { + hc_device_param_t *device_param = &opencl_ctx->devices_param[device_id]; + + if (device_param->skipped) continue; + + cl_int CL_err = CL_SUCCESS; + + myfree (device_param->pws_buf); + myfree (device_param->combs_buf); + myfree (device_param->hooks_buf); + + if (device_param->d_pws_buf) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_pws_buf); + if (device_param->d_pws_amp_buf) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_pws_amp_buf); + if (device_param->d_rules) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_rules); + if (device_param->d_rules_c) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_rules_c); + if (device_param->d_combs) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_combs); + if (device_param->d_combs_c) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_combs_c); + if (device_param->d_bfs) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_bfs); + if (device_param->d_bfs_c) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_bfs_c); + if (device_param->d_bitmap_s1_a) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_bitmap_s1_a); + if (device_param->d_bitmap_s1_b) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_bitmap_s1_b); + if (device_param->d_bitmap_s1_c) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_bitmap_s1_c); + if (device_param->d_bitmap_s1_d) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_bitmap_s1_d); + if (device_param->d_bitmap_s2_a) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_bitmap_s2_a); + if (device_param->d_bitmap_s2_b) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_bitmap_s2_b); + if (device_param->d_bitmap_s2_c) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_bitmap_s2_c); + if (device_param->d_bitmap_s2_d) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_bitmap_s2_d); + if (device_param->d_plain_bufs) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_plain_bufs); + if (device_param->d_digests_buf) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_digests_buf); + if (device_param->d_digests_shown) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_digests_shown); + if (device_param->d_salt_bufs) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_salt_bufs); + if (device_param->d_esalt_bufs) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_esalt_bufs); + if (device_param->d_tmps) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_tmps); + if (device_param->d_hooks) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_hooks); + if (device_param->d_result) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_result); + if (device_param->d_scryptV0_buf) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_scryptV0_buf); + if (device_param->d_scryptV1_buf) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_scryptV1_buf); + if (device_param->d_scryptV2_buf) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_scryptV2_buf); + if (device_param->d_scryptV3_buf) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_scryptV3_buf); + if (device_param->d_root_css_buf) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_root_css_buf); + if (device_param->d_markov_css_buf) CL_err |= hc_clReleaseMemObject (opencl_ctx->ocl, device_param->d_markov_css_buf); + if (device_param->d_tm_c) CL_err |= hc_clReleaseMemObject (opencl_ctx->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 (opencl_ctx->ocl, device_param->kernel1); + if (device_param->kernel12) CL_err |= hc_clReleaseKernel (opencl_ctx->ocl, device_param->kernel12); + if (device_param->kernel2) CL_err |= hc_clReleaseKernel (opencl_ctx->ocl, device_param->kernel2); + if (device_param->kernel23) CL_err |= hc_clReleaseKernel (opencl_ctx->ocl, device_param->kernel23); + if (device_param->kernel3) CL_err |= hc_clReleaseKernel (opencl_ctx->ocl, device_param->kernel3); + if (device_param->kernel_mp) CL_err |= hc_clReleaseKernel (opencl_ctx->ocl, device_param->kernel_mp); + if (device_param->kernel_mp_l) CL_err |= hc_clReleaseKernel (opencl_ctx->ocl, device_param->kernel_mp_l); + if (device_param->kernel_mp_r) CL_err |= hc_clReleaseKernel (opencl_ctx->ocl, device_param->kernel_mp_r); + if (device_param->kernel_tm) CL_err |= hc_clReleaseKernel (opencl_ctx->ocl, device_param->kernel_tm); + if (device_param->kernel_amp) CL_err |= hc_clReleaseKernel (opencl_ctx->ocl, device_param->kernel_amp); + if (device_param->kernel_memset) CL_err |= hc_clReleaseKernel (opencl_ctx->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 (opencl_ctx->ocl, device_param->program); + if (device_param->program_mp) CL_err |= hc_clReleaseProgram (opencl_ctx->ocl, device_param->program_mp); + if (device_param->program_amp) CL_err |= hc_clReleaseProgram (opencl_ctx->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 (opencl_ctx->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 (opencl_ctx->ocl, device_param->context); + + if (CL_err != CL_SUCCESS) + { + log_error ("ERROR: hc_clReleaseContext(): %s\n", val2cstr_cl (CL_err)); + + return -1; + } + } + + return 0; +} diff --git a/src/session.c b/src/session.c new file mode 100644 index 000000000..85e72ac6d --- /dev/null +++ b/src/session.c @@ -0,0 +1,94 @@ +/** + * Author......: See docs/credits.txt + * License.....: MIT + */ + +#include "common.h" +#include "types.h" +#include "session.h" + +void session_ctx_init (session_ctx_t *session_ctx, const bool quiet, const bool force, const bool benchmark, const u32 scrypt_tmto, char *cwd, char *install_dir, char *profile_dir, char *session_dir, char *shared_dir, char *cpath_real, const u32 wordlist_mode, char *rule_buf_l, char *rule_buf_r, const int rule_len_l, const int rule_len_r, const u32 kernel_rules_cnt, kernel_rule_t *kernel_rules_buf, const u32 attack_mode, const u32 attack_kern, const u32 bitmap_size, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, u32 *bitmap_s1_a, u32 *bitmap_s1_b, u32 *bitmap_s1_c, u32 *bitmap_s1_d, u32 *bitmap_s2_a, u32 *bitmap_s2_b, u32 *bitmap_s2_c, u32 *bitmap_s2_d) +{ + session_ctx->quiet = quiet; + session_ctx->force = force; + session_ctx->benchmark = benchmark; + + session_ctx->scrypt_tmto = scrypt_tmto; + + session_ctx->cwd = cwd; + session_ctx->install_dir = install_dir; + session_ctx->profile_dir = profile_dir; + session_ctx->session_dir = session_dir; + session_ctx->shared_dir = shared_dir; + session_ctx->cpath_real = cpath_real; + + session_ctx->wordlist_mode = wordlist_mode; + + session_ctx->rule_buf_l = rule_buf_l; + session_ctx->rule_buf_r = rule_buf_r; + session_ctx->rule_len_l = rule_len_l; + session_ctx->rule_len_r = rule_len_r; + + session_ctx->kernel_rules_cnt = kernel_rules_cnt; + session_ctx->kernel_rules_buf = kernel_rules_buf; + + session_ctx->attack_mode = attack_mode; + session_ctx->attack_kern = attack_kern; + + session_ctx->bitmap_size = bitmap_size; + session_ctx->bitmap_mask = bitmap_mask; + session_ctx->bitmap_shift1 = bitmap_shift1; + session_ctx->bitmap_shift2 = bitmap_shift2; + + session_ctx->bitmap_s1_a = bitmap_s1_a; + session_ctx->bitmap_s1_b = bitmap_s1_b; + session_ctx->bitmap_s1_c = bitmap_s1_c; + session_ctx->bitmap_s1_d = bitmap_s1_d; + session_ctx->bitmap_s2_a = bitmap_s2_a; + session_ctx->bitmap_s2_b = bitmap_s2_b; + session_ctx->bitmap_s2_c = bitmap_s2_c; + session_ctx->bitmap_s2_d = bitmap_s2_d; +} + +void session_ctx_destroy (session_ctx_t *session_ctx) +{ + session_ctx->quiet = false; + session_ctx->force = false; + session_ctx->benchmark = false; + + session_ctx->scrypt_tmto = 0; + + session_ctx->cwd = NULL; + session_ctx->install_dir = NULL; + session_ctx->profile_dir = NULL; + session_ctx->session_dir = NULL; + session_ctx->shared_dir = NULL; + session_ctx->cpath_real = NULL; + + session_ctx->wordlist_mode = 0; + + session_ctx->rule_buf_l = NULL; + session_ctx->rule_buf_r = NULL; + session_ctx->rule_len_l = 0; + session_ctx->rule_len_r = 0; + + session_ctx->kernel_rules_buf = NULL; + session_ctx->kernel_rules_cnt = 0; + + session_ctx->attack_mode = 0; + session_ctx->attack_kern = 0; + + session_ctx->bitmap_size = 0; + session_ctx->bitmap_mask = 0; + session_ctx->bitmap_shift1 = 0; + session_ctx->bitmap_shift2 = 0; + + session_ctx->bitmap_s1_a = NULL; + session_ctx->bitmap_s1_b = NULL; + session_ctx->bitmap_s1_c = NULL; + session_ctx->bitmap_s1_d = NULL; + session_ctx->bitmap_s2_a = NULL; + session_ctx->bitmap_s2_b = NULL; + session_ctx->bitmap_s2_c = NULL; + session_ctx->bitmap_s2_d = NULL; +}