arno/hashcat
1
0
Fork 0
mirror of https://github.com/hashcat/hashcat.git synced 2025-07-23 15:08:37 +00:00
Code Issues Releases Wiki Activity

Merge branch 'master' into hashInfo2int

Browse Source
This commit is contained in:
Gabriele Gristina 2025-06-30 13:42:29 +02:00 committed by GitHub
commit 3b12c6b79d
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
70 changed files with 1383 additions and 397 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
OpenCL/inc_types.h
View File

@ -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;

3
Python/hcshared.py
View File

@ -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, \

3
docs/changes.txt
View File

@ -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

2
include/ext_OpenCL.h
View File

@ -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);

27
include/ext_cuda.h
View File

@ -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);

15
include/ext_hip.h
View File

@ -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);

18
include/types.h
View File

@ -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;

30
src/autotune.c
View File

@ -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;

395
src/backend.c
View File

File diff suppressed because it is too large Load Diff

145
src/ext_cuda.c
View File

@ -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;

143
src/ext_hip.c
View File

@ -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;

14
src/hashes.c
View File

@ -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)
{

1
src/modules/module_01300.c
View File

@ -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

1
src/modules/module_01800.c
View File

@ -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;

1
src/modules/module_06400.c
View File

@ -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;

1
src/modules/module_06800.c
View File

@ -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;

11
src/modules/module_07100.c
View File

@ -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;

1
src/modules/module_07700.c
View File

@ -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

1
src/modules/module_07701.c
View File

@ -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

1
src/modules/module_08200.c
View File

@ -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

3
src/modules/module_08300.c
View File

@ -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

3
src/modules/module_13751.c
View File

@ -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

3
src/modules/module_13752.c
View File

@ -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

3
src/modules/module_13753.c
View File

@ -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

3
src/modules/module_13761.c
View File

@ -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

3
src/modules/module_13762.c
View File

@ -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

3
src/modules/module_13763.c
View File

@ -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

1
src/modules/module_14800.c
View File

@ -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

3
src/modules/module_14900.c
View File

@ -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;

1
src/modules/module_15600.c
View File

@ -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

1
src/modules/module_16300.c
View File

@ -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

1
src/modules/module_16700.c
View File

@ -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;

1
src/modules/module_16900.c
View File

@ -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;

1
src/modules/module_18100.c
View File

@ -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

1
src/modules/module_20800.c
View File

@ -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

1
src/modules/module_21300.c
View File

@ -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

1
src/modules/module_21500.c
View File

@ -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;

1
src/modules/module_21700.c
View File

@ -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

3
src/modules/module_22100.c
View File

@ -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;

1
src/modules/module_22300.c
View File

@ -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

3
src/modules/module_22921.c
View File

@ -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;

3
src/modules/module_22941.c
View File

@ -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;

1
src/modules/module_23400.c
View File

@ -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

1
src/modules/module_23600.c
View File

@ -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

3
src/modules/module_23800.c
View File

@ -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

1
src/modules/module_24200.c
View File

@ -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

1
src/modules/module_24420.c
View File

@ -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

1
src/modules/module_25500.c
View File

@ -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;

3
src/modules/module_25900.c
View File

@ -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;

1
src/modules/module_26000.c
View File

@ -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;

1
src/modules/module_26100.c
View File

@ -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;

1
src/modules/module_26600.c
View File

@ -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;

3
src/modules/module_26700.c
View File

@ -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;

3
src/modules/module_26800.c
View File

@ -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;

1
src/modules/module_26900.c
View File

@ -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

1
src/modules/module_27300.c
View File

@ -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

1
src/modules/module_27400.c
View File

@ -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

1
src/modules/module_27500.c
View File

@ -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

1
src/modules/module_27600.c
View File

@ -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

3
src/modules/module_29451.c
View File

@ -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

3
src/modules/module_29452.c
View File

@ -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

3
src/modules/module_29453.c
View File

@ -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

3
src/modules/module_29461.c
View File

@ -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

3
src/modules/module_29462.c
View File

@ -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

3
src/modules/module_29463.c
View File

@ -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

2
src/modules/module_70000.c
View File

@ -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; }

167
src/selftest.c
View File

@ -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)
{

6
src/shared.c
View File

@ -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;

702
src/terminal.c
View File

@ -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 ("}");
}
}

2
tools/test_modules/m70000.pm
View File

@ -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;