diff --git a/include/shared.h b/include/shared.h index e22f66654..346aac75b 100644 --- a/include/shared.h +++ b/include/shared.h @@ -62,7 +62,7 @@ void hc_string_trim_trailing (char *s); void hc_string_trim_leading (char *s); size_t hc_fread (void *ptr, size_t size, size_t nmemb, FILE *stream); -void hc_fwrite (const void *ptr, size_t size, size_t nmemb, FILE *stream); +size_t hc_fwrite (const void *ptr, size_t size, size_t nmemb, FILE *stream); bool hc_same_files (char *file1, char *file2); diff --git a/src/opencl.c b/src/opencl.c index 75c07cdac..d6ea5c0e2 100644 --- a/src/opencl.c +++ b/src/opencl.c @@ -4375,6 +4375,8 @@ static u32 get_kernel_threads (hashcat_ctx_t *hashcat_ctx, const hc_device_param kernel_threads_max = device_maxworkgroup_size; } + u32 kernel_threads = kernel_threads_max; + // complicated kernel tend to confuse OpenCL runtime suggestions for maximum thread size // let's workaround that by sticking to their device specific preferred thread size @@ -4390,8 +4392,7 @@ static u32 get_kernel_threads (hashcat_ctx_t *hashcat_ctx, const hc_device_param if ((kernel_preferred_wgs_multiple1 >= kernel_threads_min) && (kernel_preferred_wgs_multiple1 <= kernel_threads_max)) { - kernel_threads_min = kernel_preferred_wgs_multiple1; - kernel_threads_max = kernel_preferred_wgs_multiple1; + kernel_threads = kernel_preferred_wgs_multiple1; } } } @@ -4403,8 +4404,7 @@ static u32 get_kernel_threads (hashcat_ctx_t *hashcat_ctx, const hc_device_param if ((kernel_preferred_wgs_multiple4 >= kernel_threads_min) && (kernel_preferred_wgs_multiple4 <= kernel_threads_max)) { - kernel_threads_min = kernel_preferred_wgs_multiple4; - kernel_threads_max = kernel_preferred_wgs_multiple4; + kernel_threads = kernel_preferred_wgs_multiple4; } } } @@ -4417,8 +4417,7 @@ static u32 get_kernel_threads (hashcat_ctx_t *hashcat_ctx, const hc_device_param if ((kernel_preferred_wgs_multiple2 >= kernel_threads_min) && (kernel_preferred_wgs_multiple2 <= kernel_threads_max)) { - kernel_threads_min = kernel_preferred_wgs_multiple2; - kernel_threads_max = kernel_preferred_wgs_multiple2; + kernel_threads = kernel_preferred_wgs_multiple2; } } } @@ -4435,8 +4434,7 @@ static u32 get_kernel_threads (hashcat_ctx_t *hashcat_ctx, const hc_device_param if ((kernel_preferred_wgs_multiple1 >= kernel_threads_min) && (kernel_preferred_wgs_multiple1 <= kernel_threads_max)) { - kernel_threads_min = kernel_preferred_wgs_multiple1; - kernel_threads_max = kernel_preferred_wgs_multiple1; + kernel_threads = kernel_preferred_wgs_multiple1; } } } @@ -4448,8 +4446,7 @@ static u32 get_kernel_threads (hashcat_ctx_t *hashcat_ctx, const hc_device_param if ((kernel_preferred_wgs_multiple4 >= kernel_threads_min) && (kernel_preferred_wgs_multiple4 <= kernel_threads_max)) { - kernel_threads_min = kernel_preferred_wgs_multiple4; - kernel_threads_max = kernel_preferred_wgs_multiple4; + kernel_threads = kernel_preferred_wgs_multiple4; } } } @@ -4462,14 +4459,13 @@ static u32 get_kernel_threads (hashcat_ctx_t *hashcat_ctx, const hc_device_param if ((kernel_preferred_wgs_multiple2 >= kernel_threads_min) && (kernel_preferred_wgs_multiple2 <= kernel_threads_max)) { - kernel_threads_min = kernel_preferred_wgs_multiple2; - kernel_threads_max = kernel_preferred_wgs_multiple2; + kernel_threads = kernel_preferred_wgs_multiple2; } } } } - return kernel_threads_max; + return kernel_threads; } int opencl_session_begin (hashcat_ctx_t *hashcat_ctx) @@ -4858,6 +4854,8 @@ int opencl_session_begin (hashcat_ctx_t *hashcat_ctx) build_options_len += snprintf (build_options_buf + build_options_len, build_options_sz - build_options_len, "-D LOCAL_MEM_TYPE=%u -D VENDOR_ID=%u -D CUDA_ARCH=%u -D HAS_VPERM=%u -D HAS_VADD3=%u -D HAS_VBFE=%u -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 -w ", device_param->device_local_mem_type, device_param->platform_vendor_id, (device_param->sm_major * 100) + device_param->sm_minor, device_param->has_vperm, device_param->has_vadd3, device_param->has_vbfe, 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, kern_type); #endif + build_options_buf[build_options_len] = 0; + /* if (device_param->device_type & CL_DEVICE_TYPE_CPU) { @@ -4884,6 +4882,8 @@ int opencl_session_begin (hashcat_ctx_t *hashcat_ctx) } } + build_options_module_buf[build_options_module_len] = 0; + #if defined (DEBUG) if (user_options->quiet == false) event_log_warning (hashcat_ctx, "* Device #%u: build_options '%s'", device_id + 1, build_options_buf); if (user_options->quiet == false) event_log_warning (hashcat_ctx, "* Device #%u: build_options_module '%s'", device_id + 1, build_options_module_buf); diff --git a/src/outfile_check.c b/src/outfile_check.c index 8a4ec38c7..28fa4ff11 100644 --- a/src/outfile_check.c +++ b/src/outfile_check.c @@ -209,17 +209,17 @@ static int outfile_remove (hashcat_ctx_t *hashcat_ctx) if (line_hash_len == 0) continue; - if (hashconfig->is_salted == true) + if (hash_buf.salt) { memset (hash_buf.salt, 0, sizeof (salt_t)); } - if (hashconfig->esalt_size > 0) + if (hash_buf.esalt) { memset (hash_buf.esalt, 0, hashconfig->esalt_size); } - if (hashconfig->hook_salt_size > 0) + if (hash_buf.hook_salt) { memset (hash_buf.hook_salt, 0, hashconfig->hook_salt_size); } diff --git a/src/potfile.c b/src/potfile.c index c6ff47273..2e5262a86 100644 --- a/src/potfile.c +++ b/src/potfile.c @@ -316,9 +316,12 @@ void potfile_update_hash (hashcat_ctx_t *hashcat_ctx, hash_t *found, char *line_ found->pw_buf = (char *) hcmalloc (pw_len + 1); found->pw_len = pw_len; - memcpy (found->pw_buf, pw_buf, pw_len); + if (pw_buf) + { + memcpy (found->pw_buf, pw_buf, pw_len); - found->pw_buf[found->pw_len] = 0; + found->pw_buf[found->pw_len] = 0; + } found->cracked = 1; @@ -546,17 +549,17 @@ int potfile_remove_parse (hashcat_ctx_t *hashcat_ctx) if (line_hash_len == 0) continue; - if (hashconfig->is_salted == true) + if (hash_buf.salt) { memset (hash_buf.salt, 0, sizeof (salt_t)); } - if (hashconfig->esalt_size > 0) + if (hash_buf.esalt) { memset (hash_buf.esalt, 0, hashconfig->esalt_size); } - if (hashconfig->hook_salt_size > 0) + if (hash_buf.hook_salt) { memset (hash_buf.hook_salt, 0, hashconfig->hook_salt_size); } diff --git a/src/shared.c b/src/shared.c index d034d87dc..930e93024 100644 --- a/src/shared.c +++ b/src/shared.c @@ -592,11 +592,9 @@ size_t hc_fread (void *ptr, size_t size, size_t nmemb, FILE *stream) return fread (ptr, size, nmemb, stream); } -void hc_fwrite (const void *ptr, size_t size, size_t nmemb, FILE *stream) +size_t hc_fwrite (const void *ptr, size_t size, size_t nmemb, FILE *stream) { - size_t rc = fwrite (ptr, size, nmemb, stream); - - if (rc == 0) rc = 0; + return fwrite (ptr, size, nmemb, stream); } bool hc_same_files (char *file1, char *file2)