diff --git a/OpenCL/inc_types.h b/OpenCL/inc_types.h index 233e28958..a13c89b8e 100644 --- a/OpenCL/inc_types.h +++ b/OpenCL/inc_types.h @@ -2008,6 +2008,7 @@ typedef struct salt u32 salt_len_pc; u32 salt_iter; u32 salt_iter2; + u32 salt_dimy; u32 salt_sign[2]; u32 salt_repeats; diff --git a/Python/hcshared.py b/Python/hcshared.py index 2fcf56c70..ad4390083 100644 --- a/Python/hcshared.py +++ b/Python/hcshared.py @@ -6,13 +6,14 @@ import sys def extract_salts(salts_buf) -> list: salts=[] - for salt_buf, salt_buf_pc, salt_len, salt_len_pc, salt_iter, salt_iter2, salt_sign, salt_repeats, orig_pos, digests_cnt, digests_done, digests_offset, scrypt_N, scrypt_r, scrypt_p in struct.iter_unpack("256s 256s I I I I 8s I I I I I I I I", salts_buf): + for salt_buf, salt_buf_pc, salt_len, salt_len_pc, salt_iter, salt_iter2, salt_dimy, salt_sign, salt_repeats, orig_pos, digests_cnt, digests_done, digests_offset, scrypt_N, scrypt_r, scrypt_p in struct.iter_unpack("256s 256s I I I I I 8s I I I I I I I I", salts_buf): salt_buf = salt_buf[0:salt_len] salt_buf_pc = salt_buf_pc[0:salt_len_pc] salts.append({ "salt_buf": salt_buf, \ "salt_buf_pc": salt_buf_pc, \ "salt_iter": salt_iter, \ "salt_iter2": salt_iter2, \ + "salt_dimy": salt_dimy, \ "salt_sign": salt_sign, \ "salt_repeats": salt_repeats, \ "orig_pos": orig_pos, \ diff --git a/docs/changes.txt b/docs/changes.txt index 96cbd3317..ca97518aa 100644 --- a/docs/changes.txt +++ b/docs/changes.txt @@ -92,6 +92,7 @@ - Fixed bug in grep out-of-memory workaround on Unit Test - Fixed bug in input_tokenizer when TOKEN_ATTR_FIXED_LENGTH is used and refactor modules - Fixed bug in --stdout that caused certain rules to malfunction +- Fixed race condition in selftest_init on OpenCL with non-blocking write - Fixed build failed for 10700 optimized with Apple Metal - Fixed build failed for 13772 and 13773 with Apple Metal - Fixed build failed for 18400 with Apple Metal @@ -128,6 +129,7 @@ - Apple Driver: Updated requirements to use Apple OpenCL API to macOS 13.0 - use - Backend: Updated filename chksum format to prevent invalid cache on Apple Silicon when switching arch - Backend Checks: Describe workaround in error message when detecting more than 64 backend devices +- Backend Info: Added --machine-readable format - Brain: Added sanity check and corresponding error message for invalid --brain-port values - Dependencies: Added sse2neon v1.8.0 (commit 658eeac) - Dependencies: Updated LZMA SDK to 24.09 @@ -140,6 +142,7 @@ - Modules: Updated module_unstable_warning - Open Document Format: Added support for small documents with content length < 1024 - OpenCL Backend: added workaround to set device_available_memory from CUDA/HIP alias device +- Selftest: rename selftest function to process_selftest and splitting into 3 smaller functions - Status Code: Add specific return code for self-test fail (-11) - Scrypt: Increase buffer sizes in module for hash mode 8900 to allow longer scrypt digests - Unicode: Update UTF-8 to UTF-16 conversion to match RFC 3629 diff --git a/include/ext_OpenCL.h b/include/ext_OpenCL.h index d49f82c79..7044ac667 100644 --- a/include/ext_OpenCL.h +++ b/include/ext_OpenCL.h @@ -51,7 +51,7 @@ typedef cl_int (CL_API_CALL *OCL_CLENQUEUEFILLBUFFER) (cl_comman typedef cl_int (CL_API_CALL *OCL_CLENQUEUECOPYBUFFER) (cl_command_queue, cl_mem, cl_mem, size_t, size_t, size_t, cl_uint, const cl_event *, cl_event *); typedef void * (CL_API_CALL *OCL_CLENQUEUEMAPBUFFER) (cl_command_queue, cl_mem, cl_bool, cl_map_flags, size_t, size_t, cl_uint, const cl_event *, cl_event *, cl_int *); typedef cl_int (CL_API_CALL *OCL_CLENQUEUENDRANGEKERNEL) (cl_command_queue, cl_kernel, cl_uint, const size_t *, const size_t *, const size_t *, cl_uint, const cl_event *, cl_event *); -typedef cl_int (CL_API_CALL *OCL_CLENQUEUEREADBUFFER) (cl_command_queue, cl_mem, cl_bool, size_t, size_t, const void *, cl_uint, const cl_event *, cl_event *); +typedef cl_int (CL_API_CALL *OCL_CLENQUEUEREADBUFFER) (cl_command_queue, cl_mem, cl_bool, size_t, size_t, void *, cl_uint, const cl_event *, cl_event *); typedef cl_int (CL_API_CALL *OCL_CLENQUEUEUNMAPMEMOBJECT) (cl_command_queue, cl_mem, void *, cl_uint, const cl_event *, cl_event *); typedef cl_int (CL_API_CALL *OCL_CLENQUEUEWRITEBUFFER) (cl_command_queue, cl_mem, cl_bool, size_t, size_t, const void *, cl_uint, const cl_event *, cl_event *); typedef cl_int (CL_API_CALL *OCL_CLFINISH) (cl_command_queue); diff --git a/include/ext_cuda.h b/include/ext_cuda.h index 44ae01cd6..a80af5367 100644 --- a/include/ext_cuda.h +++ b/include/ext_cuda.h @@ -1154,14 +1154,19 @@ typedef CUresult (CUDA_API_CALL *CUDA_CUINIT) (unsigned int); typedef CUresult (CUDA_API_CALL *CUDA_CULAUNCHKERNEL) (CUfunction, unsigned int, unsigned int, unsigned int, unsigned int, unsigned int, unsigned int, unsigned int, CUstream, void **, void **); typedef CUresult (CUDA_API_CALL *CUDA_CUMEMALLOC) (CUdeviceptr *, size_t); typedef CUresult (CUDA_API_CALL *CUDA_CUMEMALLOCHOST) (void **, size_t); +typedef CUresult (CUDA_API_CALL *CUDA_CUMEMCPYDTOD) (CUdeviceptr, CUdeviceptr, size_t); +typedef CUresult (CUDA_API_CALL *CUDA_CUMEMCPYDTOH) (void *, CUdeviceptr, size_t); +typedef CUresult (CUDA_API_CALL *CUDA_CUMEMCPYHTOD) (CUdeviceptr, const void *, size_t); +typedef CUresult (CUDA_API_CALL *CUDA_CUMEMSETD32) (CUdeviceptr, unsigned int, size_t); +typedef CUresult (CUDA_API_CALL *CUDA_CUMEMSETD8) (CUdeviceptr, unsigned char, size_t); typedef CUresult (CUDA_API_CALL *CUDA_CUMEMCPYDTODASYNC) (CUdeviceptr, CUdeviceptr, size_t, CUstream); typedef CUresult (CUDA_API_CALL *CUDA_CUMEMCPYDTOHASYNC) (void *, CUdeviceptr, size_t, CUstream); typedef CUresult (CUDA_API_CALL *CUDA_CUMEMCPYHTODASYNC) (CUdeviceptr, const void *, size_t, CUstream); +typedef CUresult (CUDA_API_CALL *CUDA_CUMEMSETD32ASYNC) (CUdeviceptr, unsigned int, size_t, CUstream); +typedef CUresult (CUDA_API_CALL *CUDA_CUMEMSETD8ASYNC) (CUdeviceptr, unsigned char, size_t, CUstream); typedef CUresult (CUDA_API_CALL *CUDA_CUMEMFREE) (CUdeviceptr); typedef CUresult (CUDA_API_CALL *CUDA_CUMEMFREEHOST) (void *); typedef CUresult (CUDA_API_CALL *CUDA_CUMEMGETINFO) (size_t *, size_t *); -typedef CUresult (CUDA_API_CALL *CUDA_CUMEMSETD32ASYNC) (CUdeviceptr, unsigned int, size_t, CUstream); -typedef CUresult (CUDA_API_CALL *CUDA_CUMEMSETD8ASYNC) (CUdeviceptr, unsigned char, size_t, CUstream); typedef CUresult (CUDA_API_CALL *CUDA_CUMODULEGETFUNCTION) (CUfunction *, CUmodule, const char *); typedef CUresult (CUDA_API_CALL *CUDA_CUMODULEGETGLOBAL) (CUdeviceptr *, size_t *, CUmodule, const char *); typedef CUresult (CUDA_API_CALL *CUDA_CUMODULELOAD) (CUmodule *, const char *); @@ -1217,14 +1222,19 @@ typedef struct hc_cuda_lib CUDA_CULAUNCHKERNEL cuLaunchKernel; CUDA_CUMEMALLOC cuMemAlloc; CUDA_CUMEMALLOCHOST cuMemAllocHost; + CUDA_CUMEMCPYDTOD cuMemcpyDtoD; + CUDA_CUMEMCPYDTOH cuMemcpyDtoH; + CUDA_CUMEMCPYHTOD cuMemcpyHtoD; + CUDA_CUMEMSETD32 cuMemsetD32; + CUDA_CUMEMSETD8 cuMemsetD8; CUDA_CUMEMCPYDTODASYNC cuMemcpyDtoDAsync; CUDA_CUMEMCPYDTOHASYNC cuMemcpyDtoHAsync; CUDA_CUMEMCPYHTODASYNC cuMemcpyHtoDAsync; + CUDA_CUMEMSETD32ASYNC cuMemsetD32Async; + CUDA_CUMEMSETD8ASYNC cuMemsetD8Async; CUDA_CUMEMFREE cuMemFree; CUDA_CUMEMFREEHOST cuMemFreeHost; CUDA_CUMEMGETINFO cuMemGetInfo; - CUDA_CUMEMSETD32ASYNC cuMemsetD32Async; - CUDA_CUMEMSETD8ASYNC cuMemsetD8Async; CUDA_CUMODULEGETFUNCTION cuModuleGetFunction; CUDA_CUMODULEGETGLOBAL cuModuleGetGlobal; CUDA_CUMODULELOAD cuModuleLoad; @@ -1272,13 +1282,18 @@ int hc_cuFuncSetAttribute (void *hashcat_ctx, CUfunction hfunc, CUfunction_ int hc_cuInit (void *hashcat_ctx, unsigned int Flags); int hc_cuLaunchKernel (void *hashcat_ctx, CUfunction f, unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ, unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ, unsigned int sharedMemBytes, CUstream hStream, void **kernelParams, void **extra); int hc_cuMemAlloc (void *hashcat_ctx, CUdeviceptr *dptr, size_t bytesize); +int hc_cuMemcpyDtoD (void *hashcat_ctx, CUdeviceptr dstDevice, CUdeviceptr srcDevice, size_t ByteCount); +int hc_cuMemcpyDtoH (void *hashcat_ctx, void *dstHost, CUdeviceptr srcDevice, size_t ByteCount); +int hc_cuMemcpyHtoD (void *hashcat_ctx, CUdeviceptr dstDevice, const void *srcHost, size_t ByteCount); +int hc_cuMemsetD32 (void *hashcat_ctx, CUdeviceptr dstDevice, unsigned int ui, size_t N); +int hc_cuMemsetD8 (void *hashcat_ctx, CUdeviceptr dstDevice, unsigned char uc, size_t N); int hc_cuMemcpyDtoDAsync (void *hashcat_ctx, CUdeviceptr dstDevice, CUdeviceptr srcDevice, size_t ByteCount, CUstream hStream); int hc_cuMemcpyDtoHAsync (void *hashcat_ctx, void *dstHost, CUdeviceptr srcDevice, size_t ByteCount, CUstream hStream); int hc_cuMemcpyHtoDAsync (void *hashcat_ctx, CUdeviceptr dstDevice, const void *srcHost, size_t ByteCount, CUstream hStream); -int hc_cuMemFree (void *hashcat_ctx, CUdeviceptr dptr); -int hc_cuMemGetInfo (void *hashcat_ctx, size_t *free, size_t *total); int hc_cuMemsetD32Async (void *hashcat_ctx, CUdeviceptr dstDevice, unsigned int ui, size_t N, CUstream hStream); int hc_cuMemsetD8Async (void *hashcat_ctx, CUdeviceptr dstDevice, unsigned char uc, size_t N, CUstream hStream); +int hc_cuMemFree (void *hashcat_ctx, CUdeviceptr dptr); +int hc_cuMemGetInfo (void *hashcat_ctx, size_t *free, size_t *total); int hc_cuModuleGetFunction (void *hashcat_ctx, CUfunction *hfunc, CUmodule hmod, const char *name); int hc_cuModuleGetGlobal (void *hashcat_ctx, CUdeviceptr *dptr, size_t *bytes, CUmodule hmod, const char *name); int hc_cuModuleLoadDataEx (void *hashcat_ctx, CUmodule *module, const void *image, unsigned int numOptions, CUjit_option *options, void **optionValues); diff --git a/include/ext_hip.h b/include/ext_hip.h index 4cc6fb797..d0f53d173 100644 --- a/include/ext_hip.h +++ b/include/ext_hip.h @@ -588,6 +588,11 @@ typedef hipError_t (HIP_API_CALL *HIP_HIPLAUNCHKERNEL) (hipFunction_t, typedef hipError_t (HIP_API_CALL *HIP_HIPMEMALLOC) (hipDeviceptr_t *, size_t); typedef hipError_t (HIP_API_CALL *HIP_HIPMEMFREE) (hipDeviceptr_t); typedef hipError_t (HIP_API_CALL *HIP_HIPMEMGETINFO) (size_t *, size_t *); +typedef hipError_t (HIP_API_CALL *HIP_HIPMEMCPYDTOD) (hipDeviceptr_t, hipDeviceptr_t, size_t); +typedef hipError_t (HIP_API_CALL *HIP_HIPMEMCPYDTOH) (void *, hipDeviceptr_t, size_t); +typedef hipError_t (HIP_API_CALL *HIP_HIPMEMCPYHTOD) (hipDeviceptr_t, const void *, size_t); +typedef hipError_t (HIP_API_CALL *HIP_HIPMEMSETD32) (hipDeviceptr_t, unsigned int, size_t); +typedef hipError_t (HIP_API_CALL *HIP_HIPMEMSETD8) (hipDeviceptr_t, unsigned char, size_t); typedef hipError_t (HIP_API_CALL *HIP_HIPMEMCPYDTODASYNC) (hipDeviceptr_t, hipDeviceptr_t, size_t, hipStream_t); typedef hipError_t (HIP_API_CALL *HIP_HIPMEMCPYDTOHASYNC) (void *, hipDeviceptr_t, size_t, hipStream_t); typedef hipError_t (HIP_API_CALL *HIP_HIPMEMCPYHTODASYNC) (hipDeviceptr_t, const void *, size_t, hipStream_t); @@ -633,6 +638,11 @@ typedef struct hc_hip_lib HIP_HIPMEMALLOC hipMemAlloc; HIP_HIPMEMFREE hipMemFree; HIP_HIPMEMGETINFO hipMemGetInfo; + HIP_HIPMEMCPYDTOD hipMemcpyDtoD; + HIP_HIPMEMCPYDTOH hipMemcpyDtoH; + HIP_HIPMEMCPYHTOD hipMemcpyHtoD; + HIP_HIPMEMSETD32 hipMemsetD32; + HIP_HIPMEMSETD8 hipMemsetD8; HIP_HIPMEMCPYDTODASYNC hipMemcpyDtoDAsync; HIP_HIPMEMCPYDTOHASYNC hipMemcpyDtoHAsync; HIP_HIPMEMCPYHTODASYNC hipMemcpyHtoDAsync; @@ -680,6 +690,11 @@ int hc_hipLaunchKernel (void *hashcat_ctx, hipFunction_t f, unsigned int int hc_hipMemAlloc (void *hashcat_ctx, hipDeviceptr_t *dptr, size_t bytesize); int hc_hipMemFree (void *hashcat_ctx, hipDeviceptr_t dptr); int hc_hipMemGetInfo (void *hashcat_ctx, size_t *free, size_t *total); +int hc_hipMemcpyDtoD (void *hashcat_ctx, hipDeviceptr_t dstDevice, hipDeviceptr_t srcDevice, size_t ByteCount); +int hc_hipMemcpyDtoH (void *hashcat_ctx, void *dstHost, hipDeviceptr_t srcDevice, size_t ByteCount); +int hc_hipMemcpyHtoD (void *hashcat_ctx, hipDeviceptr_t dstDevice, const void *srcHost, size_t ByteCount); +int hc_hipMemsetD32 (void *hashcat_ctx, hipDeviceptr_t dstDevice, unsigned int ui, size_t N); +int hc_hipMemsetD8 (void *hashcat_ctx, hipDeviceptr_t dstDevice, unsigned char uc, size_t N); int hc_hipMemcpyDtoDAsync (void *hashcat_ctx, hipDeviceptr_t dstDevice, hipDeviceptr_t srcDevice, size_t ByteCount, hipStream_t hStream); int hc_hipMemcpyDtoHAsync (void *hashcat_ctx, void *dstHost, hipDeviceptr_t srcDevice, size_t ByteCount, hipStream_t hStream); int hc_hipMemcpyHtoDAsync (void *hashcat_ctx, hipDeviceptr_t dstDevice, const void *srcHost, size_t ByteCount, hipStream_t hStream); diff --git a/include/types.h b/include/types.h index a62c21e27..de74243af 100644 --- a/include/types.h +++ b/include/types.h @@ -412,6 +412,9 @@ typedef enum opti_type OPTI_TYPE_REGISTER_LIMIT = (1 << 20), // We'll limit the register count to 128 OPTI_TYPE_SLOW_HASH_SIMD_INIT2 = (1 << 21), OPTI_TYPE_SLOW_HASH_SIMD_LOOP2 = (1 << 22), + OPTI_TYPE_SLOW_HASH_DIMY_INIT = (1 << 23), + OPTI_TYPE_SLOW_HASH_DIMY_LOOP = (1 << 24), + OPTI_TYPE_SLOW_HASH_DIMY_COMP = (1 << 25), } opti_type_t; @@ -476,14 +479,17 @@ typedef enum opts_type OPTS_TYPE_DYNAMIC_SHARED = (1ULL << 53), // use dynamic shared memory (note: needs special kernel changes) OPTS_TYPE_SELF_TEST_DISABLE = (1ULL << 54), // some algos use JiT in combinations with a salt or create too much startup time OPTS_TYPE_MP_MULTI_DISABLE = (1ULL << 55), // do not multiply the kernel-accel with the multiprocessor count per device to allow more fine-tuned workload settings - OPTS_TYPE_NATIVE_THREADS = (1ULL << 56), // forces "native" thread count: CPU=1, GPU-Intel=8, GPU-AMD=64 (wavefront), GPU-NV=32 (warps) - OPTS_TYPE_MAXIMUM_THREADS = (1ULL << 57), // disable else branch in pre-compilation thread count optimization setting - OPTS_TYPE_POST_AMP_UTF16LE = (1ULL << 58), // run the utf8 to utf16le conversion kernel after they have been processed from amplifiers + OPTS_TYPE_THREAD_MULTI_DISABLE // do not multiply the kernel-power with the thread count per device for super slow algos + = (1ULL << 56), + OPTS_TYPE_NATIVE_THREADS = (1ULL << 57), // forces "native" thread count: CPU=1, GPU-Intel=8, GPU-AMD=64 (wavefront), GPU-NV=32 (warps) + OPTS_TYPE_MAXIMUM_THREADS = (1ULL << 58), // disable else branch in pre-compilation thread count optimization setting + OPTS_TYPE_POST_AMP_UTF16LE = (1ULL << 59), // run the utf8 to utf16le conversion kernel after they have been processed from amplifiers OPTS_TYPE_AUTODETECT_DISABLE - = (1ULL << 59), // skip autodetect engine - OPTS_TYPE_STOCK_MODULE = (1ULL << 60), // module included with hashcat default distribution + = (1ULL << 60), // skip autodetect engine + OPTS_TYPE_STOCK_MODULE = (1ULL << 61), // module included with hashcat default distribution OPTS_TYPE_MULTIHASH_DESPITE_ESALT - = (1ULL << 61), // overrule multihash cracking check same salt but not same esalt + = (1ULL << 62), // overrule multihash cracking check same salt but not same esalt + OPTS_TYPE_MAXIMUM_ACCEL = (1ULL << 63) // try to maximize kernel-accel during autotune } opts_type_t; diff --git a/src/autotune.c b/src/autotune.c index 87637b29d..a599e65be 100644 --- a/src/autotune.c +++ b/src/autotune.c @@ -43,7 +43,8 @@ static double try_run (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_par device_param->kernel_param.loop_cnt = kernel_loops; // not a bug, both need to be set device_param->kernel_param.il_cnt = kernel_loops; // because there's two variables for inner iters for slow and fast hashes - const u32 hardware_power = ((hashconfig->opts_type & OPTS_TYPE_MP_MULTI_DISABLE) ? 1 : device_param->device_processors) * kernel_threads; + const u32 hardware_power = ((hashconfig->opts_type & OPTS_TYPE_MP_MULTI_DISABLE) ? 1 : device_param->device_processors) + * ((hashconfig->opts_type & OPTS_TYPE_THREAD_MULTI_DISABLE) ? 1 : kernel_threads); u32 kernel_power_try = hardware_power * kernel_accel; @@ -133,7 +134,8 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param device_param->kernel_accel = kernel_accel_min; device_param->kernel_loops = kernel_loops_min; device_param->kernel_threads = kernel_threads_min; - device_param->hardware_power = ((hashconfig->opts_type & OPTS_TYPE_MP_MULTI_DISABLE) ? 1 : device_param->device_processors) * kernel_threads_min; + device_param->hardware_power = ((hashconfig->opts_type & OPTS_TYPE_MP_MULTI_DISABLE) ? 1 : device_param->device_processors) + * ((hashconfig->opts_type & OPTS_TYPE_THREAD_MULTI_DISABLE) ? 1 : kernel_threads_min); device_param->kernel_power = device_param->hardware_power * kernel_accel_min; } @@ -212,7 +214,8 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param // from here it's clear we are allowed to autotune // so let's init some fake words - const u32 hardware_power_max = ((hashconfig->opts_type & OPTS_TYPE_MP_MULTI_DISABLE) ? 1 : device_param->device_processors) * kernel_threads_max; + const u32 hardware_power_max = ((hashconfig->opts_type & OPTS_TYPE_MP_MULTI_DISABLE) ? 1 : device_param->device_processors) + * ((hashconfig->opts_type & OPTS_TYPE_THREAD_MULTI_DISABLE) ? 1 : kernel_threads_max); u32 kernel_power_max = hardware_power_max * kernel_accel_max; @@ -265,12 +268,12 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (device_param->is_cuda == true) { - if (hc_cuMemcpyDtoDAsync (hashcat_ctx, device_param->cuda_d_rules_c, device_param->cuda_d_rules, MIN (kernel_loops_max, KERNEL_RULES) * sizeof (kernel_rule_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyDtoD (hashcat_ctx, device_param->cuda_d_rules_c, device_param->cuda_d_rules, MIN (kernel_loops_max, KERNEL_RULES) * sizeof (kernel_rule_t)) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyDtoDAsync (hashcat_ctx, device_param->hip_d_rules_c, device_param->hip_d_rules, MIN (kernel_loops_max, KERNEL_RULES) * sizeof (kernel_rule_t), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyDtoD (hashcat_ctx, device_param->hip_d_rules_c, device_param->hip_d_rules, MIN (kernel_loops_max, KERNEL_RULES) * sizeof (kernel_rule_t)) == -1) return -1; } #if defined (__APPLE__) @@ -298,13 +301,13 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param { const u32 kernel_threads_sav = device_param->kernel_threads; - device_param->kernel_threads = device_param->kernel_wgs1; + device_param->kernel_threads = MIN (device_param->kernel_wgs1, kernel_threads_max); run_kernel (hashcat_ctx, device_param, KERN_RUN_1, 0, kernel_power_max, false, 0, true); if (hashconfig->opts_type & OPTS_TYPE_LOOP_PREPARE) { - device_param->kernel_threads = device_param->kernel_wgs2p; + device_param->kernel_threads = MIN (device_param->kernel_wgs2p, kernel_threads_max); run_kernel (hashcat_ctx, device_param, KERN_RUN_2P, 0, kernel_power_max, false, 0, true); } @@ -330,8 +333,6 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param // v7 autotuner is a lot more straight forward - u32 kernel_loops_min_start = kernel_loops_min; - if (hashes && hashes->st_salts_buf) { u32 start = kernel_loops_max; @@ -343,17 +344,17 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param start = MIN (start, smallest_repeat_double (hashes->st_salts_buf->salt_iter)); start = MIN (start, smallest_repeat_double (hashes->st_salts_buf->salt_iter + 1)); - if ((hashes->st_salts_buf->salt_iter % 125) == 0) start = MIN (start, 125); - if ((hashes->st_salts_buf->salt_iter + 1 % 125) == 0) start = MIN (start, 125); + if (((hashes->st_salts_buf->salt_iter + 0) % 125) == 0) start = MIN (start, 125); + if (((hashes->st_salts_buf->salt_iter + 1) % 125) == 0) start = MIN (start, 125); if ((start >= kernel_loops_min) && (start <= kernel_loops_max)) { - kernel_loops_min_start = start; + kernel_loops = start; } } } - for (u32 kernel_loops_test = kernel_loops_min_start; kernel_loops_test <= kernel_loops_max; kernel_loops_test <<= 1) + for (u32 kernel_loops_test = kernel_loops; kernel_loops_test <= kernel_loops_max; kernel_loops_test <<= 1) { double exec_msec = try_run_times (hashcat_ctx, device_param, kernel_accel_min, kernel_loops_test, kernel_threads_min, 2); @@ -564,7 +565,8 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param device_param->kernel_loops = kernel_loops; device_param->kernel_threads = kernel_threads; - const u32 hardware_power = ((hashconfig->opts_type & OPTS_TYPE_MP_MULTI_DISABLE) ? 1 : device_param->device_processors) * device_param->kernel_threads; + const u32 hardware_power = ((hashconfig->opts_type & OPTS_TYPE_MP_MULTI_DISABLE) ? 1 : device_param->device_processors) + * ((hashconfig->opts_type & OPTS_TYPE_THREAD_MULTI_DISABLE) ? 1 : device_param->kernel_threads); device_param->hardware_power = hardware_power; diff --git a/src/backend.c b/src/backend.c index bd8a82220..8761ab0d2 100644 --- a/src/backend.c +++ b/src/backend.c @@ -985,7 +985,7 @@ int gidd_to_pw_t (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, c { if (hc_cuCtxPushCurrent (hashcat_ctx, device_param->cuda_context) == -1) return -1; - if (hc_cuMemcpyDtoHAsync (hashcat_ctx, &pw_idx, device_param->cuda_d_pws_idx + (gidd * sizeof (pw_idx_t)), sizeof (pw_idx_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyDtoH (hashcat_ctx, &pw_idx, device_param->cuda_d_pws_idx + (gidd * sizeof (pw_idx_t)), sizeof (pw_idx_t)) == -1) return -1; if (hc_cuStreamSynchronize (hashcat_ctx, device_param->cuda_stream) == -1) return -1; } @@ -994,7 +994,7 @@ int gidd_to_pw_t (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, c { if (hc_hipCtxPushCurrent (hashcat_ctx, device_param->hip_context) == -1) return -1; - if (hc_hipMemcpyDtoHAsync (hashcat_ctx, &pw_idx, device_param->hip_d_pws_idx + (gidd * sizeof (pw_idx_t)), sizeof (pw_idx_t), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyDtoH (hashcat_ctx, &pw_idx, device_param->hip_d_pws_idx + (gidd * sizeof (pw_idx_t)), sizeof (pw_idx_t)) == -1) return -1; if (hc_hipStreamSynchronize (hashcat_ctx, device_param->hip_stream) == -1) return -1; } @@ -1020,14 +1020,14 @@ int gidd_to_pw_t (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, c { if (device_param->is_cuda == true) { - if (hc_cuMemcpyDtoHAsync (hashcat_ctx, pw->i, device_param->cuda_d_pws_comp_buf + (off * sizeof (u32)), cnt * sizeof (u32), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyDtoH (hashcat_ctx, pw->i, device_param->cuda_d_pws_comp_buf + (off * sizeof (u32)), cnt * sizeof (u32)) == -1) return -1; if (hc_cuStreamSynchronize (hashcat_ctx, device_param->cuda_stream) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyDtoHAsync (hashcat_ctx, pw->i, device_param->hip_d_pws_comp_buf + (off * sizeof (u32)), cnt * sizeof (u32), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyDtoH (hashcat_ctx, pw->i, device_param->hip_d_pws_comp_buf + (off * sizeof (u32)), cnt * sizeof (u32)) == -1) return -1; if (hc_hipStreamSynchronize (hashcat_ctx, device_param->hip_stream) == -1) return -1; } @@ -1072,7 +1072,7 @@ int copy_pws_idx (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, u { if (hc_cuCtxPushCurrent (hashcat_ctx, device_param->cuda_context) == -1) return -1; - if (hc_cuMemcpyDtoHAsync (hashcat_ctx, dest, device_param->cuda_d_pws_idx + (gidd * sizeof (pw_idx_t)), (cnt * sizeof (pw_idx_t)), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyDtoH (hashcat_ctx, dest, device_param->cuda_d_pws_idx + (gidd * sizeof (pw_idx_t)), (cnt * sizeof (pw_idx_t))) == -1) return -1; if (hc_cuStreamSynchronize (hashcat_ctx, device_param->cuda_stream) == -1) return -1; @@ -1083,7 +1083,7 @@ int copy_pws_idx (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, u { if (hc_hipCtxPushCurrent (hashcat_ctx, device_param->hip_context) == -1) return -1; - if (hc_hipMemcpyDtoHAsync (hashcat_ctx, dest, device_param->hip_d_pws_idx + (gidd * sizeof (pw_idx_t)), (cnt * sizeof (pw_idx_t)), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyDtoH (hashcat_ctx, dest, device_param->hip_d_pws_idx + (gidd * sizeof (pw_idx_t)), (cnt * sizeof (pw_idx_t))) == -1) return -1; if (hc_hipStreamSynchronize (hashcat_ctx, device_param->hip_stream) == -1) return -1; @@ -1112,7 +1112,7 @@ int copy_pws_comp (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, { if (hc_cuCtxPushCurrent (hashcat_ctx, device_param->cuda_context) == -1) return -1; - if (hc_cuMemcpyDtoHAsync (hashcat_ctx, dest, device_param->cuda_d_pws_comp_buf + (off * sizeof (u32)), cnt * sizeof (u32), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyDtoH (hashcat_ctx, dest, device_param->cuda_d_pws_comp_buf + (off * sizeof (u32)), cnt * sizeof (u32)) == -1) return -1; if (hc_cuStreamSynchronize (hashcat_ctx, device_param->cuda_stream) == -1) return -1; @@ -1123,7 +1123,7 @@ int copy_pws_comp (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, { if (hc_hipCtxPushCurrent (hashcat_ctx, device_param->hip_context) == -1) return -1; - if (hc_hipMemcpyDtoHAsync (hashcat_ctx, dest, device_param->hip_d_pws_comp_buf + (off * sizeof (u32)), cnt * sizeof (u32), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyDtoH (hashcat_ctx, dest, device_param->hip_d_pws_comp_buf + (off * sizeof (u32)), cnt * sizeof (u32)) == -1) return -1; if (hc_hipStreamSynchronize (hashcat_ctx, device_param->hip_stream) == -1) return -1; @@ -1199,12 +1199,12 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, if (device_param->is_cuda == true) { - if (hc_cuMemcpyDtoDAsync (hashcat_ctx, device_param->cuda_d_bfs_c, device_param->cuda_d_tm_c, size_tm, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyDtoD (hashcat_ctx, device_param->cuda_d_bfs_c, device_param->cuda_d_tm_c, size_tm) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyDtoDAsync (hashcat_ctx, device_param->hip_d_bfs_c, device_param->hip_d_tm_c, size_tm, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyDtoD (hashcat_ctx, device_param->hip_d_bfs_c, device_param->hip_d_tm_c, size_tm) == -1) return -1; } #if defined (__APPLE__) @@ -1309,12 +1309,12 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, { if (device_param->is_cuda == true) { - if (hc_cuMemcpyDtoDAsync (hashcat_ctx, device_param->cuda_d_pws_buf, device_param->cuda_d_pws_amp_buf, pws_cnt * sizeof (pw_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyDtoD (hashcat_ctx, device_param->cuda_d_pws_buf, device_param->cuda_d_pws_amp_buf, pws_cnt * sizeof (pw_t)) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyDtoDAsync (hashcat_ctx, device_param->hip_d_pws_buf, device_param->hip_d_pws_amp_buf, pws_cnt * sizeof (pw_t), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyDtoD (hashcat_ctx, device_param->hip_d_pws_buf, device_param->hip_d_pws_amp_buf, pws_cnt * sizeof (pw_t)) == -1) return -1; } #if defined (__APPLE__) @@ -1373,14 +1373,14 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, if (device_param->is_cuda == true) { - if (hc_cuMemcpyDtoHAsync (hashcat_ctx, device_param->hooks_buf, device_param->cuda_d_hooks, pws_cnt * hashconfig->hook_size, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyDtoH (hashcat_ctx, device_param->hooks_buf, device_param->cuda_d_hooks, pws_cnt * hashconfig->hook_size) == -1) return -1; if (hc_cuStreamSynchronize (hashcat_ctx, device_param->cuda_stream) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyDtoHAsync (hashcat_ctx, device_param->hooks_buf, device_param->hip_d_hooks, pws_cnt * hashconfig->hook_size, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyDtoH (hashcat_ctx, device_param->hooks_buf, device_param->hip_d_hooks, pws_cnt * hashconfig->hook_size) == -1) return -1; if (hc_hipStreamSynchronize (hashcat_ctx, device_param->hip_stream) == -1) return -1; } @@ -1432,12 +1432,12 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_hooks, device_param->hooks_buf, pws_cnt * hashconfig->hook_size, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_hooks, device_param->hooks_buf, pws_cnt * hashconfig->hook_size) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_hooks, device_param->hooks_buf, pws_cnt * hashconfig->hook_size, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_hooks, device_param->hooks_buf, pws_cnt * hashconfig->hook_size) == -1) return -1; } #if defined (__APPLE__) @@ -1534,14 +1534,14 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, { if (device_param->is_cuda == true) { - if (hc_cuMemcpyDtoHAsync (hashcat_ctx, device_param->h_tmps, device_param->cuda_d_tmps, pws_cnt * hashconfig->tmp_size, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyDtoH (hashcat_ctx, device_param->h_tmps, device_param->cuda_d_tmps, pws_cnt * hashconfig->tmp_size) == -1) return -1; if (hc_cuStreamSynchronize (hashcat_ctx, device_param->cuda_stream) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyDtoHAsync (hashcat_ctx, device_param->h_tmps, device_param->hip_d_tmps, pws_cnt * hashconfig->tmp_size, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyDtoH (hashcat_ctx, device_param->h_tmps, device_param->hip_d_tmps, pws_cnt * hashconfig->tmp_size) == -1) return -1; if (hc_hipStreamSynchronize (hashcat_ctx, device_param->hip_stream) == -1) return -1; } @@ -1563,14 +1563,14 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_tmps, device_param->h_tmps, pws_cnt * hashconfig->tmp_size, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_tmps, device_param->h_tmps, pws_cnt * hashconfig->tmp_size) == -1) return -1; if (hc_cuStreamSynchronize (hashcat_ctx, device_param->cuda_stream) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_tmps, device_param->h_tmps, pws_cnt * hashconfig->tmp_size, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_tmps, device_param->h_tmps, pws_cnt * hashconfig->tmp_size) == -1) return -1; if (hc_hipStreamSynchronize (hashcat_ctx, device_param->hip_stream) == -1) return -1; } @@ -1625,14 +1625,14 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, if (device_param->is_cuda == true) { - if (hc_cuMemcpyDtoHAsync (hashcat_ctx, device_param->hooks_buf, device_param->cuda_d_hooks, pws_cnt * hashconfig->hook_size, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyDtoH (hashcat_ctx, device_param->hooks_buf, device_param->cuda_d_hooks, pws_cnt * hashconfig->hook_size) == -1) return -1; if (hc_cuStreamSynchronize (hashcat_ctx, device_param->cuda_stream) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyDtoHAsync (hashcat_ctx, device_param->hooks_buf, device_param->hip_d_hooks, pws_cnt * hashconfig->hook_size, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyDtoH (hashcat_ctx, device_param->hooks_buf, device_param->hip_d_hooks, pws_cnt * hashconfig->hook_size) == -1) return -1; if (hc_hipStreamSynchronize (hashcat_ctx, device_param->hip_stream) == -1) return -1; } @@ -1684,12 +1684,12 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_hooks, device_param->hooks_buf, pws_cnt * hashconfig->hook_size, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_hooks, device_param->hooks_buf, pws_cnt * hashconfig->hook_size) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_hooks, device_param->hooks_buf, pws_cnt * hashconfig->hook_size, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_hooks, device_param->hooks_buf, pws_cnt * hashconfig->hook_size) == -1) return -1; } #if defined (__APPLE__) @@ -1774,14 +1774,14 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, { if (device_param->is_cuda == true) { - if (hc_cuMemcpyDtoHAsync (hashcat_ctx, device_param->h_tmps, device_param->cuda_d_tmps, pws_cnt * hashconfig->tmp_size, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyDtoH (hashcat_ctx, device_param->h_tmps, device_param->cuda_d_tmps, pws_cnt * hashconfig->tmp_size) == -1) return -1; if (hc_cuStreamSynchronize (hashcat_ctx, device_param->cuda_stream) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyDtoHAsync (hashcat_ctx, device_param->h_tmps, device_param->hip_d_tmps, pws_cnt * hashconfig->tmp_size, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyDtoH (hashcat_ctx, device_param->h_tmps, device_param->hip_d_tmps, pws_cnt * hashconfig->tmp_size) == -1) return -1; if (hc_hipStreamSynchronize (hashcat_ctx, device_param->hip_stream) == -1) return -1; } @@ -1803,14 +1803,14 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_tmps, device_param->h_tmps, pws_cnt * hashconfig->tmp_size, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_tmps, device_param->h_tmps, pws_cnt * hashconfig->tmp_size) == -1) return -1; if (hc_cuStreamSynchronize (hashcat_ctx, device_param->cuda_stream) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_tmps, device_param->h_tmps, pws_cnt * hashconfig->tmp_size, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_tmps, device_param->h_tmps, pws_cnt * hashconfig->tmp_size) == -1) return -1; if (hc_hipStreamSynchronize (hashcat_ctx, device_param->hip_stream) == -1) return -1; } @@ -2059,12 +2059,12 @@ int run_cuda_kernel_utf8toutf16le (hashcat_ctx_t *hashcat_ctx, hc_device_param_t return 0; } -int run_cuda_kernel_memset (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, CUdeviceptr buf, const u64 offset, const u8 value, const u64 size) +int run_cuda_kernel_memset (hashcat_ctx_t *hashcat_ctx, MAYBE_UNUSED hc_device_param_t *device_param, CUdeviceptr buf, const u64 offset, const u8 value, const u64 size) { - return hc_cuMemsetD8Async (hashcat_ctx, buf + offset, value, size, device_param->cuda_stream); + return hc_cuMemsetD8 (hashcat_ctx, buf + offset, value, size); } -int run_cuda_kernel_memset32 (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, CUdeviceptr buf, const u64 offset, const u32 value, const u64 size) +int run_cuda_kernel_memset32 (hashcat_ctx_t *hashcat_ctx, MAYBE_UNUSED hc_device_param_t *device_param, CUdeviceptr buf, const u64 offset, const u32 value, const u64 size) { /* check that the size is multiple of element size */ if (size % 4 != 0) @@ -2072,7 +2072,7 @@ int run_cuda_kernel_memset32 (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *dev return CUDA_ERROR_INVALID_VALUE; } - return hc_cuMemsetD32Async (hashcat_ctx, buf + offset, value, size / 4, device_param->cuda_stream); + return hc_cuMemsetD32 (hashcat_ctx, buf + offset, value, size / 4); } int run_cuda_kernel_bzero (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, CUdeviceptr buf, const u64 size) @@ -2096,7 +2096,7 @@ int run_cuda_kernel_bzero (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device if (num16m) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, buf + (num16d * 16), bzeros, num16m, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, buf + (num16d * 16), bzeros, num16m) == -1) return -1; } return 0; @@ -2138,12 +2138,12 @@ int run_hip_kernel_utf8toutf16le (hashcat_ctx_t *hashcat_ctx, hc_device_param_t return 0; } -int run_hip_kernel_memset (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, hipDeviceptr_t buf, const u64 offset, const u8 value, const u64 size) +int run_hip_kernel_memset (hashcat_ctx_t *hashcat_ctx, MAYBE_UNUSED hc_device_param_t *device_param, hipDeviceptr_t buf, const u64 offset, const u8 value, const u64 size) { - return hc_hipMemsetD8Async (hashcat_ctx, buf + offset, value, size, device_param->hip_stream); + return hc_hipMemsetD8 (hashcat_ctx, buf + offset, value, size); } -int run_hip_kernel_memset32 (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, hipDeviceptr_t buf, const u64 offset, const u32 value, const u64 size) +int run_hip_kernel_memset32 (hashcat_ctx_t *hashcat_ctx, MAYBE_UNUSED hc_device_param_t *device_param, hipDeviceptr_t buf, const u64 offset, const u32 value, const u64 size) { /* check that the size is multiple of element size */ if (size % 4 != 0) @@ -2151,7 +2151,7 @@ int run_hip_kernel_memset32 (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *devi return hipErrorInvalidValue; } - return hc_hipMemsetD32Async (hashcat_ctx, buf + offset, value, size / 4, device_param->hip_stream); + return hc_hipMemsetD32 (hashcat_ctx, buf + offset, value, size / 4); } int run_hip_kernel_bzero (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, hipDeviceptr_t buf, const u64 size) @@ -2175,7 +2175,7 @@ int run_hip_kernel_bzero (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_ if (num16m) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, buf + (num16d * 16), bzeros, num16m, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, buf + (num16d * 16), bzeros, num16m) == -1) return -1; } return 0; @@ -2466,7 +2466,7 @@ int run_opencl_kernel_bzero (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *devi } else { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, buf, CL_FALSE, num16d * 16, num16m, bzeros, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, buf, CL_TRUE, num16d * 16, num16m, bzeros, 0, NULL, NULL) == -1) return -1; } } @@ -2592,13 +2592,16 @@ int run_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, con case KERN_RUN_AUX4: cuda_function = device_param->cuda_function_aux4; break; } - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_kernel_param, &device_param->kernel_param, device_param->size_kernel_params, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_kernel_param, &device_param->kernel_param, device_param->size_kernel_params) == -1) return -1; if (hc_cuFuncSetAttribute (hashcat_ctx, cuda_function, CU_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES, dynamic_shared_mem) == -1) return -1; if (kernel_threads == 0) kernel_threads = 1; - num_elements = CEILDIV (num_elements, kernel_threads); + if ((hashconfig->opts_type & OPTS_TYPE_THREAD_MULTI_DISABLE) == 0) + { + num_elements = CEILDIV (num_elements, kernel_threads); + } if (kern_run == KERN_RUN_1) { @@ -2636,14 +2639,29 @@ int run_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, con } } + u32 gridDimX = num_elements; + u32 gridDimY = 1; + u32 gridDimZ = 1; + + u32 blockDimX = kernel_threads; + u32 blockDimY = 1; + u32 blockDimZ = 1; + + if ((hashconfig->opti_type & OPTI_TYPE_SLOW_HASH_DIMY_INIT) && (kern_run == KERN_RUN_1)) + blockDimY = hashcat_ctx->hashes->salts_buf->salt_dimy; + if ((hashconfig->opti_type & OPTI_TYPE_SLOW_HASH_DIMY_LOOP) && (kern_run == KERN_RUN_2)) + blockDimY = hashcat_ctx->hashes->salts_buf->salt_dimy; + if ((hashconfig->opti_type & OPTI_TYPE_SLOW_HASH_DIMY_COMP) && (kern_run == KERN_RUN_3)) + blockDimY = hashcat_ctx->hashes->salts_buf->salt_dimy; + if (is_autotune == true) { - if (hc_cuLaunchKernel (hashcat_ctx, cuda_function, num_elements, 1, 1, kernel_threads, 1, 1, dynamic_shared_mem, device_param->cuda_stream, device_param->kernel_params, NULL) == -1) return -1; + if (hc_cuLaunchKernel (hashcat_ctx, cuda_function, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, dynamic_shared_mem, device_param->cuda_stream, device_param->kernel_params, NULL) == -1) return -1; } if (hc_cuEventRecord (hashcat_ctx, device_param->cuda_event1, device_param->cuda_stream) == -1) return -1; - if (hc_cuLaunchKernel (hashcat_ctx, cuda_function, num_elements, 1, 1, kernel_threads, 1, 1, dynamic_shared_mem, device_param->cuda_stream, device_param->kernel_params, NULL) == -1) return -1; + if (hc_cuLaunchKernel (hashcat_ctx, cuda_function, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, dynamic_shared_mem, device_param->cuda_stream, device_param->kernel_params, NULL) == -1) return -1; if (hc_cuEventRecord (hashcat_ctx, device_param->cuda_event2, device_param->cuda_stream) == -1) return -1; @@ -2693,13 +2711,16 @@ int run_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, con case KERN_RUN_AUX4: hip_function = device_param->hip_function_aux4; break; } - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_kernel_param, &device_param->kernel_param, device_param->size_kernel_params, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_kernel_param, &device_param->kernel_param, device_param->size_kernel_params) == -1) return -1; //if (hc_hipFuncSetAttribute (hashcat_ctx, hip_function, HIP_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES, dynamic_shared_mem) == -1) return -1; if (kernel_threads == 0) kernel_threads = 1; - num_elements = CEILDIV (num_elements, kernel_threads); + if ((hashconfig->opts_type & OPTS_TYPE_THREAD_MULTI_DISABLE) == 0) + { + num_elements = CEILDIV (num_elements, kernel_threads); + } if (kern_run == KERN_RUN_1) { @@ -2737,14 +2758,31 @@ int run_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, con } } + u32 gridDimX = num_elements; + u32 gridDimY = 1; + u32 gridDimZ = 1; + + u32 blockDimX = kernel_threads; + u32 blockDimY = 1; + u32 blockDimZ = 1; + + if ((hashconfig->opti_type & OPTI_TYPE_SLOW_HASH_DIMY_INIT) && (kern_run == KERN_RUN_1)) + blockDimY = hashcat_ctx->hashes->salts_buf->salt_dimy; + if ((hashconfig->opti_type & OPTI_TYPE_SLOW_HASH_DIMY_LOOP) && (kern_run == KERN_RUN_2)) + blockDimY = hashcat_ctx->hashes->salts_buf->salt_dimy; + if ((hashconfig->opti_type & OPTI_TYPE_SLOW_HASH_DIMY_COMP) && (kern_run == KERN_RUN_3)) + blockDimY = hashcat_ctx->hashes->salts_buf->salt_dimy; + + //printf ("%d %d %d %d %d %d %d\n", kern_run, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ); + if (is_autotune == true) { - if (hc_hipLaunchKernel (hashcat_ctx, hip_function, num_elements, 1, 1, kernel_threads, 1, 1, dynamic_shared_mem, device_param->hip_stream, device_param->kernel_params, NULL) == -1) return -1; + if (hc_hipLaunchKernel (hashcat_ctx, hip_function, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, dynamic_shared_mem, device_param->hip_stream, device_param->kernel_params, NULL) == -1) return -1; } if (hc_hipEventRecord (hashcat_ctx, device_param->hip_event1, device_param->hip_stream) == -1) return -1; - if (hc_hipLaunchKernel (hashcat_ctx, hip_function, num_elements, 1, 1, kernel_threads, 1, 1, dynamic_shared_mem, device_param->hip_stream, device_param->kernel_params, NULL) == -1) return -1; + if (hc_hipLaunchKernel (hashcat_ctx, hip_function, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, dynamic_shared_mem, device_param->hip_stream, device_param->kernel_params, NULL) == -1) return -1; if (hc_hipEventRecord (hashcat_ctx, device_param->hip_event2, device_param->hip_stream) == -1) return -1; @@ -2944,7 +2982,7 @@ int run_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, con if (hc_clSetKernelArg (hashcat_ctx, opencl_kernel, i, sizeof (cl_mem), device_param->kernel_params[i]) == -1) return -1; } - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_kernel_param, CL_FALSE, 0, device_param->size_kernel_params, &device_param->kernel_param, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_kernel_param, CL_TRUE, 0, device_param->size_kernel_params, &device_param->kernel_param, 0, NULL, NULL) == -1) return -1; /* for (u32 i = 24; i <= 34; i++) @@ -2984,17 +3022,44 @@ int run_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, con } } - num_elements = round_up_multiple_64 (num_elements, kernel_threads); + if ((hashconfig->opts_type & OPTS_TYPE_THREAD_MULTI_DISABLE) == 0) + { + num_elements = round_up_multiple_64 (num_elements, kernel_threads); + } + else + { + num_elements = num_elements * kernel_threads; + } - const size_t global_work_size[3] = { num_elements, 1, 1 }; - const size_t local_work_size[3] = { kernel_threads, 1, 1 }; + size_t global_work_size[3] = { num_elements, 1, 1 }; + size_t local_work_size[3] = { kernel_threads, 1, 1 }; + + cl_uint work_dim = 1; + + if ((hashconfig->opti_type & OPTI_TYPE_SLOW_HASH_DIMY_INIT) && (kern_run == KERN_RUN_1)) + { + global_work_size[1] = local_work_size[1] = hashcat_ctx->hashes->salts_buf->salt_dimy; + work_dim = 2; + } + + if ((hashconfig->opti_type & OPTI_TYPE_SLOW_HASH_DIMY_LOOP) && (kern_run == KERN_RUN_2)) + { + global_work_size[1] = local_work_size[1] = hashcat_ctx->hashes->salts_buf->salt_dimy; + work_dim = 2; + } + + if ((hashconfig->opti_type & OPTI_TYPE_SLOW_HASH_DIMY_COMP) && (kern_run == KERN_RUN_3)) + { + global_work_size[1] = local_work_size[1] = hashcat_ctx->hashes->salts_buf->salt_dimy; + work_dim = 2; + } if (is_autotune == true) { - if (hc_clEnqueueNDRangeKernel (hashcat_ctx, device_param->opencl_command_queue, opencl_kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, &opencl_event) == -1) return -1; + if (hc_clEnqueueNDRangeKernel (hashcat_ctx, device_param->opencl_command_queue, opencl_kernel, work_dim, NULL, global_work_size, local_work_size, 0, NULL, &opencl_event) == -1) return -1; } - if (hc_clEnqueueNDRangeKernel (hashcat_ctx, device_param->opencl_command_queue, opencl_kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, &opencl_event) == -1) return -1; + if (hc_clEnqueueNDRangeKernel (hashcat_ctx, device_param->opencl_command_queue, opencl_kernel, work_dim, NULL, global_work_size, local_work_size, 0, NULL, &opencl_event) == -1) return -1; // spin damper section @@ -3551,7 +3616,7 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const { if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t)) == -1) return -1; const pw_idx_t *pw_idx = device_param->pws_idx + pws_cnt; @@ -3559,13 +3624,13 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const if (off) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32)) == -1) return -1; } } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t)) == -1) return -1; const pw_idx_t *pw_idx = device_param->pws_idx + pws_cnt; @@ -3573,7 +3638,7 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const if (off) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32)) == -1) return -1; } } @@ -3595,7 +3660,7 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const if (device_param->is_opencl == true) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_idx, CL_FALSE, 0, pws_cnt * sizeof (pw_idx_t), device_param->pws_idx, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_idx, CL_TRUE, 0, pws_cnt * sizeof (pw_idx_t), device_param->pws_idx, 0, NULL, NULL) == -1) return -1; const pw_idx_t *pw_idx = device_param->pws_idx + pws_cnt; @@ -3603,7 +3668,7 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const if (off) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_comp_buf, CL_FALSE, 0, off * sizeof (u32), device_param->pws_comp, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_comp_buf, CL_TRUE, 0, off * sizeof (u32), device_param->pws_comp, 0, NULL, NULL) == -1) return -1; } } @@ -3615,7 +3680,7 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const { if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t)) == -1) return -1; const pw_idx_t *pw_idx = device_param->pws_idx + pws_cnt; @@ -3623,13 +3688,13 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const if (off) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32)) == -1) return -1; } } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t)) == -1) return -1; const pw_idx_t *pw_idx = device_param->pws_idx + pws_cnt; @@ -3637,7 +3702,7 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const if (off) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32)) == -1) return -1; } } @@ -3659,7 +3724,7 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const if (device_param->is_opencl == true) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_idx, CL_FALSE, 0, pws_cnt * sizeof (pw_idx_t), device_param->pws_idx, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_idx, CL_TRUE, 0, pws_cnt * sizeof (pw_idx_t), device_param->pws_idx, 0, NULL, NULL) == -1) return -1; const pw_idx_t *pw_idx = device_param->pws_idx + pws_cnt; @@ -3667,7 +3732,7 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const if (off) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_comp_buf, CL_FALSE, 0, off * sizeof (u32), device_param->pws_comp, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_comp_buf, CL_TRUE, 0, off * sizeof (u32), device_param->pws_comp, 0, NULL, NULL) == -1) return -1; } } @@ -3713,7 +3778,7 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t)) == -1) return -1; const pw_idx_t *pw_idx = device_param->pws_idx + pws_cnt; @@ -3721,13 +3786,13 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const if (off) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32)) == -1) return -1; } } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t)) == -1) return -1; const pw_idx_t *pw_idx = device_param->pws_idx + pws_cnt; @@ -3735,7 +3800,7 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const if (off) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32)) == -1) return -1; } } @@ -3757,7 +3822,7 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const if (device_param->is_opencl == true) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_idx, CL_FALSE, 0, pws_cnt * sizeof (pw_idx_t), device_param->pws_idx, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_idx, CL_TRUE, 0, pws_cnt * sizeof (pw_idx_t), device_param->pws_idx, 0, NULL, NULL) == -1) return -1; const pw_idx_t *pw_idx = device_param->pws_idx + pws_cnt; @@ -3765,7 +3830,7 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const if (off) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_comp_buf, CL_FALSE, 0, off * sizeof (u32), device_param->pws_comp, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_comp_buf, CL_TRUE, 0, off * sizeof (u32), device_param->pws_comp, 0, NULL, NULL) == -1) return -1; } } @@ -3777,7 +3842,7 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const { if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t)) == -1) return -1; const pw_idx_t *pw_idx = device_param->pws_idx + pws_cnt; @@ -3785,13 +3850,13 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const if (off) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32)) == -1) return -1; } } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t)) == -1) return -1; const pw_idx_t *pw_idx = device_param->pws_idx + pws_cnt; @@ -3799,7 +3864,7 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const if (off) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32)) == -1) return -1; } } @@ -3821,7 +3886,7 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const if (device_param->is_opencl == true) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_idx, CL_FALSE, 0, pws_cnt * sizeof (pw_idx_t), device_param->pws_idx, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_idx, CL_TRUE, 0, pws_cnt * sizeof (pw_idx_t), device_param->pws_idx, 0, NULL, NULL) == -1) return -1; const pw_idx_t *pw_idx = device_param->pws_idx + pws_cnt; @@ -3829,7 +3894,7 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const if (off) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_comp_buf, CL_FALSE, 0, off * sizeof (u32), device_param->pws_comp, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_comp_buf, CL_TRUE, 0, off * sizeof (u32), device_param->pws_comp, 0, NULL, NULL) == -1) return -1; } } @@ -3839,7 +3904,7 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const { if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t)) == -1) return -1; const pw_idx_t *pw_idx = device_param->pws_idx + pws_cnt; @@ -3847,13 +3912,13 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const if (off) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32)) == -1) return -1; } } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t)) == -1) return -1; const pw_idx_t *pw_idx = device_param->pws_idx + pws_cnt; @@ -3861,7 +3926,7 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const if (off) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32)) == -1) return -1; } } @@ -3883,7 +3948,7 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const if (device_param->is_opencl == true) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_idx, CL_FALSE, 0, pws_cnt * sizeof (pw_idx_t), device_param->pws_idx, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_idx, CL_TRUE, 0, pws_cnt * sizeof (pw_idx_t), device_param->pws_idx, 0, NULL, NULL) == -1) return -1; const pw_idx_t *pw_idx = device_param->pws_idx + pws_cnt; @@ -3891,7 +3956,7 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const if (off) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_comp_buf, CL_FALSE, 0, off * sizeof (u32), device_param->pws_comp, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_comp_buf, CL_TRUE, 0, off * sizeof (u32), device_param->pws_comp, 0, NULL, NULL) == -1) return -1; } } @@ -4125,12 +4190,12 @@ int run_cracker (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, co { if (device_param->is_cuda == true) { - if (hc_cuMemcpyDtoDAsync (hashcat_ctx, device_param->cuda_d_rules_c, device_param->cuda_d_rules + (innerloop_pos * sizeof (kernel_rule_t)), innerloop_left * sizeof (kernel_rule_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyDtoD (hashcat_ctx, device_param->cuda_d_rules_c, device_param->cuda_d_rules + (innerloop_pos * sizeof (kernel_rule_t)), innerloop_left * sizeof (kernel_rule_t)) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyDtoDAsync (hashcat_ctx, device_param->hip_d_rules_c, device_param->hip_d_rules + (innerloop_pos * sizeof (kernel_rule_t)), innerloop_left * sizeof (kernel_rule_t), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyDtoD (hashcat_ctx, device_param->hip_d_rules_c, device_param->hip_d_rules + (innerloop_pos * sizeof (kernel_rule_t)), innerloop_left * sizeof (kernel_rule_t)) == -1) return -1; } #if defined (__APPLE__) @@ -4257,12 +4322,12 @@ int run_cracker (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, co if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_combs_c, device_param->combs_buf, innerloop_left * sizeof (pw_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_combs_c, device_param->combs_buf, innerloop_left * sizeof (pw_t)) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_combs_c, device_param->combs_buf, innerloop_left * sizeof (pw_t), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_combs_c, device_param->combs_buf, innerloop_left * sizeof (pw_t)) == -1) return -1; } #if defined (__APPLE__) @@ -4274,7 +4339,7 @@ int run_cracker (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, co if (device_param->is_opencl == true) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_combs_c, CL_FALSE, 0, innerloop_left * sizeof (pw_t), device_param->combs_buf, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_combs_c, CL_TRUE, 0, innerloop_left * sizeof (pw_t), device_param->combs_buf, 0, NULL, NULL) == -1) return -1; } } else if (user_options->attack_mode == ATTACK_MODE_HYBRID1) @@ -4287,12 +4352,12 @@ int run_cracker (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, co if (device_param->is_cuda == true) { - if (hc_cuMemcpyDtoDAsync (hashcat_ctx, device_param->cuda_d_combs_c, device_param->cuda_d_combs, innerloop_left * sizeof (pw_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyDtoD (hashcat_ctx, device_param->cuda_d_combs_c, device_param->cuda_d_combs, innerloop_left * sizeof (pw_t)) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyDtoDAsync (hashcat_ctx, device_param->hip_d_combs_c, device_param->hip_d_combs, innerloop_left * sizeof (pw_t), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyDtoD (hashcat_ctx, device_param->hip_d_combs_c, device_param->hip_d_combs, innerloop_left * sizeof (pw_t)) == -1) return -1; } #if defined (__APPLE__) @@ -4317,12 +4382,12 @@ int run_cracker (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, co if (device_param->is_cuda == true) { - if (hc_cuMemcpyDtoDAsync (hashcat_ctx, device_param->cuda_d_combs_c, device_param->cuda_d_combs, innerloop_left * sizeof (pw_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyDtoD (hashcat_ctx, device_param->cuda_d_combs_c, device_param->cuda_d_combs, innerloop_left * sizeof (pw_t)) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyDtoDAsync (hashcat_ctx, device_param->hip_d_combs_c, device_param->hip_d_combs, innerloop_left * sizeof (pw_t), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyDtoD (hashcat_ctx, device_param->hip_d_combs_c, device_param->hip_d_combs, innerloop_left * sizeof (pw_t)) == -1) return -1; } #if defined (__APPLE__) @@ -4450,12 +4515,12 @@ int run_cracker (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, co if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_combs_c, device_param->combs_buf, innerloop_left * sizeof (pw_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_combs_c, device_param->combs_buf, innerloop_left * sizeof (pw_t)) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_combs_c, device_param->combs_buf, innerloop_left * sizeof (pw_t), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_combs_c, device_param->combs_buf, innerloop_left * sizeof (pw_t)) == -1) return -1; } #if defined (__APPLE__) @@ -4467,7 +4532,7 @@ int run_cracker (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, co if (device_param->is_opencl == true) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_combs_c, CL_FALSE, 0, innerloop_left * sizeof (pw_t), device_param->combs_buf, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_combs_c, CL_TRUE, 0, innerloop_left * sizeof (pw_t), device_param->combs_buf, 0, NULL, NULL) == -1) return -1; } } else if (user_options->attack_mode == ATTACK_MODE_HYBRID1) @@ -4480,12 +4545,12 @@ int run_cracker (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, co if (device_param->is_cuda == true) { - if (hc_cuMemcpyDtoDAsync (hashcat_ctx, device_param->cuda_d_combs_c, device_param->cuda_d_combs, innerloop_left * sizeof (pw_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyDtoD (hashcat_ctx, device_param->cuda_d_combs_c, device_param->cuda_d_combs, innerloop_left * sizeof (pw_t)) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyDtoDAsync (hashcat_ctx, device_param->hip_d_combs_c, device_param->hip_d_combs, innerloop_left * sizeof (pw_t), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyDtoD (hashcat_ctx, device_param->hip_d_combs_c, device_param->hip_d_combs, innerloop_left * sizeof (pw_t)) == -1) return -1; } #if defined (__APPLE__) @@ -4512,12 +4577,12 @@ int run_cracker (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, co if (device_param->is_cuda == true) { - if (hc_cuMemcpyDtoDAsync (hashcat_ctx, device_param->cuda_d_bfs_c, device_param->cuda_d_bfs, innerloop_left * sizeof (bf_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyDtoD (hashcat_ctx, device_param->cuda_d_bfs_c, device_param->cuda_d_bfs, innerloop_left * sizeof (bf_t)) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyDtoDAsync (hashcat_ctx, device_param->hip_d_bfs_c, device_param->hip_d_bfs, innerloop_left * sizeof (bf_t), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyDtoD (hashcat_ctx, device_param->hip_d_bfs_c, device_param->hip_d_bfs, innerloop_left * sizeof (bf_t)) == -1) return -1; } #if defined (__APPLE__) @@ -8952,7 +9017,8 @@ void backend_ctx_devices_sync_tuning (hashcat_ctx_t *hashcat_ctx) device_param_dst->kernel_loops = device_param_src->kernel_loops; device_param_dst->kernel_threads = device_param_src->kernel_threads; - const u32 hardware_power = ((hashconfig->opts_type & OPTS_TYPE_MP_MULTI_DISABLE) ? 1 : device_param_dst->device_processors) * device_param_dst->kernel_threads; + const u32 hardware_power = ((hashconfig->opts_type & OPTS_TYPE_MP_MULTI_DISABLE) ? 1 : device_param_dst->device_processors) + * ((hashconfig->opts_type & OPTS_TYPE_THREAD_MULTI_DISABLE) ? 1 : device_param_dst->kernel_threads); device_param_dst->hardware_power = hardware_power; @@ -9522,7 +9588,11 @@ static bool load_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_p hc_asprintf (&hiprtc_options[hiprtc_options_idx++], "-D MAX_THREADS_PER_BLOCK=%d", (user_options->kernel_threads_chgd == true) ? user_options->kernel_threads : device_param->kernel_threads_max); hc_asprintf (&hiprtc_options[hiprtc_options_idx++], "--gpu-architecture=%s", device_param->gcnArchName); - hc_asprintf (&hiprtc_options[hiprtc_options_idx++], "--gpu-max-threads-per-block=%d", (user_options->kernel_threads_chgd == true) ? user_options->kernel_threads : device_param->kernel_threads_max); + + if ((hashconfig->opts_type & OPTS_TYPE_THREAD_MULTI_DISABLE) == 0) + { + hc_asprintf (&hiprtc_options[hiprtc_options_idx++], "--gpu-max-threads-per-block=%d", (user_options->kernel_threads_chgd == true) ? user_options->kernel_threads : device_param->kernel_threads_max); + } // untested but it should work #if defined (_WIN) || defined (__CYGWIN__) || defined (__MSYS__) @@ -10436,7 +10506,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) * device properties */ - const u32 device_processors = device_param->device_processors; + //const u32 device_processors = device_param->device_processors; /** * device threads @@ -11793,16 +11863,16 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (hc_cuMemAlloc (hashcat_ctx, &device_param->cuda_d_st_salts_buf, size_st_salts) == -1) return -1; if (hc_cuMemAlloc (hashcat_ctx, &device_param->cuda_d_kernel_param, size_kernel_params) == -1) return -1; - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_bitmap_s1_a, bitmap_ctx->bitmap_s1_a, bitmap_ctx->bitmap_size, device_param->cuda_stream) == -1) return -1; - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_bitmap_s1_b, bitmap_ctx->bitmap_s1_b, bitmap_ctx->bitmap_size, device_param->cuda_stream) == -1) return -1; - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_bitmap_s1_c, bitmap_ctx->bitmap_s1_c, bitmap_ctx->bitmap_size, device_param->cuda_stream) == -1) return -1; - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_bitmap_s1_d, bitmap_ctx->bitmap_s1_d, bitmap_ctx->bitmap_size, device_param->cuda_stream) == -1) return -1; - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_bitmap_s2_a, bitmap_ctx->bitmap_s2_a, bitmap_ctx->bitmap_size, device_param->cuda_stream) == -1) return -1; - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_bitmap_s2_b, bitmap_ctx->bitmap_s2_b, bitmap_ctx->bitmap_size, device_param->cuda_stream) == -1) return -1; - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_bitmap_s2_c, bitmap_ctx->bitmap_s2_c, bitmap_ctx->bitmap_size, device_param->cuda_stream) == -1) return -1; - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_bitmap_s2_d, bitmap_ctx->bitmap_s2_d, bitmap_ctx->bitmap_size, device_param->cuda_stream) == -1) return -1; - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_digests_buf, hashes->digests_buf, size_digests, device_param->cuda_stream) == -1) return -1; - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_salt_bufs, hashes->salts_buf, size_salts, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_bitmap_s1_a, bitmap_ctx->bitmap_s1_a, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_bitmap_s1_b, bitmap_ctx->bitmap_s1_b, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_bitmap_s1_c, bitmap_ctx->bitmap_s1_c, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_bitmap_s1_d, bitmap_ctx->bitmap_s1_d, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_bitmap_s2_a, bitmap_ctx->bitmap_s2_a, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_bitmap_s2_b, bitmap_ctx->bitmap_s2_b, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_bitmap_s2_c, bitmap_ctx->bitmap_s2_c, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_bitmap_s2_d, bitmap_ctx->bitmap_s2_d, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_digests_buf, hashes->digests_buf, size_digests) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_salt_bufs, hashes->salts_buf, size_salts) == -1) return -1; /** * special buffers @@ -11829,7 +11899,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (hc_cuMemAlloc (hashcat_ctx, &device_param->cuda_d_rules_c, size_rules_c) == -1) return -1; } - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_rules, straight_ctx->kernel_rules_buf, size_rules_src, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_rules, straight_ctx->kernel_rules_buf, size_rules_src) == -1) return -1; } else if (user_options_extra->attack_kern == ATTACK_KERN_COMBI) { @@ -11864,19 +11934,19 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) { if (hc_cuMemAlloc (hashcat_ctx, &device_param->cuda_d_esalt_bufs, size_esalts) == -1) return -1; - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_esalt_bufs, hashes->esalts_buf, size_esalts, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_esalt_bufs, hashes->esalts_buf, size_esalts) == -1) return -1; } if (hashconfig->st_hash != NULL) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_st_digests_buf, hashes->st_digests_buf, size_st_digests, device_param->cuda_stream) == -1) return -1; - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_st_salts_buf, hashes->st_salts_buf, size_st_salts, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_st_digests_buf, hashes->st_digests_buf, size_st_digests) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_st_salts_buf, hashes->st_salts_buf, size_st_salts) == -1) return -1; if (size_esalts) { if (hc_cuMemAlloc (hashcat_ctx, &device_param->cuda_d_st_esalts_buf, size_st_esalts) == -1) return -1; - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_st_esalts_buf, hashes->st_esalts_buf, size_st_esalts, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_st_esalts_buf, hashes->st_esalts_buf, size_st_esalts) == -1) return -1; } } } @@ -11904,16 +11974,16 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_st_salts_buf, size_st_salts) == -1) return -1; if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_kernel_param, size_kernel_params) == -1) return -1; - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_bitmap_s1_a, bitmap_ctx->bitmap_s1_a, bitmap_ctx->bitmap_size, device_param->hip_stream) == -1) return -1; - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_bitmap_s1_b, bitmap_ctx->bitmap_s1_b, bitmap_ctx->bitmap_size, device_param->hip_stream) == -1) return -1; - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_bitmap_s1_c, bitmap_ctx->bitmap_s1_c, bitmap_ctx->bitmap_size, device_param->hip_stream) == -1) return -1; - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_bitmap_s1_d, bitmap_ctx->bitmap_s1_d, bitmap_ctx->bitmap_size, device_param->hip_stream) == -1) return -1; - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_bitmap_s2_a, bitmap_ctx->bitmap_s2_a, bitmap_ctx->bitmap_size, device_param->hip_stream) == -1) return -1; - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_bitmap_s2_b, bitmap_ctx->bitmap_s2_b, bitmap_ctx->bitmap_size, device_param->hip_stream) == -1) return -1; - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_bitmap_s2_c, bitmap_ctx->bitmap_s2_c, bitmap_ctx->bitmap_size, device_param->hip_stream) == -1) return -1; - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_bitmap_s2_d, bitmap_ctx->bitmap_s2_d, bitmap_ctx->bitmap_size, device_param->hip_stream) == -1) return -1; - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_digests_buf, hashes->digests_buf, size_digests, device_param->hip_stream) == -1) return -1; - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_salt_bufs, hashes->salts_buf, size_salts, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_bitmap_s1_a, bitmap_ctx->bitmap_s1_a, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_bitmap_s1_b, bitmap_ctx->bitmap_s1_b, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_bitmap_s1_c, bitmap_ctx->bitmap_s1_c, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_bitmap_s1_d, bitmap_ctx->bitmap_s1_d, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_bitmap_s2_a, bitmap_ctx->bitmap_s2_a, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_bitmap_s2_b, bitmap_ctx->bitmap_s2_b, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_bitmap_s2_c, bitmap_ctx->bitmap_s2_c, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_bitmap_s2_d, bitmap_ctx->bitmap_s2_d, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_digests_buf, hashes->digests_buf, size_digests) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_salt_bufs, hashes->salts_buf, size_salts) == -1) return -1; /** * special buffers @@ -11940,7 +12010,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_rules_c, size_rules_c) == -1) return -1; } - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_rules, straight_ctx->kernel_rules_buf, size_rules_src, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_rules, straight_ctx->kernel_rules_buf, size_rules_src) == -1) return -1; } else if (user_options_extra->attack_kern == ATTACK_KERN_COMBI) { @@ -11975,19 +12045,19 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) { if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_esalt_bufs, size_esalts) == -1) return -1; - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_esalt_bufs, hashes->esalts_buf, size_esalts, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_esalt_bufs, hashes->esalts_buf, size_esalts) == -1) return -1; } if (hashconfig->st_hash != NULL) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_st_digests_buf, hashes->st_digests_buf, size_st_digests, device_param->hip_stream) == -1) return -1; - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_st_salts_buf, hashes->st_salts_buf, size_st_salts, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_st_digests_buf, hashes->st_digests_buf, size_st_digests) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_st_salts_buf, hashes->st_salts_buf, size_st_salts) == -1) return -1; if (size_esalts) { if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_st_esalts_buf, size_st_esalts) == -1) return -1; - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_st_esalts_buf, hashes->st_esalts_buf, size_st_esalts, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_st_esalts_buf, hashes->st_esalts_buf, size_st_esalts) == -1) return -1; } } } @@ -12125,15 +12195,15 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (hc_clCreateBuffer (hashcat_ctx, device_param->opencl_context, CL_MEM_READ_ONLY, size_st_salts, NULL, &device_param->opencl_d_st_salts_buf) == -1) return -1; if (hc_clCreateBuffer (hashcat_ctx, device_param->opencl_context, CL_MEM_READ_ONLY, size_kernel_params, NULL, &device_param->opencl_d_kernel_param) == -1) return -1; - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_bitmap_s1_a, CL_FALSE, 0, bitmap_ctx->bitmap_size, bitmap_ctx->bitmap_s1_a, 0, NULL, NULL) == -1) return -1; - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_bitmap_s1_b, CL_FALSE, 0, bitmap_ctx->bitmap_size, bitmap_ctx->bitmap_s1_b, 0, NULL, NULL) == -1) return -1; - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_bitmap_s1_c, CL_FALSE, 0, bitmap_ctx->bitmap_size, bitmap_ctx->bitmap_s1_c, 0, NULL, NULL) == -1) return -1; - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_bitmap_s1_d, CL_FALSE, 0, bitmap_ctx->bitmap_size, bitmap_ctx->bitmap_s1_d, 0, NULL, NULL) == -1) return -1; - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_bitmap_s2_a, CL_FALSE, 0, bitmap_ctx->bitmap_size, bitmap_ctx->bitmap_s2_a, 0, NULL, NULL) == -1) return -1; - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_bitmap_s2_b, CL_FALSE, 0, bitmap_ctx->bitmap_size, bitmap_ctx->bitmap_s2_b, 0, NULL, NULL) == -1) return -1; - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_bitmap_s2_c, CL_FALSE, 0, bitmap_ctx->bitmap_size, bitmap_ctx->bitmap_s2_c, 0, NULL, NULL) == -1) return -1; - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_bitmap_s2_d, CL_FALSE, 0, bitmap_ctx->bitmap_size, bitmap_ctx->bitmap_s2_d, 0, NULL, NULL) == -1) return -1; - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_digests_buf, CL_FALSE, 0, size_digests, hashes->digests_buf, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_bitmap_s1_a, CL_TRUE, 0, bitmap_ctx->bitmap_size, bitmap_ctx->bitmap_s1_a, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_bitmap_s1_b, CL_TRUE, 0, bitmap_ctx->bitmap_size, bitmap_ctx->bitmap_s1_b, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_bitmap_s1_c, CL_TRUE, 0, bitmap_ctx->bitmap_size, bitmap_ctx->bitmap_s1_c, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_bitmap_s1_d, CL_TRUE, 0, bitmap_ctx->bitmap_size, bitmap_ctx->bitmap_s1_d, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_bitmap_s2_a, CL_TRUE, 0, bitmap_ctx->bitmap_size, bitmap_ctx->bitmap_s2_a, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_bitmap_s2_b, CL_TRUE, 0, bitmap_ctx->bitmap_size, bitmap_ctx->bitmap_s2_b, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_bitmap_s2_c, CL_TRUE, 0, bitmap_ctx->bitmap_size, bitmap_ctx->bitmap_s2_c, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_bitmap_s2_d, CL_TRUE, 0, bitmap_ctx->bitmap_size, bitmap_ctx->bitmap_s2_d, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_digests_buf, CL_TRUE, 0, size_digests, hashes->digests_buf, 0, NULL, NULL) == -1) return -1; if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_salt_bufs, CL_FALSE, 0, size_salts, hashes->salts_buf, 0, NULL, NULL) == -1) return -1; /** @@ -12151,7 +12221,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (hc_clCreateBuffer (hashcat_ctx, device_param->opencl_context, CL_MEM_READ_ONLY, size_rules, NULL, &device_param->opencl_d_rules) == -1) return -1; if (hc_clCreateBuffer (hashcat_ctx, device_param->opencl_context, CL_MEM_READ_ONLY, size_rules_c, NULL, &device_param->opencl_d_rules_c) == -1) return -1; - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_rules, CL_FALSE, 0, size_rules_src, straight_ctx->kernel_rules_buf, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_rules, CL_TRUE, 0, size_rules_src, straight_ctx->kernel_rules_buf, 0, NULL, NULL) == -1) return -1; } else if (user_options_extra->attack_kern == ATTACK_KERN_COMBI) { @@ -12174,7 +12244,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) { if (hc_clCreateBuffer (hashcat_ctx, device_param->opencl_context, CL_MEM_READ_ONLY, size_esalts, NULL, &device_param->opencl_d_esalt_bufs) == -1) return -1; - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_esalt_bufs, CL_FALSE, 0, size_esalts, hashes->esalts_buf, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_esalt_bufs, CL_TRUE, 0, size_esalts, hashes->esalts_buf, 0, NULL, NULL) == -1) return -1; } if (hashconfig->st_hash != NULL) @@ -12186,7 +12256,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) { if (hc_clCreateBuffer (hashcat_ctx, device_param->opencl_context, CL_MEM_READ_ONLY, size_st_esalts, NULL, &device_param->opencl_d_st_esalts_buf) == -1) return -1; - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_st_esalts_buf, CL_FALSE, 0, size_st_esalts, hashes->st_esalts_buf, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_st_esalts_buf, CL_TRUE, 0, size_st_esalts, hashes->st_esalts_buf, 0, NULL, NULL) == -1) return -1; } } @@ -15948,6 +16018,12 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) { u32 threads_per_block_with_regs = (floor) ((float) device_param->regsPerBlock / num_regs); + if (threads_per_block_with_regs == 0) + { + // prevent threads_per_block from resulting in 0 due to a bug on the runtime + threads_per_block_with_regs = threads_per_block; + } + if (threads_per_block_with_regs > device_param->kernel_preferred_wgs_multiple) threads_per_block_with_regs -= threads_per_block_with_regs % device_param->kernel_preferred_wgs_multiple; threads_per_block = MIN (threads_per_block, threads_per_block_with_regs); @@ -15967,6 +16043,14 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) { u32 threads_per_block_with_regs = (floor) ((float) device_param->regsPerBlock / num_regs); + if (threads_per_block_with_regs == 0) + { + // https://rocm.docs.amd.com/projects/HIP/en/docs-develop/doxygen/html/bug.html + // HIP-Clang always returns 0 for regsPerBlock due to a known bug + // prevent threads_per_block from resulting in 0, otherwise hashcat crashes + threads_per_block_with_regs = threads_per_block; + } + if (threads_per_block_with_regs > device_param->kernel_preferred_wgs_multiple) threads_per_block_with_regs -= threads_per_block_with_regs % device_param->kernel_preferred_wgs_multiple; threads_per_block = MIN (threads_per_block, threads_per_block_with_regs); @@ -16045,7 +16129,8 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) // device_param->kernel_threads = kernel_threads; device_param->kernel_threads = 0; - u32 hardware_power_max = ((hashconfig->opts_type & OPTS_TYPE_MP_MULTI_DISABLE) ? 1 : device_processors) * device_param->kernel_threads_max; + const u32 hardware_power_max = ((hashconfig->opts_type & OPTS_TYPE_MP_MULTI_DISABLE) ? 1 : device_param->device_processors) + * ((hashconfig->opts_type & OPTS_TYPE_THREAD_MULTI_DISABLE) ? 1 : device_param->kernel_threads_max); u32 kernel_accel_min = device_param->kernel_accel_min; u32 kernel_accel_max = device_param->kernel_accel_max; @@ -17452,14 +17537,14 @@ int backend_session_update_mp (hashcat_ctx_t *hashcat_ctx) if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_root_css_buf, mask_ctx->root_css_buf, device_param->size_root_css, device_param->cuda_stream) == -1) return -1; - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_markov_css_buf, mask_ctx->markov_css_buf, device_param->size_markov_css, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_root_css_buf, mask_ctx->root_css_buf, device_param->size_root_css) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_markov_css_buf, mask_ctx->markov_css_buf, device_param->size_markov_css) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_root_css_buf, mask_ctx->root_css_buf, device_param->size_root_css, device_param->hip_stream) == -1) return -1; - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_markov_css_buf, mask_ctx->markov_css_buf, device_param->size_markov_css, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_root_css_buf, mask_ctx->root_css_buf, device_param->size_root_css) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_markov_css_buf, mask_ctx->markov_css_buf, device_param->size_markov_css) == -1) return -1; } #if defined (__APPLE__) @@ -17473,7 +17558,7 @@ int backend_session_update_mp (hashcat_ctx_t *hashcat_ctx) if (device_param->is_opencl == true) { if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_root_css_buf, CL_FALSE, 0, device_param->size_root_css, mask_ctx->root_css_buf, 0, NULL, NULL) == -1) return -1; - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_markov_css_buf, CL_FALSE, 0, device_param->size_markov_css, mask_ctx->markov_css_buf, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_markov_css_buf, CL_TRUE, 0, device_param->size_markov_css, mask_ctx->markov_css_buf, 0, NULL, NULL) == -1) return -1; if (hc_clFlush (hashcat_ctx, device_param->opencl_command_queue) == -1) return -1; } @@ -17508,14 +17593,14 @@ int backend_session_update_mp_rl (hashcat_ctx_t *hashcat_ctx, const u32 css_cnt_ if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_root_css_buf, mask_ctx->root_css_buf, device_param->size_root_css, device_param->cuda_stream) == -1) return -1; - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_markov_css_buf, mask_ctx->markov_css_buf, device_param->size_markov_css, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_root_css_buf, mask_ctx->root_css_buf, device_param->size_root_css) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_markov_css_buf, mask_ctx->markov_css_buf, device_param->size_markov_css) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_root_css_buf, mask_ctx->root_css_buf, device_param->size_root_css, device_param->hip_stream) == -1) return -1; - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_markov_css_buf, mask_ctx->markov_css_buf, device_param->size_markov_css, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_root_css_buf, mask_ctx->root_css_buf, device_param->size_root_css) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_markov_css_buf, mask_ctx->markov_css_buf, device_param->size_markov_css) == -1) return -1; } #if defined (__APPLE__) @@ -17529,7 +17614,7 @@ int backend_session_update_mp_rl (hashcat_ctx_t *hashcat_ctx, const u32 css_cnt_ if (device_param->is_opencl == true) { if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_root_css_buf, CL_FALSE, 0, device_param->size_root_css, mask_ctx->root_css_buf, 0, NULL, NULL) == -1) return -1; - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_markov_css_buf, CL_FALSE, 0, device_param->size_markov_css, mask_ctx->markov_css_buf, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_markov_css_buf, CL_TRUE, 0, device_param->size_markov_css, mask_ctx->markov_css_buf, 0, NULL, NULL) == -1) return -1; if (hc_clFlush (hashcat_ctx, device_param->opencl_command_queue) == -1) return -1; } diff --git a/src/ext_cuda.c b/src/ext_cuda.c index 4c0db9ffb..878977ade 100644 --- a/src/ext_cuda.c +++ b/src/ext_cuda.c @@ -87,14 +87,19 @@ int cuda_init (void *hashcat_ctx) HC_LOAD_FUNC_CUDA (cuda, cuLaunchKernel, cuLaunchKernel, CUDA_CULAUNCHKERNEL, CUDA, 1); HC_LOAD_FUNC_CUDA (cuda, cuMemAlloc, cuMemAlloc_v2, CUDA_CUMEMALLOC, CUDA, 1); HC_LOAD_FUNC_CUDA (cuda, cuMemAllocHost, cuMemAllocHost_v2, CUDA_CUMEMALLOCHOST, CUDA, 1); + HC_LOAD_FUNC_CUDA (cuda, cuMemcpyDtoD, cuMemcpyDtoD_v2, CUDA_CUMEMCPYDTOD, CUDA, 1); + HC_LOAD_FUNC_CUDA (cuda, cuMemcpyDtoH, cuMemcpyDtoH_v2, CUDA_CUMEMCPYDTOH, CUDA, 1); + HC_LOAD_FUNC_CUDA (cuda, cuMemcpyHtoD, cuMemcpyHtoD_v2, CUDA_CUMEMCPYHTOD, CUDA, 1); + HC_LOAD_FUNC_CUDA (cuda, cuMemsetD32, cuMemsetD32, CUDA_CUMEMSETD32, CUDA, 1); + HC_LOAD_FUNC_CUDA (cuda, cuMemsetD8, cuMemsetD8, CUDA_CUMEMSETD8, CUDA, 1); HC_LOAD_FUNC_CUDA (cuda, cuMemcpyDtoDAsync, cuMemcpyDtoDAsync_v2, CUDA_CUMEMCPYDTODASYNC, CUDA, 1); HC_LOAD_FUNC_CUDA (cuda, cuMemcpyDtoHAsync, cuMemcpyDtoHAsync_v2, CUDA_CUMEMCPYDTOHASYNC, CUDA, 1); HC_LOAD_FUNC_CUDA (cuda, cuMemcpyHtoDAsync, cuMemcpyHtoDAsync_v2, CUDA_CUMEMCPYHTODASYNC, CUDA, 1); + HC_LOAD_FUNC_CUDA (cuda, cuMemsetD32Async, cuMemsetD32Async, CUDA_CUMEMSETD32ASYNC, CUDA, 1); + HC_LOAD_FUNC_CUDA (cuda, cuMemsetD8Async, cuMemsetD8Async, CUDA_CUMEMSETD8ASYNC, CUDA, 1); HC_LOAD_FUNC_CUDA (cuda, cuMemFree, cuMemFree_v2, CUDA_CUMEMFREE, CUDA, 1); HC_LOAD_FUNC_CUDA (cuda, cuMemFreeHost, cuMemFreeHost, CUDA_CUMEMFREEHOST, CUDA, 1); HC_LOAD_FUNC_CUDA (cuda, cuMemGetInfo, cuMemGetInfo_v2, CUDA_CUMEMGETINFO, CUDA, 1); - HC_LOAD_FUNC_CUDA (cuda, cuMemsetD32Async, cuMemsetD32Async, CUDA_CUMEMSETD32ASYNC, CUDA, 1); - HC_LOAD_FUNC_CUDA (cuda, cuMemsetD8Async, cuMemsetD8Async, CUDA_CUMEMSETD8ASYNC, CUDA, 1); HC_LOAD_FUNC_CUDA (cuda, cuModuleGetFunction, cuModuleGetFunction, CUDA_CUMODULEGETFUNCTION, CUDA, 1); HC_LOAD_FUNC_CUDA (cuda, cuModuleGetGlobal, cuModuleGetGlobal_v2, CUDA_CUMODULEGETGLOBAL, CUDA, 1); HC_LOAD_FUNC_CUDA (cuda, cuModuleLoad, cuModuleLoad, CUDA_CUMODULELOAD, CUDA, 1); @@ -517,6 +522,142 @@ int hc_cuMemFree (void *hashcat_ctx, CUdeviceptr dptr) return 0; } + +int hc_cuMemcpyDtoH (void *hashcat_ctx, void *dstHost, CUdeviceptr srcDevice, size_t ByteCount) +{ + backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; + + CUDA_PTR *cuda = (CUDA_PTR *) backend_ctx->cuda; + + const CUresult CU_err = cuda->cuMemcpyDtoH (dstHost, srcDevice, ByteCount); + + if (CU_err != CUDA_SUCCESS) + { + const char *pStr = NULL; + + if (cuda->cuGetErrorString (CU_err, &pStr) == CUDA_SUCCESS) + { + event_log_error (hashcat_ctx, "cuMemcpyDtoH(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "cuMemcpyDtoH(): %d", CU_err); + } + + return -1; + } + + return 0; +} + +int hc_cuMemcpyDtoD (void *hashcat_ctx, CUdeviceptr dstDevice, CUdeviceptr srcDevice, size_t ByteCount) +{ + backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; + + CUDA_PTR *cuda = (CUDA_PTR *) backend_ctx->cuda; + + const CUresult CU_err = cuda->cuMemcpyDtoD (dstDevice, srcDevice, ByteCount); + + if (CU_err != CUDA_SUCCESS) + { + const char *pStr = NULL; + + if (cuda->cuGetErrorString (CU_err, &pStr) == CUDA_SUCCESS) + { + event_log_error (hashcat_ctx, "cuMemcpyDtoD(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "cuMemcpyDtoD(): %d", CU_err); + } + + return -1; + } + + return 0; +} + +int hc_cuMemcpyHtoD (void *hashcat_ctx, CUdeviceptr dstDevice, const void *srcHost, size_t ByteCount) +{ + backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; + + CUDA_PTR *cuda = (CUDA_PTR *) backend_ctx->cuda; + + const CUresult CU_err = cuda->cuMemcpyHtoD (dstDevice, srcHost, ByteCount); + + if (CU_err != CUDA_SUCCESS) + { + const char *pStr = NULL; + + if (cuda->cuGetErrorString (CU_err, &pStr) == CUDA_SUCCESS) + { + event_log_error (hashcat_ctx, "cuMemcpyHtoD(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "cuMemcpyHtoD(): %d", CU_err); + } + + return -1; + } + + return 0; +} + +int hc_cuMemsetD32 (void *hashcat_ctx, CUdeviceptr dstDevice, unsigned int ui, size_t N) +{ + backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; + + CUDA_PTR *cuda = (CUDA_PTR *) backend_ctx->cuda; + + const CUresult CU_err = cuda->cuMemsetD32 (dstDevice, ui, N); + + if (CU_err != CUDA_SUCCESS) + { + const char *pStr = NULL; + + if (cuda->cuGetErrorString (CU_err, &pStr) == CUDA_SUCCESS) + { + event_log_error (hashcat_ctx, "cuMemsetD32(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "cuMemsetD32(): %d", CU_err); + } + + return -1; + } + + return 0; +} + +int hc_cuMemsetD8 (void *hashcat_ctx, CUdeviceptr dstDevice, unsigned char uc, size_t N) +{ + backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; + + CUDA_PTR *cuda = (CUDA_PTR *) backend_ctx->cuda; + + const CUresult CU_err = cuda->cuMemsetD8 (dstDevice, uc, N); + + if (CU_err != CUDA_SUCCESS) + { + const char *pStr = NULL; + + if (cuda->cuGetErrorString (CU_err, &pStr) == CUDA_SUCCESS) + { + event_log_error (hashcat_ctx, "cuMemsetD8(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "cuMemsetD8(): %d", CU_err); + } + + return -1; + } + + return 0; +} + int hc_cuMemcpyDtoHAsync (void *hashcat_ctx, void *dstHost, CUdeviceptr srcDevice, size_t ByteCount, CUstream hStream) { backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; diff --git a/src/ext_hip.c b/src/ext_hip.c index decfa987c..4a26a7e9f 100644 --- a/src/ext_hip.c +++ b/src/ext_hip.c @@ -140,12 +140,16 @@ int hip_init (void *hashcat_ctx) HC_LOAD_FUNC_HIP (hip, hipMemAlloc, hipMalloc, HIP_HIPMEMALLOC, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipMemFree, hipFree, HIP_HIPMEMFREE, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipMemGetInfo, hipMemGetInfo, HIP_HIPMEMGETINFO, HIP, 1); + HC_LOAD_FUNC_HIP (hip, hipMemcpyDtoD, hipMemcpyDtoD, HIP_HIPMEMCPYDTOD, HIP, 1); + HC_LOAD_FUNC_HIP (hip, hipMemcpyDtoH, hipMemcpyDtoH, HIP_HIPMEMCPYDTOH, HIP, 1); + HC_LOAD_FUNC_HIP (hip, hipMemcpyHtoD, hipMemcpyHtoD, HIP_HIPMEMCPYHTOD, HIP, 1); + HC_LOAD_FUNC_HIP (hip, hipMemsetD32, hipMemsetD32, HIP_HIPMEMSETD32, HIP, 1); + HC_LOAD_FUNC_HIP (hip, hipMemsetD8, hipMemsetD8, HIP_HIPMEMSETD8, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipMemcpyDtoDAsync, hipMemcpyDtoDAsync, HIP_HIPMEMCPYDTODASYNC, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipMemcpyDtoHAsync, hipMemcpyDtoHAsync, HIP_HIPMEMCPYDTOHASYNC, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipMemcpyHtoDAsync, hipMemcpyHtoDAsync, HIP_HIPMEMCPYHTODASYNC, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipMemsetD32Async, hipMemsetD32Async, HIP_HIPMEMSETD32ASYNC, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipMemsetD8Async, hipMemsetD8Async, HIP_HIPMEMSETD8ASYNC, HIP, 1); - HC_LOAD_FUNC_HIP (hip, hipMemcpyHtoDAsync, hipMemcpyHtoDAsync, HIP_HIPMEMCPYHTODASYNC, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipModuleGetFunction, hipModuleGetFunction, HIP_HIPMODULEGETFUNCTION, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipModuleGetGlobal, hipModuleGetGlobal, HIP_HIPMODULEGETGLOBAL, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipModuleLoadDataEx, hipModuleLoadDataEx, HIP_HIPMODULELOADDATAEX, HIP, 1); @@ -800,6 +804,143 @@ int hc_hipMemGetInfo (void *hashcat_ctx, size_t *free, size_t *total) return 0; } + + +int hc_hipMemcpyDtoH (void *hashcat_ctx, void *dstHost, hipDeviceptr_t srcDevice, size_t ByteCount) +{ + backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const hipError_t HIP_err = hip->hipMemcpyDtoH (dstHost, srcDevice, ByteCount); + + if (HIP_err != hipSuccess) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) + { + event_log_error (hashcat_ctx, "hipMemcpyDtoH(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipMemcpyDtoH(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipMemcpyDtoD (void *hashcat_ctx, hipDeviceptr_t dstDevice, hipDeviceptr_t srcDevice, size_t ByteCount) +{ + backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const hipError_t HIP_err = hip->hipMemcpyDtoD (dstDevice, srcDevice, ByteCount); + + if (HIP_err != hipSuccess) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) + { + event_log_error (hashcat_ctx, "hipMemcpyDtoD(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipMemcpyDtoD(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipMemcpyHtoD (void *hashcat_ctx, hipDeviceptr_t dstDevice, const void *srcHost, size_t ByteCount) +{ + backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const hipError_t HIP_err = hip->hipMemcpyHtoD (dstDevice, srcHost, ByteCount); + + if (HIP_err != hipSuccess) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) + { + event_log_error (hashcat_ctx, "hipMemcpyHtoD(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipMemcpyHtoD(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipMemsetD32 (void *hashcat_ctx, hipDeviceptr_t dstDevice, unsigned int ui, size_t N) +{ + backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const hipError_t HIP_err = hip->hipMemsetD32 (dstDevice, ui, N); + + if (HIP_err != hipSuccess) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) + { + event_log_error (hashcat_ctx, "hipMemsetD32(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipMemsetD32(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipMemsetD8 (void *hashcat_ctx, hipDeviceptr_t dstDevice, unsigned char uc, size_t N) +{ + backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const hipError_t HIP_err = hip->hipMemsetD8 (dstDevice, uc, N); + + if (HIP_err != hipSuccess) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) + { + event_log_error (hashcat_ctx, "hipMemsetD8(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipMemsetD8(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + int hc_hipMemcpyDtoHAsync (void *hashcat_ctx, void *dstHost, hipDeviceptr_t srcDevice, size_t ByteCount, hipStream_t hStream) { backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; diff --git a/src/hashes.c b/src/hashes.c index 78a6d5b90..72ab14433 100644 --- a/src/hashes.c +++ b/src/hashes.c @@ -334,7 +334,7 @@ int check_hash (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, pla if (device_param->is_cuda == true) { - rc = hc_cuMemcpyDtoHAsync (hashcat_ctx, tmps, device_param->cuda_d_tmps + (plain->gidvid * hashconfig->tmp_size), hashconfig->tmp_size, device_param->cuda_stream); + rc = hc_cuMemcpyDtoH (hashcat_ctx, tmps, device_param->cuda_d_tmps + (plain->gidvid * hashconfig->tmp_size), hashconfig->tmp_size); if (rc == 0) { @@ -351,7 +351,7 @@ int check_hash (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, pla if (device_param->is_hip == true) { - rc = hc_hipMemcpyDtoHAsync (hashcat_ctx, tmps, device_param->hip_d_tmps + (plain->gidvid * hashconfig->tmp_size), hashconfig->tmp_size, device_param->hip_stream); + rc = hc_hipMemcpyDtoH (hashcat_ctx, tmps, device_param->hip_d_tmps + (plain->gidvid * hashconfig->tmp_size), hashconfig->tmp_size); if (rc == 0) { @@ -382,7 +382,7 @@ int check_hash (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, pla if (device_param->is_opencl == true) { - rc = hc_clEnqueueReadBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_tmps, CL_FALSE, plain->gidvid * hashconfig->tmp_size, hashconfig->tmp_size, tmps, 0, NULL, &opencl_event); + rc = hc_clEnqueueReadBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_tmps, CL_TRUE, plain->gidvid * hashconfig->tmp_size, hashconfig->tmp_size, tmps, 0, NULL, &opencl_event); if (rc == 0) { @@ -587,14 +587,14 @@ int check_cracked (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param) if (device_param->is_cuda == true) { - if (hc_cuMemcpyDtoHAsync (hashcat_ctx, &num_cracked, device_param->cuda_d_result, sizeof (u32), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyDtoH (hashcat_ctx, &num_cracked, device_param->cuda_d_result, sizeof (u32)) == -1) return -1; if (hc_cuStreamSynchronize (hashcat_ctx, device_param->cuda_stream) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyDtoHAsync (hashcat_ctx, &num_cracked, device_param->hip_d_result, sizeof (u32), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyDtoH (hashcat_ctx, &num_cracked, device_param->hip_d_result, sizeof (u32)) == -1) return -1; if (hc_hipStreamSynchronize (hashcat_ctx, device_param->hip_stream) == -1) return -1; } @@ -624,7 +624,7 @@ int check_cracked (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param) if (device_param->is_cuda == true) { - rc = hc_cuMemcpyDtoHAsync (hashcat_ctx, cracked, device_param->cuda_d_plain_bufs, num_cracked * sizeof (plain_t), device_param->cuda_stream); + rc = hc_cuMemcpyDtoH (hashcat_ctx, cracked, device_param->cuda_d_plain_bufs, num_cracked * sizeof (plain_t)); if (rc == 0) { @@ -641,7 +641,7 @@ int check_cracked (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param) if (device_param->is_hip == true) { - rc = hc_hipMemcpyDtoHAsync (hashcat_ctx, cracked, device_param->hip_d_plain_bufs, num_cracked * sizeof (plain_t), device_param->hip_stream); + rc = hc_hipMemcpyDtoH (hashcat_ctx, cracked, device_param->hip_d_plain_bufs, num_cracked * sizeof (plain_t)); if (rc == 0) { diff --git a/src/modules/module_01300.c b/src/modules/module_01300.c index b50a0e99c..c79a3493c 100644 --- a/src/modules/module_01300.c +++ b/src/modules/module_01300.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_RAW_HASH; static const char *HASH_NAME = "SHA2-224"; static const u64 KERN_TYPE = 1300; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_PRECOMPUTE_INIT | OPTI_TYPE_EARLY_SKIP | OPTI_TYPE_NOT_ITERATED diff --git a/src/modules/module_01800.c b/src/modules/module_01800.c index fea4851c6..f0d2e9623 100644 --- a/src/modules/module_01800.c +++ b/src/modules/module_01800.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_OS; static const char *HASH_NAME = "sha512crypt $6$, SHA512 (Unix)"; static const u64 KERN_TYPE = 1800; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_USES_BITS_64; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE; diff --git a/src/modules/module_06400.c b/src/modules/module_06400.c index b7ac9d87f..a553846c7 100644 --- a/src/modules/module_06400.c +++ b/src/modules/module_06400.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_OS; static const char *HASH_NAME = "AIX {ssha256}"; static const u64 KERN_TYPE = 6400; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE; diff --git a/src/modules/module_06800.c b/src/modules/module_06800.c index 215f26992..e6cf96fd6 100644 --- a/src/modules/module_06800.c +++ b/src/modules/module_06800.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_PASSWORD_MANAGER; static const char *HASH_NAME = "LastPass + LastPass sniffed"; static const u64 KERN_TYPE = 6800; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE; diff --git a/src/modules/module_07100.c b/src/modules/module_07100.c index ec3924cec..78c8f8081 100644 --- a/src/modules/module_07100.c +++ b/src/modules/module_07100.c @@ -85,15 +85,6 @@ u32 module_pw_max (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED con return pw_max; } -char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra, MAYBE_UNUSED const hashes_t *hashes, MAYBE_UNUSED const hc_device_param_t *device_param) -{ - char *jit_build_options = NULL; - - hc_asprintf (&jit_build_options, "-D NO_UNROLL"); - - return jit_build_options; -} - int module_hash_decode (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED void *digest_buf, MAYBE_UNUSED salt_t *salt, MAYBE_UNUSED void *esalt_buf, MAYBE_UNUSED void *hook_salt_buf, MAYBE_UNUSED hashinfo_t *hash_info, const char *line_buf, MAYBE_UNUSED const int line_len) { u64 *digest = (u64 *) digest_buf; @@ -398,7 +389,7 @@ void module_init (module_ctx_t *module_ctx) module_ctx->module_hook23 = MODULE_DEFAULT; module_ctx->module_hook_salt_size = MODULE_DEFAULT; module_ctx->module_hook_size = MODULE_DEFAULT; - module_ctx->module_jit_build_options = module_jit_build_options; + module_ctx->module_jit_build_options = MODULE_DEFAULT; module_ctx->module_jit_cache_disable = MODULE_DEFAULT; module_ctx->module_kernel_accel_max = MODULE_DEFAULT; module_ctx->module_kernel_accel_min = MODULE_DEFAULT; diff --git a/src/modules/module_07700.c b/src/modules/module_07700.c index 2642a0bb7..1d783d351 100644 --- a/src/modules/module_07700.c +++ b/src/modules/module_07700.c @@ -21,6 +21,7 @@ static const char *HASH_NAME = "SAP CODVN B (BCODE)"; static const u64 KERN_TYPE = 7700; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE | OPTI_TYPE_PRECOMPUTE_INIT + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_NOT_ITERATED; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE diff --git a/src/modules/module_07701.c b/src/modules/module_07701.c index bc852639b..232d7ca4d 100644 --- a/src/modules/module_07701.c +++ b/src/modules/module_07701.c @@ -21,6 +21,7 @@ static const char *HASH_NAME = "SAP CODVN B (BCODE) from RFC_READ_TABLE"; static const u64 KERN_TYPE = 7701; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE | OPTI_TYPE_PRECOMPUTE_INIT + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_NOT_ITERATED; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE diff --git a/src/modules/module_08200.c b/src/modules/module_08200.c index 54d5c62b5..90049b1ed 100644 --- a/src/modules/module_08200.c +++ b/src/modules/module_08200.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_PASSWORD_MANAGER; static const char *HASH_NAME = "1Password, cloudkeychain"; static const u64 KERN_TYPE = 8200; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_USES_BITS_64 | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE diff --git a/src/modules/module_08300.c b/src/modules/module_08300.c index c499667ae..80dd3475d 100644 --- a/src/modules/module_08300.c +++ b/src/modules/module_08300.c @@ -19,7 +19,8 @@ static const u32 DGST_SIZE = DGST_SIZE_4_5; static const u32 HASH_CATEGORY = HASH_CATEGORY_NETWORK_SERVER; static const char *HASH_NAME = "DNSSEC (NSEC3)"; static const u64 KERN_TYPE = 8300; -static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE; +static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_BE | OPTS_TYPE_ST_HEX diff --git a/src/modules/module_13751.c b/src/modules/module_13751.c index dac31a633..ee27acfd6 100644 --- a/src/modules/module_13751.c +++ b/src/modules/module_13751.c @@ -23,7 +23,8 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "VeraCrypt SHA256 + XTS 512 bit (legacy)"; static const u64 KERN_TYPE = 13751; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE - | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; + | OPTI_TYPE_SLOW_HASH_SIMD_LOOP + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_BINARY_HASHFILE diff --git a/src/modules/module_13752.c b/src/modules/module_13752.c index e959a80af..095758dd2 100644 --- a/src/modules/module_13752.c +++ b/src/modules/module_13752.c @@ -23,7 +23,8 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "VeraCrypt SHA256 + XTS 1024 bit (legacy)"; static const u64 KERN_TYPE = 13752; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE - | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; + | OPTI_TYPE_SLOW_HASH_SIMD_LOOP + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_BINARY_HASHFILE diff --git a/src/modules/module_13753.c b/src/modules/module_13753.c index 277afaf02..4eee98625 100644 --- a/src/modules/module_13753.c +++ b/src/modules/module_13753.c @@ -23,7 +23,8 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "VeraCrypt SHA256 + XTS 1536 bit (legacy)"; static const u64 KERN_TYPE = 13753; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE - | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; + | OPTI_TYPE_SLOW_HASH_SIMD_LOOP + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_BINARY_HASHFILE diff --git a/src/modules/module_13761.c b/src/modules/module_13761.c index e33a693e2..5b1b82b27 100644 --- a/src/modules/module_13761.c +++ b/src/modules/module_13761.c @@ -23,7 +23,8 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "VeraCrypt SHA256 + XTS 512 bit + boot-mode (legacy)"; static const u64 KERN_TYPE = 13751; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE - | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; + | OPTI_TYPE_SLOW_HASH_SIMD_LOOP + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_BINARY_HASHFILE diff --git a/src/modules/module_13762.c b/src/modules/module_13762.c index 0f5f18545..6f1a27929 100644 --- a/src/modules/module_13762.c +++ b/src/modules/module_13762.c @@ -23,7 +23,8 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "VeraCrypt SHA256 + XTS 1024 bit + boot-mode (legacy)"; static const u64 KERN_TYPE = 13752; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE - | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; + | OPTI_TYPE_SLOW_HASH_SIMD_LOOP + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_BINARY_HASHFILE diff --git a/src/modules/module_13763.c b/src/modules/module_13763.c index c9e28d4c7..fb50002a2 100644 --- a/src/modules/module_13763.c +++ b/src/modules/module_13763.c @@ -23,7 +23,8 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "VeraCrypt SHA256 + XTS 1536 bit + boot-mode (legacy)"; static const u64 KERN_TYPE = 13753; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE - | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; + | OPTI_TYPE_SLOW_HASH_SIMD_LOOP + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_BINARY_HASHFILE diff --git a/src/modules/module_14800.c b/src/modules/module_14800.c index 9f3c1ca0f..1e4b91a15 100644 --- a/src/modules/module_14800.c +++ b/src/modules/module_14800.c @@ -21,6 +21,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_ARCHIVE; static const char *HASH_NAME = "iTunes backup >= 10.0"; static const u64 KERN_TYPE = 14800; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP | OPTI_TYPE_SLOW_HASH_SIMD_LOOP2; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE diff --git a/src/modules/module_14900.c b/src/modules/module_14900.c index cbe009024..fcc639688 100644 --- a/src/modules/module_14900.c +++ b/src/modules/module_14900.c @@ -19,7 +19,8 @@ static const u32 DGST_SIZE = DGST_SIZE_4_4; static const u32 HASH_CATEGORY = HASH_CATEGORY_RAW_CIPHER_KPA; static const char *HASH_NAME = "Skip32 (PT = $salt, key = $pass)"; static const u64 KERN_TYPE = 14900; -static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE; +static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_SUGGEST_KG; diff --git a/src/modules/module_15600.c b/src/modules/module_15600.c index 1f2b7b9a3..66469d9bf 100644 --- a/src/modules/module_15600.c +++ b/src/modules/module_15600.c @@ -21,6 +21,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_CRYPTOCURRENCY_WALLET; static const char *HASH_NAME = "Ethereum Wallet, PBKDF2-HMAC-SHA256"; static const u64 KERN_TYPE = 15600; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE diff --git a/src/modules/module_16300.c b/src/modules/module_16300.c index e7507f1fa..6a12a1b39 100644 --- a/src/modules/module_16300.c +++ b/src/modules/module_16300.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_CRYPTOCURRENCY_WALLET; static const char *HASH_NAME = "Ethereum Pre-Sale Wallet, PBKDF2-HMAC-SHA256"; static const u64 KERN_TYPE = 16300; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE diff --git a/src/modules/module_16700.c b/src/modules/module_16700.c index 9dce98e62..2fec2f557 100644 --- a/src/modules/module_16700.c +++ b/src/modules/module_16700.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "FileVault 2"; static const u64 KERN_TYPE = 16200; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE; diff --git a/src/modules/module_16900.c b/src/modules/module_16900.c index 443da6007..70eca3f16 100644 --- a/src/modules/module_16900.c +++ b/src/modules/module_16900.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_PASSWORD_MANAGER; static const char *HASH_NAME = "Ansible Vault"; static const u64 KERN_TYPE = 16900; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE; diff --git a/src/modules/module_18100.c b/src/modules/module_18100.c index 46dadfe60..678cb8265 100644 --- a/src/modules/module_18100.c +++ b/src/modules/module_18100.c @@ -21,6 +21,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_OTP; static const char *HASH_NAME = "TOTP (HMAC-SHA1)"; static const u64 KERN_TYPE = 18100; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_NOT_ITERATED; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_BE diff --git a/src/modules/module_20800.c b/src/modules/module_20800.c index 6cece615d..1330b99bc 100644 --- a/src/modules/module_20800.c +++ b/src/modules/module_20800.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_RAW_HASH_SALTED; static const char *HASH_NAME = "sha256(md5($pass))"; static const u64 KERN_TYPE = 20800; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_PRECOMPUTE_INIT | OPTI_TYPE_EARLY_SKIP | OPTI_TYPE_NOT_ITERATED diff --git a/src/modules/module_21300.c b/src/modules/module_21300.c index 75e57c1f4..34ebff7f5 100644 --- a/src/modules/module_21300.c +++ b/src/modules/module_21300.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_RAW_HASH_SALTED; static const char *HASH_NAME = "md5($salt.sha1($salt.$pass))"; static const u64 KERN_TYPE = 21300; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_PRECOMPUTE_INIT | OPTI_TYPE_EARLY_SKIP | OPTI_TYPE_NOT_ITERATED diff --git a/src/modules/module_21500.c b/src/modules/module_21500.c index dba580a7c..8936be11c 100644 --- a/src/modules/module_21500.c +++ b/src/modules/module_21500.c @@ -21,6 +21,7 @@ static const char *HASH_NAME = "SolarWinds Orion"; static const u64 KERN_TYPE = 21500; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE | OPTI_TYPE_USES_BITS_64 + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE; diff --git a/src/modules/module_21700.c b/src/modules/module_21700.c index 65536027b..e00cf232f 100644 --- a/src/modules/module_21700.c +++ b/src/modules/module_21700.c @@ -22,6 +22,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_CRYPTOCURRENCY_WALLET; static const char *HASH_NAME = "Electrum Wallet (Salt-Type 4)"; static const u64 KERN_TYPE = 21700; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_USES_BITS_64 | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE diff --git a/src/modules/module_22100.c b/src/modules/module_22100.c index 344218243..4d79d30bf 100644 --- a/src/modules/module_22100.c +++ b/src/modules/module_22100.c @@ -20,7 +20,8 @@ static const u32 DGST_SIZE = DGST_SIZE_4_4; static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "BitLocker"; static const u64 KERN_TYPE = 22100; -static const u32 OPTI_TYPE = OPTI_TYPE_SLOW_HASH_SIMD_LOOP; +static const u32 OPTI_TYPE = OPTI_TYPE_SLOW_HASH_SIMD_LOOP + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_MP_MULTI_DISABLE; diff --git a/src/modules/module_22300.c b/src/modules/module_22300.c index a85d1d753..04a4bc08c 100644 --- a/src/modules/module_22300.c +++ b/src/modules/module_22300.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_RAW_HASH_SALTED; static const char *HASH_NAME = "sha256($salt.$pass.$salt)"; static const u64 KERN_TYPE = 22300; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_PRECOMPUTE_INIT | OPTI_TYPE_EARLY_SKIP | OPTI_TYPE_NOT_ITERATED diff --git a/src/modules/module_22921.c b/src/modules/module_22921.c index 163aacee9..c6471a232 100644 --- a/src/modules/module_22921.c +++ b/src/modules/module_22921.c @@ -19,7 +19,8 @@ static const u32 DGST_SIZE = DGST_SIZE_4_4; static const u32 HASH_CATEGORY = HASH_CATEGORY_PRIVATE_KEY; static const char *HASH_NAME = "RSA/DSA/EC/OpenSSH Private Keys ($6$)"; static const u64 KERN_TYPE = 22921; -static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE; +static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE; static const u32 SALT_TYPE = SALT_TYPE_EMBEDDED; diff --git a/src/modules/module_22941.c b/src/modules/module_22941.c index 4b4bf09ac..65e7d7c2b 100644 --- a/src/modules/module_22941.c +++ b/src/modules/module_22941.c @@ -19,7 +19,8 @@ static const u32 DGST_SIZE = DGST_SIZE_4_4; static const u32 HASH_CATEGORY = HASH_CATEGORY_PRIVATE_KEY; static const char *HASH_NAME = "RSA/DSA/EC/OpenSSH Private Keys ($4$)"; static const u64 KERN_TYPE = 22941; -static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE; +static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE; static const u32 SALT_TYPE = SALT_TYPE_EMBEDDED; diff --git a/src/modules/module_23400.c b/src/modules/module_23400.c index 13743815f..c2cc2a2e2 100644 --- a/src/modules/module_23400.c +++ b/src/modules/module_23400.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_PASSWORD_MANAGER; static const char *HASH_NAME = "Bitwarden"; static const u64 KERN_TYPE = 23400; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE diff --git a/src/modules/module_23600.c b/src/modules/module_23600.c index fdb3467d1..0b8994ff7 100644 --- a/src/modules/module_23600.c +++ b/src/modules/module_23600.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_ARCHIVE; static const char *HASH_NAME = "AxCrypt 2 AES-256"; static const u64 KERN_TYPE = 23600; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_USES_BITS_64 | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE diff --git a/src/modules/module_23800.c b/src/modules/module_23800.c index c6c8fec61..c75a9aaaf 100644 --- a/src/modules/module_23800.c +++ b/src/modules/module_23800.c @@ -20,7 +20,8 @@ static const u32 DGST_SIZE = DGST_SIZE_4_4; // actually only DGST_SIZE_4_ static const u32 HASH_CATEGORY = HASH_CATEGORY_ARCHIVE; static const char *HASH_NAME = "RAR3-p (Compressed)"; static const u64 KERN_TYPE = 23800; -static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE; +static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_HOOK23 diff --git a/src/modules/module_24200.c b/src/modules/module_24200.c index 8b6a8e5c8..20ea6ccec 100644 --- a/src/modules/module_24200.c +++ b/src/modules/module_24200.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_DATABASE_SERVER; static const char *HASH_NAME = "MongoDB ServerKey SCRAM-SHA-256"; static const u64 KERN_TYPE = 24200; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE diff --git a/src/modules/module_24420.c b/src/modules/module_24420.c index 39820834a..544e07ee1 100644 --- a/src/modules/module_24420.c +++ b/src/modules/module_24420.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_PRIVATE_KEY; static const char *HASH_NAME = "PKCS#8 Private Keys (PBKDF2-HMAC-SHA256 + 3DES/AES)"; static const u64 KERN_TYPE = 24420; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE diff --git a/src/modules/module_25500.c b/src/modules/module_25500.c index afc451fb9..06a1e795d 100644 --- a/src/modules/module_25500.c +++ b/src/modules/module_25500.c @@ -21,6 +21,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_CRYPTOCURRENCY_WALLET; static const char *HASH_NAME = "Stargazer Stellar Wallet XLM"; static const u64 KERN_TYPE = 25500; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE; diff --git a/src/modules/module_25900.c b/src/modules/module_25900.c index 7ad951f6d..a460a3bd8 100644 --- a/src/modules/module_25900.c +++ b/src/modules/module_25900.c @@ -19,7 +19,8 @@ static const u32 DGST_SIZE = DGST_SIZE_4_4; static const u32 HASH_CATEGORY = HASH_CATEGORY_NETWORK_SERVER; static const char *HASH_NAME = "KNX IP Secure - Device Authentication Code"; static const u64 KERN_TYPE = 25900; -static const u32 OPTI_TYPE = OPTI_TYPE_SLOW_HASH_SIMD_LOOP; +static const u32 OPTI_TYPE = OPTI_TYPE_SLOW_HASH_SIMD_LOOP + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_DEEP_COMP_KERNEL; diff --git a/src/modules/module_26000.c b/src/modules/module_26000.c index 35a6dd937..8ed90d5ee 100644 --- a/src/modules/module_26000.c +++ b/src/modules/module_26000.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_PASSWORD_MANAGER; static const char *HASH_NAME = "Mozilla key3.db"; static const u64 KERN_TYPE = 26000; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_NOT_ITERATED; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_BE; diff --git a/src/modules/module_26100.c b/src/modules/module_26100.c index c42958289..3414fcbdb 100644 --- a/src/modules/module_26100.c +++ b/src/modules/module_26100.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_PASSWORD_MANAGER; static const char *HASH_NAME = "Mozilla key4.db"; static const u64 KERN_TYPE = 26100; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE; diff --git a/src/modules/module_26600.c b/src/modules/module_26600.c index 926555e94..92b5c8efc 100644 --- a/src/modules/module_26600.c +++ b/src/modules/module_26600.c @@ -21,6 +21,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_CRYPTOCURRENCY_WALLET; static const char *HASH_NAME = "MetaMask Wallet (needs all data, checks AES-GCM tag)"; static const u64 KERN_TYPE = 26600; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE; diff --git a/src/modules/module_26700.c b/src/modules/module_26700.c index 63854c0d2..47427ba32 100644 --- a/src/modules/module_26700.c +++ b/src/modules/module_26700.c @@ -21,7 +21,8 @@ static const u32 DGST_SIZE = DGST_SIZE_4_4; static const u32 HASH_CATEGORY = HASH_CATEGORY_NETWORK_PROTOCOL; static const char *HASH_NAME = "SNMPv3 HMAC-SHA224-128"; static const u64 KERN_TYPE = 26700; -static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE; +static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_NATIVE_THREADS | OPTS_TYPE_PT_GENERATE_LE; diff --git a/src/modules/module_26800.c b/src/modules/module_26800.c index 024bfa636..44d72f7df 100644 --- a/src/modules/module_26800.c +++ b/src/modules/module_26800.c @@ -21,7 +21,8 @@ static const u32 DGST_SIZE = DGST_SIZE_4_6; static const u32 HASH_CATEGORY = HASH_CATEGORY_NETWORK_PROTOCOL; static const char *HASH_NAME = "SNMPv3 HMAC-SHA256-192"; static const u64 KERN_TYPE = 26800; -static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE; +static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_NATIVE_THREADS | OPTS_TYPE_PT_GENERATE_LE; diff --git a/src/modules/module_26900.c b/src/modules/module_26900.c index 8d916ebf2..29e3cdc92 100644 --- a/src/modules/module_26900.c +++ b/src/modules/module_26900.c @@ -22,6 +22,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_NETWORK_PROTOCOL; static const char *HASH_NAME = "SNMPv3 HMAC-SHA384-256"; static const u64 KERN_TYPE = 26900; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_USES_BITS_64; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_NATIVE_THREADS diff --git a/src/modules/module_27300.c b/src/modules/module_27300.c index ff20acc9d..4d607051e 100644 --- a/src/modules/module_27300.c +++ b/src/modules/module_27300.c @@ -22,6 +22,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_NETWORK_PROTOCOL; static const char *HASH_NAME = "SNMPv3 HMAC-SHA512-384"; static const u64 KERN_TYPE = 27300; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_USES_BITS_64; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_NATIVE_THREADS diff --git a/src/modules/module_27400.c b/src/modules/module_27400.c index 7d9f0fbd4..059b50cb5 100644 --- a/src/modules/module_27400.c +++ b/src/modules/module_27400.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "VMware VMX (PBKDF2-HMAC-SHA1 + AES-256-CBC)"; static const u64 KERN_TYPE = 27400; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE diff --git a/src/modules/module_27500.c b/src/modules/module_27500.c index ebabe24e5..6b79eba3a 100644 --- a/src/modules/module_27500.c +++ b/src/modules/module_27500.c @@ -22,6 +22,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "VirtualBox (PBKDF2-HMAC-SHA256 & AES-128-XTS)"; static const u64 KERN_TYPE = 27500; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP | OPTI_TYPE_SLOW_HASH_SIMD_LOOP2; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE diff --git a/src/modules/module_27600.c b/src/modules/module_27600.c index b4cdade54..3259ce05f 100644 --- a/src/modules/module_27600.c +++ b/src/modules/module_27600.c @@ -22,6 +22,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "VirtualBox (PBKDF2-HMAC-SHA256 & AES-256-XTS)"; static const u64 KERN_TYPE = 27600; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP | OPTI_TYPE_SLOW_HASH_SIMD_LOOP2; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE diff --git a/src/modules/module_29451.c b/src/modules/module_29451.c index c8e10bdcf..fe3c0737f 100644 --- a/src/modules/module_29451.c +++ b/src/modules/module_29451.c @@ -23,7 +23,8 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "VeraCrypt SHA256 + XTS 512 bit"; static const u64 KERN_TYPE = 13751; // old kernel used here static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE - | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; + | OPTI_TYPE_SLOW_HASH_SIMD_LOOP + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_LOOP_EXTENDED diff --git a/src/modules/module_29452.c b/src/modules/module_29452.c index 395839f64..842f36f47 100644 --- a/src/modules/module_29452.c +++ b/src/modules/module_29452.c @@ -23,7 +23,8 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "VeraCrypt SHA256 + XTS 1024 bit"; static const u64 KERN_TYPE = 13752; // old kernel used here static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE - | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; + | OPTI_TYPE_SLOW_HASH_SIMD_LOOP + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_LOOP_EXTENDED diff --git a/src/modules/module_29453.c b/src/modules/module_29453.c index 97bf9a333..6162337de 100644 --- a/src/modules/module_29453.c +++ b/src/modules/module_29453.c @@ -23,7 +23,8 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "VeraCrypt SHA256 + XTS 1536 bit"; static const u64 KERN_TYPE = 13753; // old kernel used here static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE - | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; + | OPTI_TYPE_SLOW_HASH_SIMD_LOOP + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_LOOP_EXTENDED diff --git a/src/modules/module_29461.c b/src/modules/module_29461.c index be7dcfc2e..ff2705da7 100644 --- a/src/modules/module_29461.c +++ b/src/modules/module_29461.c @@ -23,7 +23,8 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "VeraCrypt SHA256 + XTS 512 bit + boot-mode"; static const u64 KERN_TYPE = 13751; // old kernel used here static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE - | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; + | OPTI_TYPE_SLOW_HASH_SIMD_LOOP + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_LOOP_EXTENDED diff --git a/src/modules/module_29462.c b/src/modules/module_29462.c index 7e9a6535c..14b53d6ad 100644 --- a/src/modules/module_29462.c +++ b/src/modules/module_29462.c @@ -23,7 +23,8 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "VeraCrypt SHA256 + XTS 1024 bit + boot-mode"; static const u64 KERN_TYPE = 13752; // old kernel used here static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE - | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; + | OPTI_TYPE_SLOW_HASH_SIMD_LOOP + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_LOOP_EXTENDED diff --git a/src/modules/module_29463.c b/src/modules/module_29463.c index e98c71289..7a5c8ac92 100644 --- a/src/modules/module_29463.c +++ b/src/modules/module_29463.c @@ -23,7 +23,8 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "VeraCrypt SHA256 + XTS 1536 bit + boot-mode"; static const u64 KERN_TYPE = 13753; // old kernel used here static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE - | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; + | OPTI_TYPE_SLOW_HASH_SIMD_LOOP + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_LOOP_EXTENDED diff --git a/src/modules/module_70000.c b/src/modules/module_70000.c index 01504ea09..57bd77d12 100644 --- a/src/modules/module_70000.c +++ b/src/modules/module_70000.c @@ -29,7 +29,7 @@ static const u64 BRIDGE_TYPE = BRIDGE_TYPE_MATCH_TUNINGS // optional - impr | BRIDGE_TYPE_REPLACE_LOOP; static const char *BRIDGE_NAME = "argon2id_reference"; static const char *ST_PASS = "hashcat"; -static const char *ST_HASH = "$argon2id$v=19$m=4096,t=3,p=1$FoIjFnZlM2JSJWYXUgMFAw$eYKMzhbW8uyT1LLtKRdRcJj2CQeRrdr2pKv/Y71YbAQ"; +static const char *ST_HASH = "$argon2id$v=19$m=65536,t=3,p=1$FBMjI4RJBhIykCgol1KEJA$2ky5GAdhT1kH4kIgPN/oERE3Taiy43vNN70a3HpiKQU"; u32 module_attack_exec (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return ATTACK_EXEC; } u32 module_dgst_pos0 (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return DGST_POS0; } diff --git a/src/selftest.c b/src/selftest.c index f1b2b2dd5..5e744f88b 100644 --- a/src/selftest.c +++ b/src/selftest.c @@ -12,18 +12,14 @@ #include "thread.h" #include "selftest.h" -static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param) +static int selftest_init (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, u32 *highest_pw_len) { - bridge_ctx_t *bridge_ctx = hashcat_ctx->bridge_ctx; - hashconfig_t *hashconfig = hashcat_ctx->hashconfig; hashes_t *hashes = hashcat_ctx->hashes; module_ctx_t *module_ctx = hashcat_ctx->module_ctx; - status_ctx_t *status_ctx = hashcat_ctx->status_ctx; + hashconfig_t *hashconfig = hashcat_ctx->hashconfig; user_options_t *user_options = hashcat_ctx->user_options; user_options_extra_t *user_options_extra = hashcat_ctx->user_options_extra; - if (hashconfig->st_hash == NULL) return 0; - // init : replace hashes with selftest hash if (device_param->is_cuda == true) @@ -82,12 +78,12 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param tmp.pw_len = (u32) tmp_len; } + cl_event opencl_event; + pw_t pw; pw_t comb; bf_t bf; - u32 highest_pw_len = 0; - if (user_options->slow_candidates == true) { if (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL) @@ -107,12 +103,12 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_pws_buf, &pw, 1 * sizeof (pw_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_buf, &pw, 1 * sizeof (pw_t)) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_pws_buf, &pw, 1 * sizeof (pw_t), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_buf, &pw, 1 * sizeof (pw_t)) == -1) return -1; } #if defined (__APPLE__) @@ -124,7 +120,7 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (device_param->is_opencl == true) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_buf, CL_FALSE, 0, 1 * sizeof (pw_t), &pw, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_buf, CL_TRUE, 0, 1 * sizeof (pw_t), &pw, 0, NULL, &opencl_event) == -1) return -1; } } else @@ -152,12 +148,12 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_pws_buf, &pw, 1 * sizeof (pw_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_buf, &pw, 1 * sizeof (pw_t)) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_pws_buf, &pw, 1 * sizeof (pw_t), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_buf, &pw, 1 * sizeof (pw_t)) == -1) return -1; } #if defined (__APPLE__) @@ -169,7 +165,7 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (device_param->is_opencl == true) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_buf, CL_FALSE, 0, 1 * sizeof (pw_t), &pw, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_buf, CL_TRUE, 0, 1 * sizeof (pw_t), &pw, 0, NULL, &opencl_event) == -1) return -1; } } else if (user_options_extra->attack_kern == ATTACK_KERN_COMBI) @@ -222,16 +218,16 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_combs_c, &comb, 1 * sizeof (pw_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_combs_c, &comb, 1 * sizeof (pw_t)) == -1) return -1; - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_pws_buf, &pw, 1 * sizeof (pw_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_buf, &pw, 1 * sizeof (pw_t)) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_combs_c, &comb, 1 * sizeof (pw_t), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_combs_c, &comb, 1 * sizeof (pw_t)) == -1) return -1; - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_pws_buf, &pw, 1 * sizeof (pw_t), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_buf, &pw, 1 * sizeof (pw_t)) == -1) return -1; } #if defined (__APPLE__) @@ -245,9 +241,9 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (device_param->is_opencl == true) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_combs_c, CL_FALSE, 0, 1 * sizeof (pw_t), &comb, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_combs_c, CL_TRUE, 0, 1 * sizeof (pw_t), &comb, 0, NULL, &opencl_event) == -1) return -1; - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_buf, CL_FALSE, 0, 1 * sizeof (pw_t), &pw, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_buf, CL_TRUE, 0, 1 * sizeof (pw_t), &pw, 0, NULL, &opencl_event) == -1) return -1; } } else if (user_options_extra->attack_kern == ATTACK_KERN_BF) @@ -273,12 +269,12 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_pws_buf, &pw, 1 * sizeof (pw_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_buf, &pw, 1 * sizeof (pw_t)) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_pws_buf, &pw, 1 * sizeof (pw_t), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_buf, &pw, 1 * sizeof (pw_t)) == -1) return -1; } #if defined (__APPLE__) @@ -290,7 +286,7 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (device_param->is_opencl == true) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_buf, CL_FALSE, 0, 1 * sizeof (pw_t), &pw, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_buf, CL_TRUE, 0, 1 * sizeof (pw_t), &pw, 0, NULL, &opencl_event) == -1) return -1; } } else @@ -334,12 +330,12 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_bfs_c, &bf, 1 * sizeof (bf_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_bfs_c, &bf, 1 * sizeof (bf_t)) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_bfs_c, &bf, 1 * sizeof (bf_t), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_bfs_c, &bf, 1 * sizeof (bf_t)) == -1) return -1; } #if defined (__APPLE__) @@ -351,7 +347,7 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (device_param->is_opencl == true) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_bfs_c, CL_FALSE, 0, 1 * sizeof (bf_t), &bf, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_bfs_c, CL_TRUE, 0, 1 * sizeof (bf_t), &bf, 0, NULL, &opencl_event) == -1) return -1; } memset (&pw, 0, sizeof (pw)); @@ -440,12 +436,12 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_pws_buf, &pw, 1 * sizeof (pw_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_buf, &pw, 1 * sizeof (pw_t)) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_pws_buf, &pw, 1 * sizeof (pw_t), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_buf, &pw, 1 * sizeof (pw_t)) == -1) return -1; } #if defined (__APPLE__) @@ -457,10 +453,10 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (device_param->is_opencl == true) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_buf, CL_FALSE, 0, 1 * sizeof (pw_t), &pw, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_buf, CL_TRUE, 0, 1 * sizeof (pw_t), &pw, 0, NULL, &opencl_event) == -1) return -1; } - highest_pw_len = pw.pw_len; + *highest_pw_len = pw.pw_len; } } } @@ -478,12 +474,12 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_pws_buf, &pw, 1 * sizeof (pw_t), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_buf, &pw, 1 * sizeof (pw_t)) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_pws_buf, &pw, 1 * sizeof (pw_t), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_buf, &pw, 1 * sizeof (pw_t)) == -1) return -1; } #if defined (__APPLE__) @@ -495,11 +491,30 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (device_param->is_opencl == true) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_buf, CL_FALSE, 0, 1 * sizeof (pw_t), &pw, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_buf, CL_TRUE, 0, 1 * sizeof (pw_t), &pw, 0, NULL, &opencl_event) == -1) return -1; } } } + // prevent race condition on OpenCL with non-blocking write + + if (device_param->is_opencl == true) + { + if (hc_clWaitForEvents (hashcat_ctx, 1, &opencl_event) == -1) return -1; + + if (hc_clReleaseEvent (hashcat_ctx, opencl_event) == -1) return -1; + } + + return 0; +} + +static int selftest_run_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, u32 highest_pw_len) +{ + bridge_ctx_t *bridge_ctx = hashcat_ctx->bridge_ctx; + hashconfig_t *hashconfig = hashcat_ctx->hashconfig; + hashes_t *hashes = hashcat_ctx->hashes; + module_ctx_t *module_ctx = hashcat_ctx->module_ctx; + // main : run the kernel const u32 kernel_threads_sav = device_param->kernel_threads; @@ -572,14 +587,14 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (device_param->is_cuda == true) { - if (hc_cuMemcpyDtoHAsync (hashcat_ctx, device_param->hooks_buf, device_param->cuda_d_hooks, device_param->size_hooks, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyDtoH (hashcat_ctx, device_param->hooks_buf, device_param->cuda_d_hooks, device_param->size_hooks) == -1) return -1; if (hc_cuStreamSynchronize (hashcat_ctx, device_param->cuda_stream) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyDtoHAsync (hashcat_ctx, device_param->hooks_buf, device_param->hip_d_hooks, device_param->size_hooks, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyDtoH (hashcat_ctx, device_param->hooks_buf, device_param->hip_d_hooks, device_param->size_hooks) == -1) return -1; if (hc_hipStreamSynchronize (hashcat_ctx, device_param->hip_stream) == -1) return -1; } @@ -601,12 +616,12 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_hooks, device_param->hooks_buf, device_param->size_hooks, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_hooks, device_param->hooks_buf, device_param->size_hooks) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_hooks, device_param->hooks_buf, device_param->size_hooks, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_hooks, device_param->hooks_buf, device_param->size_hooks) == -1) return -1; } #if defined (__APPLE__) @@ -618,7 +633,7 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (device_param->is_opencl == true) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_hooks, CL_FALSE, 0, device_param->size_hooks, device_param->hooks_buf, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_hooks, CL_TRUE, 0, device_param->size_hooks, device_param->hooks_buf, 0, NULL, NULL) == -1) return -1; } } @@ -668,14 +683,14 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param { if (device_param->is_cuda == true) { - if (hc_cuMemcpyDtoHAsync (hashcat_ctx, device_param->h_tmps, device_param->cuda_d_tmps, hashconfig->tmp_size, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyDtoH (hashcat_ctx, device_param->h_tmps, device_param->cuda_d_tmps, hashconfig->tmp_size) == -1) return -1; if (hc_cuStreamSynchronize (hashcat_ctx, device_param->cuda_stream) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyDtoHAsync (hashcat_ctx, device_param->h_tmps, device_param->hip_d_tmps, hashconfig->tmp_size, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyDtoH (hashcat_ctx, device_param->h_tmps, device_param->hip_d_tmps, hashconfig->tmp_size) == -1) return -1; if (hc_hipStreamSynchronize (hashcat_ctx, device_param->hip_stream) == -1) return -1; } @@ -709,14 +724,14 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param { if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_tmps, device_param->h_tmps, hashconfig->tmp_size, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_tmps, device_param->h_tmps, hashconfig->tmp_size) == -1) return -1; if (hc_cuStreamSynchronize (hashcat_ctx, device_param->cuda_stream) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_tmps, device_param->h_tmps, hashconfig->tmp_size, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_tmps, device_param->h_tmps, hashconfig->tmp_size) == -1) return -1; if (hc_hipStreamSynchronize (hashcat_ctx, device_param->hip_stream) == -1) return -1; } @@ -743,14 +758,14 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (device_param->is_cuda == true) { - if (hc_cuMemcpyDtoHAsync (hashcat_ctx, device_param->hooks_buf, device_param->cuda_d_hooks, device_param->size_hooks, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyDtoH (hashcat_ctx, device_param->hooks_buf, device_param->cuda_d_hooks, device_param->size_hooks) == -1) return -1; if (hc_cuStreamSynchronize (hashcat_ctx, device_param->cuda_stream) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyDtoHAsync (hashcat_ctx, device_param->hooks_buf, device_param->hip_d_hooks, device_param->size_hooks, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyDtoH (hashcat_ctx, device_param->hooks_buf, device_param->hip_d_hooks, device_param->size_hooks) == -1) return -1; if (hc_hipStreamSynchronize (hashcat_ctx, device_param->hip_stream) == -1) return -1; } @@ -772,12 +787,12 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_hooks, device_param->hooks_buf, device_param->size_hooks, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_hooks, device_param->hooks_buf, device_param->size_hooks) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_hooks, device_param->hooks_buf, device_param->size_hooks, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_hooks, device_param->hooks_buf, device_param->size_hooks) == -1) return -1; } #if defined (__APPLE__) @@ -789,7 +804,7 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (device_param->is_opencl == true) { - if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_hooks, CL_FALSE, 0, device_param->size_hooks, device_param->hooks_buf, 0, NULL, NULL) == -1) return -1; + if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_hooks, CL_TRUE, 0, device_param->size_hooks, device_param->hooks_buf, 0, NULL, NULL) == -1) return -1; } } } @@ -831,14 +846,14 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param { if (device_param->is_cuda == true) { - if (hc_cuMemcpyDtoHAsync (hashcat_ctx, device_param->h_tmps, device_param->cuda_d_tmps, hashconfig->tmp_size, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyDtoH (hashcat_ctx, device_param->h_tmps, device_param->cuda_d_tmps, hashconfig->tmp_size) == -1) return -1; if (hc_cuStreamSynchronize (hashcat_ctx, device_param->cuda_stream) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyDtoHAsync (hashcat_ctx, device_param->h_tmps, device_param->hip_d_tmps, hashconfig->tmp_size, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyDtoH (hashcat_ctx, device_param->h_tmps, device_param->hip_d_tmps, hashconfig->tmp_size) == -1) return -1; if (hc_hipStreamSynchronize (hashcat_ctx, device_param->hip_stream) == -1) return -1; } @@ -872,12 +887,12 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param { if (device_param->is_cuda == true) { - if (hc_cuMemcpyHtoDAsync (hashcat_ctx, device_param->cuda_d_tmps, device_param->h_tmps, hashconfig->tmp_size, device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_tmps, device_param->h_tmps, hashconfig->tmp_size) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyHtoDAsync (hashcat_ctx, device_param->hip_d_tmps, device_param->h_tmps, hashconfig->tmp_size, device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_tmps, device_param->h_tmps, hashconfig->tmp_size) == -1) return -1; } #if defined (__APPLE__) @@ -933,22 +948,28 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param device_param->kernel_threads = kernel_threads_sav; - // check : check if cracked + return 0; +} - u32 num_cracked = 0; +static int selftest_cleanup (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, u32 *num_cracked) +{ + user_options_t *user_options = hashcat_ctx->user_options; + user_options_extra_t *user_options_extra = hashcat_ctx->user_options_extra; + + // check : check if cracked cl_event opencl_event; if (device_param->is_cuda == true) { - if (hc_cuMemcpyDtoHAsync (hashcat_ctx, &num_cracked, device_param->cuda_d_result, sizeof (u32), device_param->cuda_stream) == -1) return -1; + if (hc_cuMemcpyDtoH (hashcat_ctx, num_cracked, device_param->cuda_d_result, sizeof (u32)) == -1) return -1; if (hc_cuEventRecord (hashcat_ctx, device_param->cuda_event3, device_param->cuda_stream) == -1) return -1; } if (device_param->is_hip == true) { - if (hc_hipMemcpyDtoHAsync (hashcat_ctx, &num_cracked, device_param->hip_d_result, sizeof (u32), device_param->hip_stream) == -1) return -1; + if (hc_hipMemcpyDtoH (hashcat_ctx, num_cracked, device_param->hip_d_result, sizeof (u32)) == -1) return -1; if (hc_hipEventRecord (hashcat_ctx, device_param->hip_event3, device_param->hip_stream) == -1) return -1; } @@ -956,13 +977,13 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param #if defined (__APPLE__) if (device_param->is_metal == true) { - if (hc_mtlMemcpyDtoH (hashcat_ctx, device_param->metal_command_queue, &num_cracked, device_param->metal_d_result, 0, sizeof (u32)) == -1) return -1; + if (hc_mtlMemcpyDtoH (hashcat_ctx, device_param->metal_command_queue, num_cracked, device_param->metal_d_result, 0, sizeof (u32)) == -1) return -1; } #endif if (device_param->is_opencl == true) { - if (hc_clEnqueueReadBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_result, CL_FALSE, 0, sizeof (u32), &num_cracked, 0, NULL, &opencl_event) == -1) return -1; + if (hc_clEnqueueReadBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_result, CL_TRUE, 0, sizeof (u32), num_cracked, 0, NULL, &opencl_event) == -1) return -1; if (hc_clFlush (hashcat_ctx, device_param->opencl_command_queue) == -1) return -1; } @@ -1154,7 +1175,27 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (hc_clReleaseEvent (hashcat_ctx, opencl_event) == -1) return -1; } + return 0; +} + +static int process_selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param) +{ + hashconfig_t *hashconfig = hashcat_ctx->hashconfig; + status_ctx_t *status_ctx = hashcat_ctx->status_ctx; + + if (hashconfig->st_hash == NULL) return 0; + + u32 highest_pw_len = 0; + u32 num_cracked = 0; + + if (selftest_init (hashcat_ctx, device_param, &highest_pw_len) == -1) return -1; + + if (selftest_run_kernel (hashcat_ctx, device_param, highest_pw_len) == -1) return -1; + + if (selftest_cleanup (hashcat_ctx, device_param, &num_cracked) == -1) return -1; + // check return + if (num_cracked == 0) { hc_thread_mutex_lock (status_ctx->mux_display); @@ -1169,20 +1210,22 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param event_log_error (hashcat_ctx, "* Device #%u: ATTENTION! HIP kernel self-test failed.", device_param->device_id + 1); } - #if defined (__APPLE__) if (device_param->is_metal == true) { event_log_error (hashcat_ctx, "* Device #%u: ATTENTION! Metal kernel self-test failed.", device_param->device_id + 1); } - #endif if (device_param->is_opencl == true) { event_log_error (hashcat_ctx, "* Device #%u: ATTENTION! OpenCL kernel self-test failed.", device_param->device_id + 1); } - event_log_warning (hashcat_ctx, "Your device driver installation is probably broken."); - event_log_warning (hashcat_ctx, "See also: https://hashcat.net/faq/wrongdriver"); + if (device_param->is_metal == false) + { + event_log_warning (hashcat_ctx, "Your device driver installation is probably broken."); + event_log_warning (hashcat_ctx, "See also: https://hashcat.net/faq/wrongdriver"); + } + event_log_warning (hashcat_ctx, NULL); hc_thread_mutex_unlock (status_ctx->mux_display); @@ -1232,7 +1275,7 @@ HC_API_CALL void *thread_selftest (void *p) if (hc_hipCtxPushCurrent (hashcat_ctx, device_param->hip_context) == -1) return NULL; } - const int rc_selftest = selftest (hashcat_ctx, device_param); + const int rc_selftest = process_selftest (hashcat_ctx, device_param); if (user_options->benchmark == true) { diff --git a/src/shared.c b/src/shared.c index a00025352..f7dbcd78d 100644 --- a/src/shared.c +++ b/src/shared.c @@ -104,6 +104,9 @@ static const char *const OPTI_STR_USES_BITS_8 = "Uses-8-Bit"; static const char *const OPTI_STR_USES_BITS_16 = "Uses-16-Bit"; static const char *const OPTI_STR_USES_BITS_32 = "Uses-32-Bit"; static const char *const OPTI_STR_USES_BITS_64 = "Uses-64-Bit"; +static const char *const OPTI_STR_SLOW_HASH_DIMY_INIT = "Slow-Hash-DimensionY-INIT"; +static const char *const OPTI_STR_SLOW_HASH_DIMY_COMP = "Slow-Hash-DimensionY-LOOP"; +static const char *const OPTI_STR_SLOW_HASH_DIMY_LOOP = "Slow-Hash-DimensionY-COMP"; static const char *const HASH_CATEGORY_UNDEFINED_STR = "Undefined"; static const char *const HASH_CATEGORY_RAW_HASH_STR = "Raw Hash"; @@ -1072,6 +1075,9 @@ const char *stroptitype (const u32 opti_type) case OPTI_TYPE_SLOW_HASH_SIMD_LOOP: return OPTI_STR_SLOW_HASH_SIMD_LOOP; case OPTI_TYPE_SLOW_HASH_SIMD_LOOP2: return OPTI_STR_SLOW_HASH_SIMD_LOOP2; case OPTI_TYPE_SLOW_HASH_SIMD_COMP: return OPTI_STR_SLOW_HASH_SIMD_COMP; + case OPTI_TYPE_SLOW_HASH_DIMY_INIT: return OPTI_STR_SLOW_HASH_DIMY_INIT; + case OPTI_TYPE_SLOW_HASH_DIMY_LOOP: return OPTI_STR_SLOW_HASH_DIMY_LOOP; + case OPTI_TYPE_SLOW_HASH_DIMY_COMP: return OPTI_STR_SLOW_HASH_DIMY_COMP; case OPTI_TYPE_USES_BITS_8: return OPTI_STR_USES_BITS_8; case OPTI_TYPE_USES_BITS_16: return OPTI_STR_USES_BITS_16; case OPTI_TYPE_USES_BITS_32: return OPTI_STR_USES_BITS_32; diff --git a/src/terminal.c b/src/terminal.c index 4a9fc45cb..e66b6ef9a 100644 --- a/src/terminal.c +++ b/src/terminal.c @@ -1231,18 +1231,43 @@ void backend_info (hashcat_ctx_t *hashcat_ctx) const user_options_t *user_options = hashcat_ctx->user_options; const folder_config_t *folder_config = hashcat_ctx->folder_config; + if (user_options->machine_readable == true) + { + printf ("{ "); + } + if (user_options->backend_info > 1) { - event_log_info (hashcat_ctx, "System Info:"); - event_log_info (hashcat_ctx, "============"); - event_log_info (hashcat_ctx, NULL); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, "System Info:"); + event_log_info (hashcat_ctx, "============"); + event_log_info (hashcat_ctx, NULL); + } + else + { + printf ("\"SystemInfo\": { "); + } #if defined (_WIN) || defined (__CYGWIN__) || defined (__MSYS__) // TODO - event_log_info (hashcat_ctx, "OS.Name......: Windows"); - event_log_info (hashcat_ctx, "OS.Release...: N/A"); - event_log_info (hashcat_ctx, "HW.Platform..: N/A"); - event_log_info (hashcat_ctx, "HW.Model.....: N/A"); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, "OS.Name......: Windows"); + event_log_info (hashcat_ctx, "OS.Release...: N/A"); + event_log_info (hashcat_ctx, "HW.Platform..: N/A"); + event_log_info (hashcat_ctx, "HW.Model.....: N/A"); + } + else + { + printf ("\"OS\": { "); + printf ("\"Name\": \"%s\", ", "Windows"); + printf ("\"Release\": \"%s\" }, ", "N/A"); + printf ("\"Hardware\": { "); + printf ("\"Platform\": \"%s\", ", "N/A"); + printf ("\"Model\": \"%s\" } ", "N/A"); + printf ("}, "); + } #else struct utsname utsbuf; @@ -1279,10 +1304,23 @@ void backend_info (hashcat_ctx_t *hashcat_ctx) rc_uname = true; } - event_log_info (hashcat_ctx, "OS.Name......: %s", (rc_uname == true) ? utsbuf.sysname : "N/A"); - event_log_info (hashcat_ctx, "OS.Release...: %s", (rc_uname == true) ? utsbuf.release : "N/A"); - event_log_info (hashcat_ctx, "HW.Model.....: %s", (rc_sysctl == true) ? hw_model_buf : "N/A"); - event_log_info (hashcat_ctx, "HW.Platform..: %s", (rc_uname == true) ? utsbuf.machine : "N/A"); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, "OS.Name......: %s", (rc_uname == true) ? utsbuf.sysname : "N/A"); + event_log_info (hashcat_ctx, "OS.Release...: %s", (rc_uname == true) ? utsbuf.release : "N/A"); + event_log_info (hashcat_ctx, "HW.Platform..: %s", (rc_uname == true) ? utsbuf.machine : "N/A"); + event_log_info (hashcat_ctx, "HW.Model.....: %s", (rc_sysctl == true) ? hw_model_buf : "N/A"); + } + else + { + printf ("\"OS\": { "); + printf ("\"Name\": \"%s\", ", (rc_uname == true) ? utsbuf.sysname : "N/A"); + printf ("\"Release\": \"%s\" }, ", (rc_uname == true) ? utsbuf.release : "N/A"); + printf ("\"Hardware\": { "); + printf ("\"Platform\": \"%s\", ", (rc_uname == true) ? utsbuf.machine : "N/A"); + printf ("\"Model\": \"%s\" } ", (rc_sysctl == true) ? hw_model_buf : "N/A"); + printf ("}, "); + } if (rc_sysctl == true) { @@ -1290,38 +1328,72 @@ void backend_info (hashcat_ctx_t *hashcat_ctx) } #endif // _WIN || __CYGWIN__ || __MSYS__ - event_log_info (hashcat_ctx, NULL); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, NULL); - event_log_info (hashcat_ctx, "Environment Info:"); - event_log_info (hashcat_ctx, "================="); - event_log_info (hashcat_ctx, NULL); + event_log_info (hashcat_ctx, "Environment Info:"); + event_log_info (hashcat_ctx, "================="); + event_log_info (hashcat_ctx, NULL); - event_log_info (hashcat_ctx, "Cur.Work.Dir.: %s", folder_config->cwd); - event_log_info (hashcat_ctx, "Install.Dir..: %s", folder_config->install_dir); - event_log_info (hashcat_ctx, "Profile.Dir..: %s", folder_config->profile_dir); - event_log_info (hashcat_ctx, "Cache.Dir....: %s", folder_config->cache_dir); - // uninitialized at this point, for instance if the user uses --session - //event_log_info (hashcat_ctx, "Session.Dir..: %s", folder_config->session_dir); - event_log_info (hashcat_ctx, "Shared.Dir...: %s", folder_config->shared_dir); - event_log_info (hashcat_ctx, "CL.Inc.Path..: %s", folder_config->cpath_real); + event_log_info (hashcat_ctx, "Cur.Work.Dir.: %s", folder_config->cwd); + event_log_info (hashcat_ctx, "Install.Dir..: %s", folder_config->install_dir); + event_log_info (hashcat_ctx, "Profile.Dir..: %s", folder_config->profile_dir); + event_log_info (hashcat_ctx, "Cache.Dir....: %s", folder_config->cache_dir); + // uninitialized at this point, for instance if the user uses --session + //event_log_info (hashcat_ctx, "Session.Dir..: %s", folder_config->session_dir); + event_log_info (hashcat_ctx, "Shared.Dir...: %s", folder_config->shared_dir); + event_log_info (hashcat_ctx, "CL.Inc.Path..: %s", folder_config->cpath_real); - event_log_info (hashcat_ctx, NULL); + event_log_info (hashcat_ctx, NULL); + } + else + { + printf ("\"EnvironmentInfo\": { "); + printf ("\"CurrentWorkingDirectory\": \"%s\", ", folder_config->cwd); + printf ("\"InstallDirectory\": \"%s\", ", folder_config->install_dir); + printf ("\"ProfileDirectory\": \"%s\", ", folder_config->profile_dir); + printf ("\"CacheDirectory\": \"%s\", ", folder_config->cache_dir); + printf ("\"SharedDirectory\": \"%s\", ", folder_config->shared_dir); + printf ("\"CLIncludePath\": \"%s\" ", folder_config->cpath_real); + printf ("}, "); + } } if (backend_ctx->cuda) { - event_log_info (hashcat_ctx, "CUDA Info:"); - event_log_info (hashcat_ctx, "=========="); - event_log_info (hashcat_ctx, NULL); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, "CUDA Info:"); + event_log_info (hashcat_ctx, "=========="); + event_log_info (hashcat_ctx, NULL); + } + else + { + printf ("\"CUDAInfo\": { "); + } int cuda_devices_cnt = backend_ctx->cuda_devices_cnt; int cuda_driver_version = backend_ctx->cuda_driver_version; - event_log_info (hashcat_ctx, "CUDA.Version.: %u.%u", cuda_driver_version / 1000, (cuda_driver_version % 100) / 10); - event_log_info (hashcat_ctx, NULL); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, "CUDA.Version.: %u.%u", cuda_driver_version / 1000, (cuda_driver_version % 100) / 10); + event_log_info (hashcat_ctx, NULL); + } + else + { + printf ("\"Version\": \"%u.%u\", ", cuda_driver_version / 1000, (cuda_driver_version % 100) / 10); + printf ("\"BackendDevices\": [ "); + } for (int cuda_devices_idx = 0; cuda_devices_idx < cuda_devices_cnt; cuda_devices_idx++) { + if (user_options->machine_readable == true) + { + printf ("{ "); + } + const int backend_devices_idx = backend_ctx->backend_device_from_cuda[cuda_devices_idx]; const hc_device_param_t *device_param = backend_ctx->devices_param + backend_devices_idx; @@ -1340,29 +1412,88 @@ void backend_info (hashcat_ctx_t *hashcat_ctx) if (device_param->device_id_alias_cnt) { - event_log_info (hashcat_ctx, "Backend Device ID #%02u (Alias: #%02u)", device_id + 1, device_param->device_id_alias_buf[0] + 1); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, "Backend Device ID #%02u (Alias: #%02u)", device_id + 1, device_param->device_id_alias_buf[0] + 1); + } + else + { + printf ("\"DeviceID\": \"%02u\", ", device_id + 1); + printf ("\"Alias\": \"%02u\", ", device_param->device_id_alias_buf[0] + 1); + } } else { - event_log_info (hashcat_ctx, "Backend Device ID #%02u", device_id + 1); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, "Backend Device ID #%02u", device_id + 1); + } + else + { + printf ("\"DeviceID\": \"%02u\", ", device_id + 1); + } } - event_log_info (hashcat_ctx, " Name...........: %s", device_name); - event_log_info (hashcat_ctx, " Processor(s)...: %u", device_processors); - event_log_info (hashcat_ctx, " Clock..........: %u", device_maxclock_frequency); - event_log_info (hashcat_ctx, " Memory.Total...: %" PRIu64 " MB", device_global_mem / 1024 / 1024); - event_log_info (hashcat_ctx, " Memory.Free....: %" PRIu64 " MB", device_available_mem / 1024 / 1024); - event_log_info (hashcat_ctx, " Local.Memory...: %" PRIu64 " KB", device_local_mem_size / 1024); - event_log_info (hashcat_ctx, " PCI.Addr.BDFe..: %04x:%02x:%02x.%u", (u16) pcie_domain, pcie_bus, pcie_device, pcie_function); - event_log_info (hashcat_ctx, NULL); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, " Name...........: %s", device_name); + event_log_info (hashcat_ctx, " Processor(s)...: %u", device_processors); + event_log_info (hashcat_ctx, " Clock..........: %u", device_maxclock_frequency); + event_log_info (hashcat_ctx, " Memory.Total...: %" PRIu64 " MB", device_global_mem / 1024 / 1024); + event_log_info (hashcat_ctx, " Memory.Free....: %" PRIu64 " MB", device_available_mem / 1024 / 1024); + event_log_info (hashcat_ctx, " Local.Memory...: %" PRIu64 " KB", device_local_mem_size / 1024); + event_log_info (hashcat_ctx, " PCI.Addr.BDFe..: %04x:%02x:%02x.%u", (u16) pcie_domain, pcie_bus, pcie_device, pcie_function); + event_log_info (hashcat_ctx, NULL); + } + else + { + printf ("\"Name\": \"%s\", ", device_name); + printf ("\"Processor(s)\": \"%u\", ", device_processors); + printf ("\"Clock\": \"%u\", ", device_maxclock_frequency); + printf ("\"MemoryTotal\": \"%" PRIu64 " MB\", ", device_global_mem / 1024 / 1024); + printf ("\"MemoryFree\": \"%" PRIu64 " MB\", ", device_available_mem / 1024 / 1024); + printf ("\"LocalMemory\": \"%" PRIu64 " MB\", ", device_local_mem_size / 1024); + printf ("\"PCI.Addr.BDFe\": \"%04x:%02x:%02x.%u\" ", (u16) pcie_domain, pcie_bus, pcie_device, pcie_function); + } + + if (user_options->machine_readable == true) + { + if ((cuda_devices_idx + 1) < cuda_devices_cnt) + { + printf ("}, "); + } + else + { + printf ("} "); + } + } + } + + if (user_options->machine_readable == true) + { + if (backend_ctx->hip || backend_ctx->mtl || backend_ctx->ocl) + { + printf ("] }, "); + } + else + { + printf ("] } "); + } } } if (backend_ctx->hip) { - event_log_info (hashcat_ctx, "HIP Info:"); - event_log_info (hashcat_ctx, "========="); - event_log_info (hashcat_ctx, NULL); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, "HIP Info:"); + event_log_info (hashcat_ctx, "========="); + event_log_info (hashcat_ctx, NULL); + } + else + { + printf ("\"HIPInfo\": { "); + } int hip_devices_cnt = backend_ctx->hip_devices_cnt; int hip_runtimeVersion = backend_ctx->hip_runtimeVersion; @@ -1373,17 +1504,41 @@ void backend_info (hashcat_ctx_t *hashcat_ctx) int hip_version_minor = (hip_runtimeVersion - (hip_version_major * 10000000)) / 100000; int hip_version_patch = (hip_runtimeVersion - (hip_version_major * 10000000) - (hip_version_minor * 100000)); - event_log_info (hashcat_ctx, "HIP.Version.: %u.%u.%u", hip_version_major, hip_version_minor, hip_version_patch); - event_log_info (hashcat_ctx, NULL); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, "HIP.Version.: %u.%u.%u", hip_version_major, hip_version_minor, hip_version_patch); + event_log_info (hashcat_ctx, NULL); + } + else + { + printf ("\"Version\": \"%u.%u.%u\", ", hip_version_major, hip_version_minor, hip_version_patch); + } } else { - event_log_info (hashcat_ctx, "HIP.Version.: %u.%u", hip_runtimeVersion / 100, hip_runtimeVersion % 10); - event_log_info (hashcat_ctx, NULL); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, "HIP.Version.: %u.%u", hip_runtimeVersion / 100, hip_runtimeVersion % 10); + event_log_info (hashcat_ctx, NULL); + } + else + { + printf ("\"Version\": \"%u.%u\", ", hip_runtimeVersion / 100, hip_runtimeVersion % 10); + } + } + + if (user_options->machine_readable == true) + { + printf ("\"BackendDevices\": [ "); } for (int hip_devices_idx = 0; hip_devices_idx < hip_devices_cnt; hip_devices_idx++) { + if (user_options->machine_readable == true) + { + printf ("{ "); + } + const int backend_devices_idx = backend_ctx->backend_device_from_hip[hip_devices_idx]; const hc_device_param_t *device_param = backend_ctx->devices_param + backend_devices_idx; @@ -1402,40 +1557,116 @@ void backend_info (hashcat_ctx_t *hashcat_ctx) if (device_param->device_id_alias_cnt) { - event_log_info (hashcat_ctx, "Backend Device ID #%02u (Alias: #%02u)", device_id + 1, device_param->device_id_alias_buf[0] + 1); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, "Backend Device ID #%02u (Alias: #%02u)", device_id + 1, device_param->device_id_alias_buf[0] + 1); + } + else + { + printf ("\"DeviceID\": \"%02u\", ", device_id + 1); + printf ("\"Alias\": \"%02u\", ", device_param->device_id_alias_buf[0] + 1); + } } else { - event_log_info (hashcat_ctx, "Backend Device ID #%02u", device_id + 1); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, "Backend Device ID #%02u", device_id + 1); + } + else + { + printf ("\"DeviceID\": \"%02u\", ", device_id + 1); + } } - event_log_info (hashcat_ctx, " Name...........: %s", device_name); - event_log_info (hashcat_ctx, " Processor(s)...: %u", device_processors); - event_log_info (hashcat_ctx, " Clock..........: %u", device_maxclock_frequency); - event_log_info (hashcat_ctx, " Memory.Total...: %" PRIu64 " MB", device_global_mem / 1024 / 1024); - event_log_info (hashcat_ctx, " Memory.Free....: %" PRIu64 " MB", device_available_mem / 1024 / 1024); - event_log_info (hashcat_ctx, " Local.Memory...: %" PRIu64 " KB", device_local_mem_size / 1024); - event_log_info (hashcat_ctx, " PCI.Addr.BDFe..: %04x:%02x:%02x.%u", (u16) pcie_domain, pcie_bus, pcie_device, pcie_function); - event_log_info (hashcat_ctx, NULL); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, " Name...........: %s", device_name); + event_log_info (hashcat_ctx, " Processor(s)...: %u", device_processors); + event_log_info (hashcat_ctx, " Clock..........: %u", device_maxclock_frequency); + event_log_info (hashcat_ctx, " Memory.Total...: %" PRIu64 " MB", device_global_mem / 1024 / 1024); + event_log_info (hashcat_ctx, " Memory.Free....: %" PRIu64 " MB", device_available_mem / 1024 / 1024); + event_log_info (hashcat_ctx, " Local.Memory...: %" PRIu64 " KB", device_local_mem_size / 1024); + event_log_info (hashcat_ctx, " PCI.Addr.BDFe..: %04x:%02x:%02x.%u", (u16) pcie_domain, pcie_bus, pcie_device, pcie_function); + event_log_info (hashcat_ctx, NULL); + } + else + { + printf ("\"Name\": \"%s\", ", device_name); + printf ("\"Processor(s)\": \"%u\", ", device_processors); + printf ("\"Clock\": \"%u\", ", device_maxclock_frequency); + printf ("\"MemoryTotal\": \"%" PRIu64 " MB\", ", device_global_mem / 1024 / 1024); + printf ("\"MemoryFree\": \"%" PRIu64 " MB\", ", device_available_mem / 1024 / 1024); + printf ("\"LocalMemory\": \"%" PRIu64 " MB\", ", device_local_mem_size / 1024); + printf ("\"PCI.Addr.BDFe\": \"%04x:%02x:%02x.%u\" ", (u16) pcie_domain, pcie_bus, pcie_device, pcie_function); + } + + if (user_options->machine_readable == true) + { + if ((hip_devices_idx + 1) < hip_devices_cnt) + { + printf ("}, "); + } + else + { + printf ("} "); + } + } + } + + if (user_options->machine_readable == true) + { + if (backend_ctx->mtl || backend_ctx->ocl) + { + printf ("] }, "); + } + else + { + printf ("] } "); + } } } #if defined (__APPLE__) if (backend_ctx->mtl) { - event_log_info (hashcat_ctx, "Metal Info:"); - event_log_info (hashcat_ctx, "==========="); - event_log_info (hashcat_ctx, NULL); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, "Metal Info:"); + event_log_info (hashcat_ctx, "==========="); + event_log_info (hashcat_ctx, NULL); + } + else + { + printf ("\"MetalInfo\": { "); + } int metal_devices_cnt = backend_ctx->metal_devices_cnt; char *metal_runtimeVersionStr = backend_ctx->metal_runtimeVersionStr; - event_log_info (hashcat_ctx, "Metal.Version.: %s", metal_runtimeVersionStr); - event_log_info (hashcat_ctx, NULL); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, "Metal.Version.: %s", metal_runtimeVersionStr); + event_log_info (hashcat_ctx, NULL); + } + else + { + printf ("\"Version\": \"%s\", ", metal_runtimeVersionStr); + } + + if (user_options->machine_readable == true) + { + printf ("\"BackendDevices\": [ "); + } for (int metal_devices_idx = 0; metal_devices_idx < metal_devices_cnt; metal_devices_idx++) { + if (user_options->machine_readable == true) + { + printf ("{ "); + } + const int backend_devices_idx = backend_ctx->backend_device_from_metal[metal_devices_idx]; const hc_device_param_t *device_param = backend_ctx->devices_param + backend_devices_idx; @@ -1466,30 +1697,111 @@ void backend_info (hashcat_ctx_t *hashcat_ctx) if (device_param->device_id_alias_cnt) { - event_log_info (hashcat_ctx, "Backend Device ID #%02u (Alias: #%02u)", device_id + 1, device_param->device_id_alias_buf[0] + 1); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, "Backend Device ID #%02u (Alias: #%02u)", device_id + 1, device_param->device_id_alias_buf[0] + 1); + } + else + { + printf ("\"DeviceID\": \"%02u\", ", device_id + 1); + printf ("\"Alias\": \"%02u\", ", device_param->device_id_alias_buf[0] + 1); + } } else { - event_log_info (hashcat_ctx, "Backend Device ID #%02u", device_id + 1); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, "Backend Device ID #%02u", device_id + 1); + } + else + { + printf ("\"DeviceID\": \"%02u\", ", device_id + 1); + } } - event_log_info (hashcat_ctx, " Type...........: %s", ((opencl_device_type & CL_DEVICE_TYPE_CPU) ? "CPU" : ((opencl_device_type & CL_DEVICE_TYPE_GPU) ? "GPU" : "Accelerator"))); - event_log_info (hashcat_ctx, " Vendor.ID......: %u", opencl_device_vendor_id); - event_log_info (hashcat_ctx, " Vendor.........: %s", opencl_device_vendor); - event_log_info (hashcat_ctx, " Name...........: %s", device_name); - event_log_info (hashcat_ctx, " Processor(s)...: %u", device_processors); - event_log_info (hashcat_ctx, " Clock..........: N/A"); - event_log_info (hashcat_ctx, " Memory.Total...: %" PRIu64 " MB (limited to %" PRIu64 " MB allocatable in one block)", device_global_mem / 1024 / 1024, device_maxmem_alloc / 1024 / 1024); - event_log_info (hashcat_ctx, " Memory.Free....: %" PRIu64 " MB", device_available_mem / 1024 / 1024); - event_log_info (hashcat_ctx, " Local.Memory...: %" PRIu64 " KB", device_local_mem_size / 1024); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, " Type...........: %s", ((opencl_device_type & CL_DEVICE_TYPE_CPU) ? "CPU" : ((opencl_device_type & CL_DEVICE_TYPE_GPU) ? "GPU" : "Accelerator"))); + event_log_info (hashcat_ctx, " Vendor.ID......: %u", opencl_device_vendor_id); + event_log_info (hashcat_ctx, " Vendor.........: %s", opencl_device_vendor); + event_log_info (hashcat_ctx, " Name...........: %s", device_name); + event_log_info (hashcat_ctx, " Processor(s)...: %u", device_processors); + event_log_info (hashcat_ctx, " Clock..........: N/A"); + event_log_info (hashcat_ctx, " Memory.Total...: %" PRIu64 " MB (limited to %" PRIu64 " MB allocatable in one block)", device_global_mem / 1024 / 1024, device_maxmem_alloc / 1024 / 1024); + event_log_info (hashcat_ctx, " Memory.Free....: %" PRIu64 " MB", device_available_mem / 1024 / 1024); + event_log_info (hashcat_ctx, " Local.Memory...: %" PRIu64 " KB", device_local_mem_size / 1024); + } + else + { + printf ("\"Type\": \"%s\", ", ((opencl_device_type & CL_DEVICE_TYPE_CPU) ? "CPU" : ((opencl_device_type & CL_DEVICE_TYPE_GPU) ? "GPU" : "Accelerator"))); + printf ("\"VendorID\": \"%u\", ", opencl_device_vendor_id); + printf ("\"Vendor\": \"%s\", ", opencl_device_vendor); + printf ("\"Name\": \"%s\", ", device_name); + printf ("\"Processor(s)\": \"%u\", ", device_processors); + printf ("\"Clock\": \"%s\", ", "N/A"); + printf ("\"MemoryTotal\": \"%" PRIu64 " MB\", ", device_global_mem / 1024 / 1024); + printf ("\"MemoryAllocPerBlock\": \"%" PRIu64 " MB\", ", device_maxmem_alloc / 1024 / 1024); + printf ("\"MemoryFree\": \"%" PRIu64 " MB\", ", device_available_mem / 1024 / 1024); + printf ("\"LocalMemory\": \"%" PRIu64 " MB\", ", device_local_mem_size / 1024); + } switch (device_physical_location) { - case MTL_DEVICE_LOCATION_BUILTIN: event_log_info (hashcat_ctx, " Phys.Location..: built-in"); break; - case MTL_DEVICE_LOCATION_SLOT: event_log_info (hashcat_ctx, " Phys.Location..: connected to slot %u", device_location_number); break; - case MTL_DEVICE_LOCATION_EXTERNAL: event_log_info (hashcat_ctx, " Phys.Location..: connected via an external interface (port %u)", device_location_number); break; - case MTL_DEVICE_LOCATION_UNSPECIFIED: event_log_info (hashcat_ctx, " Phys.Location..: unspecified"); break; - default: event_log_info (hashcat_ctx, " Phys.Location..: N/A"); break; + case MTL_DEVICE_LOCATION_BUILTIN: + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, " Phys.Location..: built-in"); + } + else + { + printf ("\"PhysicalLocation\": \"built-in\", "); + } + + break; + case MTL_DEVICE_LOCATION_SLOT: + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, " Phys.Location..: connected to slot %u", device_location_number); + } + else + { + printf ("\"PhysicalLocation\": \"connected to slot %u\", ", device_location_number); + } + + break; + case MTL_DEVICE_LOCATION_EXTERNAL: + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, " Phys.Location..: connected via an external interface (port %u)", device_location_number); + } + else + { + printf ("\"PhysicalLocation\": \"connected via an external interface (port %u)\", ", device_location_number); + } + + break; + case MTL_DEVICE_LOCATION_UNSPECIFIED: + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, " Phys.Location..: unspecified"); + } + else + { + printf ("\"PhysicalLocation\": \"unspecified\", "); + } + + break; + default: + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, " Phys.Location..: N/A"); + } + else + { + printf ("\"PhysicalLocation\": \"%s\", ", "N/A"); + } + + break; } /* @@ -1503,28 +1815,92 @@ void backend_info (hashcat_ctx_t *hashcat_ctx) } */ - event_log_info (hashcat_ctx, " Registry.ID....: %u", device_registryID); - - if (device_physical_location != MTL_DEVICE_LOCATION_BUILTIN) + if (user_options->machine_readable == false) { - event_log_info (hashcat_ctx, " Max.TX.Rate....: %u MB/sec", device_max_transfer_rate); + event_log_info (hashcat_ctx, " Registry.ID....: %u", device_registryID); } else { - event_log_info (hashcat_ctx, " Max.TX.Rate....: N/A"); + printf ("\"RegistryID\": \"%u\", ", device_registryID); } - event_log_info (hashcat_ctx, " GPU.Properties.: headless %u, low-power %u, removable %u", device_is_headless, device_is_low_power, device_is_removable); - event_log_info (hashcat_ctx, NULL); + if (device_physical_location != MTL_DEVICE_LOCATION_BUILTIN) + { + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, " Max.TX.Rate....: %u MB/sec", device_max_transfer_rate); + } + else + { + printf ("\"MaxTXRate\": \"%u MB/sec\", ", device_max_transfer_rate); + } + } + else + { + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, " Max.TX.Rate....: N/A"); + } + else + { + printf ("\"MaxTXRate\": \"%s\", ", "N/A"); + } + } + + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, " GPU.Properties.: headless %u, low-power %u, removable %u", device_is_headless, device_is_low_power, device_is_removable); + event_log_info (hashcat_ctx, NULL); + } + else + { + printf ("\"GPUProperties\": { "); + printf ("\"headless\": \"%u\", ", device_is_headless); + printf ("\"low_power\": \"%u\", ", device_is_low_power); + printf ("\"removable\": \"%u\" ", device_is_removable); + printf ("} "); + } + + if (user_options->machine_readable == true) + { + if ((metal_devices_idx + 1) < metal_devices_cnt) + { + printf ("}, "); + } + else + { + printf ("} "); + } + } + } + + if (user_options->machine_readable == true) + { + if (backend_ctx->ocl) + { + printf ("] }, "); + } + else + { + printf ("] } "); + } } } #endif if (backend_ctx->ocl) { - event_log_info (hashcat_ctx, "OpenCL Info:"); - event_log_info (hashcat_ctx, "============"); - event_log_info (hashcat_ctx, NULL); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, "OpenCL Info:"); + event_log_info (hashcat_ctx, "============"); + event_log_info (hashcat_ctx, NULL); + } + else + { + printf ("\"OpenCLInfo\": { "); + printf ("\"Platforms\": [ "); + } cl_uint opencl_platforms_cnt = backend_ctx->opencl_platforms_cnt; cl_uint *opencl_platforms_devices_cnt = backend_ctx->opencl_platforms_devices_cnt; @@ -1534,19 +1910,44 @@ void backend_info (hashcat_ctx_t *hashcat_ctx) for (cl_uint opencl_platforms_idx = 0; opencl_platforms_idx < opencl_platforms_cnt; opencl_platforms_idx++) { + if (user_options->machine_readable == true) + { + printf ("{ "); + } + char *opencl_platform_vendor = opencl_platforms_vendor[opencl_platforms_idx]; char *opencl_platform_name = opencl_platforms_name[opencl_platforms_idx]; char *opencl_platform_version = opencl_platforms_version[opencl_platforms_idx]; cl_uint opencl_platform_devices_cnt = opencl_platforms_devices_cnt[opencl_platforms_idx]; - event_log_info (hashcat_ctx, "OpenCL Platform ID #%u", opencl_platforms_idx + 1); - event_log_info (hashcat_ctx, " Vendor..: %s", opencl_platform_vendor); - event_log_info (hashcat_ctx, " Name....: %s", opencl_platform_name); - event_log_info (hashcat_ctx, " Version.: %s", opencl_platform_version); - event_log_info (hashcat_ctx, NULL); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, "OpenCL Platform ID #%u", opencl_platforms_idx + 1); + event_log_info (hashcat_ctx, " Vendor..: %s", opencl_platform_vendor); + event_log_info (hashcat_ctx, " Name....: %s", opencl_platform_name); + event_log_info (hashcat_ctx, " Version.: %s", opencl_platform_version); + event_log_info (hashcat_ctx, NULL); + } + else + { + printf ("\"PlatformID\": \"%u\", ", opencl_platforms_idx + 1); + printf ("\"Vendor\": \"%s\", ", opencl_platform_vendor); + printf ("\"Name\": \"%s\", ", opencl_platform_name); + printf ("\"Version\": \"%s\", ", opencl_platform_version); + } + + if (user_options->machine_readable == true) + { + printf ("\"BackendDevices\": [ "); + } for (cl_uint opencl_platform_devices_idx = 0; opencl_platform_devices_idx < opencl_platform_devices_cnt; opencl_platform_devices_idx++) { + if (user_options->machine_readable == true) + { + printf ("{ "); + } + const int backend_devices_idx = backend_ctx->backend_device_from_opencl_platform[opencl_platforms_idx][opencl_platform_devices_idx]; const hc_device_param_t *device_param = backend_ctx->devices_param + backend_devices_idx; @@ -1568,25 +1969,58 @@ void backend_info (hashcat_ctx_t *hashcat_ctx) if (device_param->device_id_alias_cnt) { - event_log_info (hashcat_ctx, " Backend Device ID #%02u (Alias: #%02u)", device_id + 1, device_param->device_id_alias_buf[0] + 1); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, " Backend Device ID #%02u (Alias: #%02u)", device_id + 1, device_param->device_id_alias_buf[0] + 1); + } + else + { + printf ("\"DeviceID\": \"%02u\", ", device_id + 1); + printf ("\"Alias\": \"%02u\", ", device_param->device_id_alias_buf[0] + 1); + } } else { - event_log_info (hashcat_ctx, " Backend Device ID #%02u", device_id + 1); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, " Backend Device ID #%02u", device_id + 1); + } + else + { + printf ("\"DeviceID\": \"%02u\", ", device_id + 1); + } } - event_log_info (hashcat_ctx, " Type...........: %s", ((opencl_device_type & CL_DEVICE_TYPE_CPU) ? "CPU" : ((opencl_device_type & CL_DEVICE_TYPE_GPU) ? "GPU" : "Accelerator"))); - event_log_info (hashcat_ctx, " Vendor.ID......: %u", opencl_device_vendor_id); - event_log_info (hashcat_ctx, " Vendor.........: %s", opencl_device_vendor); - event_log_info (hashcat_ctx, " Name...........: %s", device_name); - event_log_info (hashcat_ctx, " Version........: %s", opencl_device_version); - event_log_info (hashcat_ctx, " Processor(s)...: %u", device_processors); - event_log_info (hashcat_ctx, " Clock..........: %u", device_maxclock_frequency); - event_log_info (hashcat_ctx, " Memory.Total...: %" PRIu64 " MB (limited to %" PRIu64 " MB allocatable in one block)", device_global_mem / 1024 / 1024, device_maxmem_alloc / 1024 / 1024); - event_log_info (hashcat_ctx, " Memory.Free....: %" PRIu64 " MB", device_available_mem / 1024 / 1024); - event_log_info (hashcat_ctx, " Local.Memory...: %" PRIu64 " KB", device_local_mem_size / 1024); - event_log_info (hashcat_ctx, " OpenCL.Version.: %s", opencl_device_c_version); - event_log_info (hashcat_ctx, " Driver.Version.: %s", opencl_driver_version); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, " Type...........: %s", ((opencl_device_type & CL_DEVICE_TYPE_CPU) ? "CPU" : ((opencl_device_type & CL_DEVICE_TYPE_GPU) ? "GPU" : "Accelerator"))); + event_log_info (hashcat_ctx, " Vendor.ID......: %u", opencl_device_vendor_id); + event_log_info (hashcat_ctx, " Vendor.........: %s", opencl_device_vendor); + event_log_info (hashcat_ctx, " Name...........: %s", device_name); + event_log_info (hashcat_ctx, " Version........: %s", opencl_device_version); + event_log_info (hashcat_ctx, " Processor(s)...: %u", device_processors); + event_log_info (hashcat_ctx, " Clock..........: %u", device_maxclock_frequency); + event_log_info (hashcat_ctx, " Memory.Total...: %" PRIu64 " MB (limited to %" PRIu64 " MB allocatable in one block)", device_global_mem / 1024 / 1024, device_maxmem_alloc / 1024 / 1024); + event_log_info (hashcat_ctx, " Memory.Free....: %" PRIu64 " MB", device_available_mem / 1024 / 1024); + event_log_info (hashcat_ctx, " Local.Memory...: %" PRIu64 " KB", device_local_mem_size / 1024); + event_log_info (hashcat_ctx, " OpenCL.Version.: %s", opencl_device_c_version); + event_log_info (hashcat_ctx, " Driver.Version.: %s", opencl_driver_version); + } + else + { + printf ("\"Type\": \"%s\", ", ((opencl_device_type & CL_DEVICE_TYPE_CPU) ? "CPU" : ((opencl_device_type & CL_DEVICE_TYPE_GPU) ? "GPU" : "Accelerator"))); + printf ("\"VendorID\": \"%u\", ", opencl_device_vendor_id); + printf ("\"Vendor\": \"%s\", ", opencl_device_vendor); + printf ("\"Name\": \"%s\", ", device_name); + printf ("\"Processor(s)\": \"%u\", ", device_processors); + printf ("\"Clock\": \"%u\", ", device_maxclock_frequency); + printf ("\"MemoryTotal\": \"%" PRIu64 " MB\", ", device_global_mem / 1024 / 1024); + printf ("\"MemoryAllocPerBlock\": \"%" PRIu64 " MB\", ", device_maxmem_alloc / 1024 / 1024); + printf ("\"MemoryFree\": \"%" PRIu64 " MB\", ", device_available_mem / 1024 / 1024); + printf ("\"LocalMemory\": \"%" PRIu64 " MB\", ", device_local_mem_size / 1024); + printf ("\"OpenCLVersion\": \"%s\", ", opencl_device_c_version); + printf ("\"DriverVersion\": \"%s\" ", opencl_device_version); + } if (device_param->opencl_device_type & CL_DEVICE_TYPE_GPU) { @@ -1596,18 +2030,68 @@ void backend_info (hashcat_ctx_t *hashcat_ctx) if ((device_param->opencl_platform_vendor_id == VENDOR_ID_AMD) && (device_param->opencl_device_vendor_id == VENDOR_ID_AMD)) { - event_log_info (hashcat_ctx, " PCI.Addr.BDF...: %02x:%02x.%u", pcie_bus, pcie_device, pcie_function); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, " PCI.Addr.BDF...: %02x:%02x.%u", pcie_bus, pcie_device, pcie_function); + } + else + { + printf (", \"PCI.Addr.BDF\": \"%02x:%02x.%u\" ", pcie_bus, pcie_device, pcie_function); + } } if ((device_param->opencl_platform_vendor_id == VENDOR_ID_NV) && (device_param->opencl_device_vendor_id == VENDOR_ID_NV)) { - event_log_info (hashcat_ctx, " PCI.Addr.BDF...: %02x:%02x.%u", pcie_bus, pcie_device, pcie_function); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, " PCI.Addr.BDF...: %02x:%02x.%u", pcie_bus, pcie_device, pcie_function); + } + else + { + printf (", \"PCI.Addr.BDF\": \"%02x:%02x.%u\" ", pcie_bus, pcie_device, pcie_function); + } } } - event_log_info (hashcat_ctx, NULL); + if (user_options->machine_readable == false) + { + event_log_info (hashcat_ctx, NULL); + } + else + { + if ((opencl_platform_devices_idx + 1) < opencl_platform_devices_cnt) + { + printf ("}, "); + } + else + { + printf ("} "); + } + } + } + + if (user_options->machine_readable == true) + { + if ((opencl_platforms_idx + 1) < opencl_platforms_cnt) + { + printf ("] }, "); + } + else + { + printf ("] } "); + } } } + + if (user_options->machine_readable == true) + { + printf ("] } "); + } + } + + if (user_options->machine_readable == true) + { + printf ("}"); } } diff --git a/tools/test_modules/m70000.pm b/tools/test_modules/m70000.pm index 4b44334ac..72861ae49 100644 --- a/tools/test_modules/m70000.pm +++ b/tools/test_modules/m70000.pm @@ -17,7 +17,7 @@ sub module_generate_hash { my $word = shift; my $salt = shift; - my $m = shift // 4096; + my $m = shift // 65536; my $t = shift // 3; my $p = shift // 1; my $len = shift // random_number (1, 2) * 16;