diff --git a/OpenCL/inc_platform.cl b/OpenCL/inc_platform.cl index c04575066..9af10e9ba 100644 --- a/OpenCL/inc_platform.cl +++ b/OpenCL/inc_platform.cl @@ -128,108 +128,128 @@ DECLSPEC size_t get_local_size (const u32 dimindx __attribute__((unused))) DECLSPEC u32x rotl32 (const u32x a, const int n) { #if VECT_SIZE == 1 - return __funnelshift_l(a, a, n); - #endif + + return rotl32_S (a, n); + + #else u32x t = 0; #if VECT_SIZE >= 2 - t.s0 = __funnelshift_l(a.s0, a.s0, n); - t.s1 = __funnelshift_l(a.s1, a.s1, n); + t.s0 = rotl32_S (a.s0, n); + t.s1 = rotl32_S (a.s1, n); #endif #if VECT_SIZE >= 4 - t.s2 = __funnelshift_l(a.s2, a.s2, n); - t.s3 = __funnelshift_l(a.s3, a.s3, n); + t.s2 = rotl32_S (a.s2, n); + t.s3 = rotl32_S (a.s3, n); #endif #if VECT_SIZE >= 8 - t.s4 = __funnelshift_l(a.s4, a.s4, n); - t.s5 = __funnelshift_l(a.s5, a.s5, n); - t.s6 = __funnelshift_l(a.s6, a.s6, n); - t.s7 = __funnelshift_l(a.s7, a.s7, n); + t.s4 = rotl32_S (a.s4, n); + t.s5 = rotl32_S (a.s5, n); + t.s6 = rotl32_S (a.s6, n); + t.s7 = rotl32_S (a.s7, n); #endif #if VECT_SIZE >= 16 - t.s8 = __funnelshift_l(a.s8, a.s8, n); - t.s9 = __funnelshift_l(a.s9, a.s9, n); - t.sa = __funnelshift_l(a.sa, a.sa, n); - t.sb = __funnelshift_l(a.sb, a.sb, n); - t.sc = __funnelshift_l(a.sc, a.sc, n); - t.sd = __funnelshift_l(a.sd, a.sd, n); - t.se = __funnelshift_l(a.se, a.se, n); - t.sf = __funnelshift_l(a.sf, a.sf, n); + t.s8 = rotl32_S (a.s8, n); + t.s9 = rotl32_S (a.s9, n); + t.sa = rotl32_S (a.sa, n); + t.sb = rotl32_S (a.sb, n); + t.sc = rotl32_S (a.sc, n); + t.sd = rotl32_S (a.sd, n); + t.se = rotl32_S (a.se, n); + t.sf = rotl32_S (a.sf, n); #endif return t; - //return ((a << n) | ((a >> (32 - n)))); + + #endif } DECLSPEC u32x rotr32 (const u32x a, const int n) { #if VECT_SIZE == 1 - return __funnelshift_r(a, a, n); - #endif + + return rotr32_S (a, n); + + #else u32x t = 0; #if VECT_SIZE >= 2 - t.s0 = __funnelshift_r(a.s0, a.s0, n); - t.s1 = __funnelshift_r(a.s1, a.s1, n); + t.s0 = rotr32_S (a.s0, n); + t.s1 = rotr32_S (a.s1, n); #endif #if VECT_SIZE >= 4 - t.s2 = __funnelshift_r(a.s2, a.s2, n); - t.s3 = __funnelshift_r(a.s3, a.s3, n); + t.s2 = rotr32_S (a.s2, n); + t.s3 = rotr32_S (a.s3, n); #endif #if VECT_SIZE >= 8 - t.s4 = __funnelshift_r(a.s4, a.s4, n); - t.s5 = __funnelshift_r(a.s5, a.s5, n); - t.s6 = __funnelshift_r(a.s6, a.s6, n); - t.s7 = __funnelshift_r(a.s7, a.s7, n); + t.s4 = rotr32_S (a.s4, n); + t.s5 = rotr32_S (a.s5, n); + t.s6 = rotr32_S (a.s6, n); + t.s7 = rotr32_S (a.s7, n); #endif #if VECT_SIZE >= 16 - t.s8 = __funnelshift_r(a.s8, a.s8, n); - t.s9 = __funnelshift_r(a.s9, a.s9, n); - t.sa = __funnelshift_r(a.sa, a.sa, n); - t.sb = __funnelshift_r(a.sb, a.sb, n); - t.sc = __funnelshift_r(a.sc, a.sc, n); - t.sd = __funnelshift_r(a.sd, a.sd, n); - t.se = __funnelshift_r(a.se, a.se, n); - t.sf = __funnelshift_r(a.sf, a.sf, n); + t.s8 = rotr32_S (a.s8, n); + t.s9 = rotr32_S (a.s9, n); + t.sa = rotr32_S (a.sa, n); + t.sb = rotr32_S (a.sb, n); + t.sc = rotr32_S (a.sc, n); + t.sd = rotr32_S (a.sd, n); + t.se = rotr32_S (a.se, n); + t.sf = rotr32_S (a.sf, n); #endif return t; - //return ((a >> n) | ((a << (32 - n)))); + + #endif } DECLSPEC u32 rotl32_S (const u32 a, const int n) { - return __funnelshift_l(a, a, n); - //return ((a << n) | ((a >> (32 - n)))); + #ifdef USE_FUNNELSHIFT + return __funnelshift_l (a, a, n); + #else + return ((a << n) | ((a >> (32 - n)))); + #endif } DECLSPEC u32 rotr32_S (const u32 a, const int n) { - return __funnelshift_r(a, a, n); - //return ((a >> n) | ((a << (32 - n)))); + #ifdef USE_FUNNELSHIFT + return __funnelshift_r (a, a, n); + #else + return ((a >> n) | ((a << (32 - n)))); + #endif } DECLSPEC u64x rotl64 (const u64x a, const int n) { + #if VECT_SIZE == 1 + return rotl64_S (a, n); + #else return ((a << n) | ((a >> (64 - n)))); + #endif } DECLSPEC u64x rotr64 (const u64x a, const int n) { + #if VECT_SIZE == 1 + return rotr64_S (a, n); + #else return ((a >> n) | ((a << (64 - n)))); + #endif } DECLSPEC u64 rotl64_S (const u64 a, const int n) { - return ((a << n) | ((a >> (64 - n)))); + return rotr64_S (a, 64 - n); } DECLSPEC u64 rotr64_S (const u64 a, const int n) @@ -309,93 +329,105 @@ DECLSPEC size_t get_local_size (const u32 dimindx __attribute__((unused))) DECLSPEC u32x rotl32 (const u32x a, const int n) { #if VECT_SIZE == 1 - return __funnelshift_l(a, a, n); - #endif - u32x t; + return rotl32_S (a, n); + + #else + + u32x t = 0; #if VECT_SIZE >= 2 - t.s0 = __funnelshift_l(a.s0, a.s0, n); - t.s1 = __funnelshift_l(a.s1, a.s1, n); + t.s0 = rotl32_S (a.s0, n); + t.s1 = rotl32_S (a.s1, n); #endif #if VECT_SIZE >= 4 - t.s2 = __funnelshift_l(a.s2, a.s2, n); - t.s3 = __funnelshift_l(a.s3, a.s3, n); + t.s2 = rotl32_S (a.s2, n); + t.s3 = rotl32_S (a.s3, n); #endif #if VECT_SIZE >= 8 - t.s4 = __funnelshift_l(a.s4, a.s4, n); - t.s5 = __funnelshift_l(a.s5, a.s5, n); - t.s6 = __funnelshift_l(a.s6, a.s6, n); - t.s7 = __funnelshift_l(a.s7, a.s7, n); + t.s4 = rotl32_S (a.s4, n); + t.s5 = rotl32_S (a.s5, n); + t.s6 = rotl32_S (a.s6, n); + t.s7 = rotl32_S (a.s7, n); #endif #if VECT_SIZE >= 16 - t.s8 = __funnelshift_l(a.s8, a.s8, n); - t.s9 = __funnelshift_l(a.s9, a.s9, n); - t.sa = __funnelshift_l(a.sa, a.sa, n); - t.sb = __funnelshift_l(a.sb, a.sb, n); - t.sc = __funnelshift_l(a.sc, a.sc, n); - t.sd = __funnelshift_l(a.sd, a.sd, n); - t.se = __funnelshift_l(a.se, a.se, n); - t.sf = __funnelshift_l(a.sf, a.sf, n); + t.s8 = rotl32_S (a.s8, n); + t.s9 = rotl32_S (a.s9, n); + t.sa = rotl32_S (a.sa, n); + t.sb = rotl32_S (a.sb, n); + t.sc = rotl32_S (a.sc, n); + t.sd = rotl32_S (a.sd, n); + t.se = rotl32_S (a.se, n); + t.sf = rotl32_S (a.sf, n); #endif return t; - //return ((a << n) | ((a >> (32 - n)))); + + #endif } DECLSPEC u32x rotr32 (const u32x a, const int n) { #if VECT_SIZE == 1 - return __funnelshift_r(a, a, n); - #endif - u32x t; + return rotr32_S (a, n); + + #else + + u32x t = 0; #if VECT_SIZE >= 2 - t.s0 = __funnelshift_r(a.s0, a.s0, n); - t.s1 = __funnelshift_r(a.s1, a.s1, n); + t.s0 = rotr32_S (a.s0, n); + t.s1 = rotr32_S (a.s1, n); #endif #if VECT_SIZE >= 4 - t.s2 = __funnelshift_r(a.s2, a.s2, n); - t.s3 = __funnelshift_r(a.s3, a.s3, n); + t.s2 = rotr32_S (a.s2, n); + t.s3 = rotr32_S (a.s3, n); #endif #if VECT_SIZE >= 8 - t.s4 = __funnelshift_r(a.s4, a.s4, n); - t.s5 = __funnelshift_r(a.s5, a.s5, n); - t.s6 = __funnelshift_r(a.s6, a.s6, n); - t.s7 = __funnelshift_r(a.s7, a.s7, n); + t.s4 = rotr32_S (a.s4, n); + t.s5 = rotr32_S (a.s5, n); + t.s6 = rotr32_S (a.s6, n); + t.s7 = rotr32_S (a.s7, n); #endif #if VECT_SIZE >= 16 - t.s8 = __funnelshift_r(a.s8, a.s8, n); - t.s9 = __funnelshift_r(a.s9, a.s9, n); - t.sa = __funnelshift_r(a.sa, a.sa, n); - t.sb = __funnelshift_r(a.sb, a.sb, n); - t.sc = __funnelshift_r(a.sc, a.sc, n); - t.sd = __funnelshift_r(a.sd, a.sd, n); - t.se = __funnelshift_r(a.se, a.se, n); - t.sf = __funnelshift_r(a.sf, a.sf, n); + t.s8 = rotr32_S (a.s8, n); + t.s9 = rotr32_S (a.s9, n); + t.sa = rotr32_S (a.sa, n); + t.sb = rotr32_S (a.sb, n); + t.sc = rotr32_S (a.sc, n); + t.sd = rotr32_S (a.sd, n); + t.se = rotr32_S (a.se, n); + t.sf = rotr32_S (a.sf, n); #endif return t; - //return ((a >> n) | ((a << (32 - n)))); + + #endif } DECLSPEC u32 rotl32_S (const u32 a, const int n) { - return __funnelshift_l(a, a, n); - //return ((a << n) | ((a >> (32 - n)))); + #ifdef USE_FUNNELSHIFT + return __funnelshift_l (a, a, n); + #else + return ((a << n) | ((a >> (32 - n)))); + #endif } DECLSPEC u32 rotr32_S (const u32 a, const int n) { - return __funnelshift_r(a, a, n); - //return ((a >> n) | ((a << (32 - n)))); + #ifdef USE_FUNNELSHIFT + return __funnelshift_r (a, a, n); + #else + return ((a >> n) | ((a << (32 - n)))); + #endif } DECLSPEC u64x rotl64 (const u64x a, const int n) @@ -423,6 +455,7 @@ DECLSPEC u64 rotl64_S (const u64 a, const int n) DECLSPEC u64 rotr64_S (const u64 a, const int n) { + #ifdef USE_FUNNELSHIFT vconv64_t in; in.v64 = a; @@ -444,6 +477,9 @@ DECLSPEC u64 rotr64_S (const u64 a, const int n) } return out.v64; + #else + return ((a >> n) | ((a << (64 - n)))); + #endif } #define FIXED_THREAD_COUNT(n) __launch_bounds__((n), 0) diff --git a/OpenCL/inc_vendor.h b/OpenCL/inc_vendor.h index 4221b3b9f..a52b4c899 100644 --- a/OpenCL/inc_vendor.h +++ b/OpenCL/inc_vendor.h @@ -183,11 +183,13 @@ using namespace metal; #ifdef IS_CUDA #define USE_BITSELECT #define USE_ROTATE +#define USE_FUNNELSHIFT #endif #ifdef IS_HIP #define USE_BITSELECT #define USE_ROTATE +#define USE_FUNNELSHIFT #endif #ifdef IS_ROCM @@ -218,4 +220,9 @@ using namespace metal; #define s3 w #endif +// some algorithms do not like this, eg 150, 1100, ... +#ifdef NO_FUNNELSHIFT +#undef USE_FUNNELSHIFT +#endif + #endif // INC_VENDOR_H diff --git a/src/backend.c b/src/backend.c index 02cdd73b8..578431158 100644 --- a/src/backend.c +++ b/src/backend.c @@ -15769,6 +15769,19 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) { const u64 kernel_power_max = hardware_power_max * kernel_accel_max; + // size_spilling: we cannot query this directly. + // Example output: + // ptxas . 4096 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads + // However, this is very relevant for us. In theory, these numbers could reach gigabytes, + // but in practice, excessive spilling would make the kernel too slow, + // and the kernel developer would adapt accordingly. We'll assume a maximum spilling buffer + // size of 4 KiB per thread for now. + // This setting will reduce the available memory pool on a 4090: + // 4 * 1024 * 128 * 32 = 16 MiB per -n accel, + // which adds up to 2 GiB with -n 128. + + size_t size_spilling = kernel_power_max * (4 * 1024); + // size_pws size_pws = kernel_power_max * sizeof (pw_t); @@ -15904,7 +15917,8 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) + size_st_digests + size_st_salts + size_st_esalts - + size_kernel_params; + + size_kernel_params + + size_spilling; if ((size_total + EXTRA_SPACE) > device_param->device_available_mem) memory_limit_hit = 1; diff --git a/src/modules/module_00150.c b/src/modules/module_00150.c index 74c00d8c5..4a36fc8c5 100644 --- a/src/modules/module_00150.c +++ b/src/modules/module_00150.c @@ -44,6 +44,15 @@ u32 module_salt_type (MAYBE_UNUSED const hashconfig_t *hashconfig, const char *module_st_hash (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 ST_HASH; } const char *module_st_pass (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 ST_PASS; } +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_FUNNELSHIFT"); + + 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) { u32 *digest = (u32 *) digest_buf; @@ -190,7 +199,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_DEFAULT; + module_ctx->module_jit_build_options = module_jit_build_options; module_ctx->module_jit_cache_disable = MODULE_DEFAULT; module_ctx->module_kernel_accel_max = MODULE_DEFAULT; module_ctx->module_kernel_accel_min = MODULE_DEFAULT; diff --git a/src/modules/module_01100.c b/src/modules/module_01100.c index ddd4948c9..1c3cbd731 100644 --- a/src/modules/module_01100.c +++ b/src/modules/module_01100.c @@ -50,6 +50,15 @@ u32 module_salt_type (MAYBE_UNUSED const hashconfig_t *hashconfig, const char *module_st_hash (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 ST_HASH; } const char *module_st_pass (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 ST_PASS; } +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_FUNNELSHIFT"); + + 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) { u32 *digest = (u32 *) digest_buf; @@ -197,7 +206,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_DEFAULT; + module_ctx->module_jit_build_options = module_jit_build_options; module_ctx->module_jit_cache_disable = MODULE_DEFAULT; module_ctx->module_kernel_accel_max = MODULE_DEFAULT; module_ctx->module_kernel_accel_min = MODULE_DEFAULT; diff --git a/src/modules/module_06000.c b/src/modules/module_06000.c index 81aba5f37..8b118d9af 100644 --- a/src/modules/module_06000.c +++ b/src/modules/module_06000.c @@ -42,6 +42,15 @@ u32 module_salt_type (MAYBE_UNUSED const hashconfig_t *hashconfig, const char *module_st_hash (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 ST_HASH; } const char *module_st_pass (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 ST_PASS; } +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_FUNNELSHIFT"); + + 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) { u32 *digest = (u32 *) digest_buf; @@ -150,7 +159,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_DEFAULT; + module_ctx->module_jit_build_options = module_jit_build_options; module_ctx->module_jit_cache_disable = MODULE_DEFAULT; module_ctx->module_kernel_accel_max = MODULE_DEFAULT; module_ctx->module_kernel_accel_min = MODULE_DEFAULT; diff --git a/src/modules/module_06211.c b/src/modules/module_06211.c index c444626e5..d2d9ff3ce 100644 --- a/src/modules/module_06211.c +++ b/src/modules/module_06211.c @@ -77,6 +77,15 @@ typedef struct tc static const int ROUNDS_TRUECRYPT_2K = 2000; static const float MIN_SUFFICIENT_ENTROPY_FILE = 7.0f; +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_FUNNELSHIFT"); + + return jit_build_options; +} + bool module_unstable_warning (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 hc_device_param_t *device_param) { if ((device_param->opencl_platform_vendor_id == VENDOR_ID_APPLE) && (device_param->opencl_device_type & CL_DEVICE_TYPE_GPU)) @@ -298,7 +307,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_DEFAULT; + module_ctx->module_jit_build_options = module_jit_build_options; module_ctx->module_jit_cache_disable = MODULE_DEFAULT; module_ctx->module_kernel_accel_max = MODULE_DEFAULT; module_ctx->module_kernel_accel_min = MODULE_DEFAULT; diff --git a/src/modules/module_06212.c b/src/modules/module_06212.c index 7ea785b75..e5560fa3d 100644 --- a/src/modules/module_06212.c +++ b/src/modules/module_06212.c @@ -77,6 +77,15 @@ typedef struct tc static const int ROUNDS_TRUECRYPT_2K = 2000; static const float MIN_SUFFICIENT_ENTROPY_FILE = 7.0f; +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_FUNNELSHIFT"); + + return jit_build_options; +} + bool module_unstable_warning (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 hc_device_param_t *device_param) { if ((device_param->opencl_platform_vendor_id == VENDOR_ID_APPLE) && (device_param->opencl_device_type & CL_DEVICE_TYPE_GPU)) @@ -298,7 +307,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_DEFAULT; + module_ctx->module_jit_build_options = module_jit_build_options; module_ctx->module_jit_cache_disable = MODULE_DEFAULT; module_ctx->module_kernel_accel_max = MODULE_DEFAULT; module_ctx->module_kernel_accel_min = MODULE_DEFAULT; diff --git a/src/modules/module_06213.c b/src/modules/module_06213.c index 690a3faec..8ab92f5c5 100644 --- a/src/modules/module_06213.c +++ b/src/modules/module_06213.c @@ -77,6 +77,15 @@ typedef struct tc static const int ROUNDS_TRUECRYPT_2K = 2000; static const float MIN_SUFFICIENT_ENTROPY_FILE = 7.0f; +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_FUNNELSHIFT"); + + return jit_build_options; +} + bool module_unstable_warning (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 hc_device_param_t *device_param) { if ((device_param->opencl_platform_vendor_id == VENDOR_ID_APPLE) && (device_param->opencl_device_type & CL_DEVICE_TYPE_GPU)) @@ -304,7 +313,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_DEFAULT; + module_ctx->module_jit_build_options = module_jit_build_options; module_ctx->module_jit_cache_disable = MODULE_DEFAULT; module_ctx->module_kernel_accel_max = MODULE_DEFAULT; module_ctx->module_kernel_accel_min = MODULE_DEFAULT; diff --git a/src/modules/module_08600.c b/src/modules/module_08600.c index f1e83a17e..c72da5c7f 100644 --- a/src/modules/module_08600.c +++ b/src/modules/module_08600.c @@ -44,6 +44,18 @@ u32 module_salt_type (MAYBE_UNUSED const hashconfig_t *hashconfig, const char *module_st_hash (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 ST_HASH; } const char *module_st_pass (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 ST_PASS; } +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; + + if (device_param->is_cuda == true) + { + hc_asprintf (&jit_build_options, "-D FIXED_LOCAL_SIZE=%u", (u32) device_param->device_maxworkgroup_size); + } + + return jit_build_options; +} + u32 module_pw_max (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { const u32 pw_max = 16; // Lotus Notes/Domino 5 limits itself to 16 @@ -146,7 +158,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_DEFAULT; + module_ctx->module_jit_build_options = module_jit_build_options; module_ctx->module_jit_cache_disable = MODULE_DEFAULT; module_ctx->module_kernel_accel_max = MODULE_DEFAULT; module_ctx->module_kernel_accel_min = MODULE_DEFAULT; diff --git a/src/modules/module_08700.c b/src/modules/module_08700.c index 39b157196..cf4c428cf 100644 --- a/src/modules/module_08700.c +++ b/src/modules/module_08700.c @@ -43,6 +43,18 @@ u32 module_salt_type (MAYBE_UNUSED const hashconfig_t *hashconfig, const char *module_st_hash (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 ST_HASH; } const char *module_st_pass (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 ST_PASS; } +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; + + if (device_param->is_cuda == true) + { + hc_asprintf (&jit_build_options, "-D FIXED_LOCAL_SIZE=%u", (u32) device_param->device_maxworkgroup_size); + } + + return jit_build_options; +} + u32 module_pw_max (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { const bool optimized_kernel = (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL); @@ -188,7 +200,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_DEFAULT; + module_ctx->module_jit_build_options = module_jit_build_options; module_ctx->module_jit_cache_disable = MODULE_DEFAULT; module_ctx->module_kernel_accel_max = MODULE_DEFAULT; module_ctx->module_kernel_accel_min = MODULE_DEFAULT; diff --git a/src/modules/module_09710.c b/src/modules/module_09710.c index f4520d3d3..1698ba1bc 100644 --- a/src/modules/module_09710.c +++ b/src/modules/module_09710.c @@ -28,6 +28,7 @@ static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_ALWAYS_HEXIFY | OPTS_TYPE_AUTODETECT_DISABLE; static const u32 SALT_TYPE = SALT_TYPE_EMBEDDED; +static const char *BENCHMARK_MASK = "?b?b?b?b?b"; static const char *ST_PASS = "\x91\xb2\xe0\x62\xb9"; static const char *ST_HASH = "$oldoffice$0*55045061647456688860411218030058*e7e24d163fbd743992d4b8892bf3f2f7*493410dbc832557d3fe1870ace8397e2"; @@ -134,7 +135,7 @@ u32 module_pw_max (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED con const char *module_benchmark_mask (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 "?b?b?b?b?b"; + return BENCHMARK_MASK; } u32 module_forced_outfile_format (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) diff --git a/src/modules/module_09810.c b/src/modules/module_09810.c index 433616cf1..a4aa244e4 100644 --- a/src/modules/module_09810.c +++ b/src/modules/module_09810.c @@ -27,6 +27,7 @@ static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_ALWAYS_HEXIFY | OPTS_TYPE_AUTODETECT_DISABLE; static const u32 SALT_TYPE = SALT_TYPE_EMBEDDED; +static const char *BENCHMARK_MASK = "?b?b?b?b?b"; static const char *ST_PASS = "\xb8\xf6\x36\x19\xca"; static const char *ST_HASH = "$oldoffice$3*83328705222323020515404251156288*2855956a165ff6511bc7f4cd77b9e101*941861655e73a09c40f7b1e9dfd0c256ed285acd"; @@ -135,7 +136,7 @@ u32 module_pw_max (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED con const char *module_benchmark_mask (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 "?b?b?b?b?b"; + return BENCHMARK_MASK; } u32 module_forced_outfile_format (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) diff --git a/src/modules/module_10410.c b/src/modules/module_10410.c index 25f3417de..063ac4916 100644 --- a/src/modules/module_10410.c +++ b/src/modules/module_10410.c @@ -26,6 +26,7 @@ static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_ALWAYS_HEXIFY | OPTS_TYPE_AUTODETECT_DISABLE; static const u32 SALT_TYPE = SALT_TYPE_EMBEDDED; +static const char *BENCHMARK_MASK = "?b?b?b?b?b"; static const char *ST_PASS = "\x6a\x8a\xed\xcc\xb7"; static const char *ST_HASH = "$pdf$1*2*40*-1*0*16*01221086741440841668371056103222*32*27c3fecef6d46a78eb61b8b4dbc690f5f8a2912bbb9afc842c12d79481568b74*32*0000000000000000000000000000000000000000000000000000000000000000"; @@ -144,7 +145,7 @@ u32 module_pw_max (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED con const char *module_benchmark_mask (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 "?b?b?b?b?b"; + return BENCHMARK_MASK; } u32 module_forced_outfile_format (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) diff --git a/src/modules/module_11000.c b/src/modules/module_11000.c index 536e30e73..916f0ffa1 100644 --- a/src/modules/module_11000.c +++ b/src/modules/module_11000.c @@ -45,6 +45,15 @@ u32 module_salt_type (MAYBE_UNUSED const hashconfig_t *hashconfig, const char *module_st_hash (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 ST_HASH; } const char *module_st_pass (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 ST_PASS; } +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_FUNNELSHIFT"); + + return jit_build_options; +} + u32 module_salt_min (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { const u32 salt_min = 56; @@ -181,7 +190,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_DEFAULT; + module_ctx->module_jit_build_options = module_jit_build_options; module_ctx->module_jit_cache_disable = MODULE_DEFAULT; module_ctx->module_kernel_accel_max = MODULE_DEFAULT; module_ctx->module_kernel_accel_min = MODULE_DEFAULT; diff --git a/src/modules/module_14400.c b/src/modules/module_14400.c index 8512337b7..18c8d4c4a 100644 --- a/src/modules/module_14400.c +++ b/src/modules/module_14400.c @@ -43,6 +43,15 @@ u32 module_salt_type (MAYBE_UNUSED const hashconfig_t *hashconfig, const char *module_st_hash (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 ST_HASH; } const char *module_st_pass (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 ST_PASS; } +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_FUNNELSHIFT"); + + return jit_build_options; +} + u32 module_pw_max (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { const bool optimized_kernel = (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL); @@ -188,7 +197,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_DEFAULT; + module_ctx->module_jit_build_options = module_jit_build_options; module_ctx->module_jit_cache_disable = MODULE_DEFAULT; module_ctx->module_kernel_accel_max = MODULE_DEFAULT; module_ctx->module_kernel_accel_min = MODULE_DEFAULT; diff --git a/src/modules/module_14900.c b/src/modules/module_14900.c index 7de3dba45..cbe009024 100644 --- a/src/modules/module_14900.c +++ b/src/modules/module_14900.c @@ -24,6 +24,7 @@ static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_SUGGEST_KG; static const u32 SALT_TYPE = SALT_TYPE_EMBEDDED; +static const char *BENCHMARK_MASK = "?b?b?b?b?bxxxxx"; static const char *ST_PASS = "hashcat!!!"; static const char *ST_HASH = "7090b6b9:04223875"; @@ -58,7 +59,7 @@ u32 module_pw_min (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED con const char *module_benchmark_mask (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 "?b?b?b?b?bxxxxx"; + return BENCHMARK_MASK; } 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) diff --git a/src/modules/module_16000.c b/src/modules/module_16000.c index 9781e0e72..7da4eccc7 100644 --- a/src/modules/module_16000.c +++ b/src/modules/module_16000.c @@ -42,6 +42,19 @@ u32 module_salt_type (MAYBE_UNUSED const hashconfig_t *hashconfig, const char *module_st_hash (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 ST_HASH; } const char *module_st_pass (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 ST_PASS; } + +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; + + if (device_param->is_cuda == true) + { + hc_asprintf (&jit_build_options, "-D FIXED_LOCAL_SIZE=%u", (u32) device_param->device_maxworkgroup_size); + } + + return jit_build_options; +} + u32 module_pw_max (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { const u32 pw_max = 8; // Underlaying DES max @@ -166,7 +179,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_DEFAULT; + module_ctx->module_jit_build_options = module_jit_build_options; module_ctx->module_jit_cache_disable = MODULE_DEFAULT; module_ctx->module_kernel_accel_max = MODULE_DEFAULT; module_ctx->module_kernel_accel_min = MODULE_DEFAULT; diff --git a/src/modules/module_16600.c b/src/modules/module_16600.c index 89cecad59..75d6b3cb5 100644 --- a/src/modules/module_16600.c +++ b/src/modules/module_16600.c @@ -55,6 +55,18 @@ typedef struct electrum_wallet static const char *SIGNATURE_ELECTRUM_WALLET = "$electrum$"; +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; + + if (device_param->is_cuda == true) + { + hc_asprintf (&jit_build_options, "-D FIXED_LOCAL_SIZE=%u", (u32) device_param->device_maxworkgroup_size); + } + + return jit_build_options; +} + u64 module_esalt_size (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { const u64 esalt_size = (const u64) sizeof (electrum_wallet_t); @@ -232,7 +244,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_DEFAULT; + module_ctx->module_jit_build_options = module_jit_build_options; module_ctx->module_jit_cache_disable = MODULE_DEFAULT; module_ctx->module_kernel_accel_max = MODULE_DEFAULT; module_ctx->module_kernel_accel_min = MODULE_DEFAULT; diff --git a/src/modules/module_20510.c b/src/modules/module_20510.c index 394a69a3b..8fa2d2f35 100644 --- a/src/modules/module_20510.c +++ b/src/modules/module_20510.c @@ -108,6 +108,18 @@ typedef struct pkzip_extra } pkzip_extra_t; +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; + + if (device_param->is_cuda == true) + { + hc_asprintf (&jit_build_options, "-D FIXED_LOCAL_SIZE=%u", (u32) device_param->device_maxworkgroup_size); + } + + return jit_build_options; +} + u32 module_hashes_count_max (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { const u64 tmp_size = 1; @@ -243,7 +255,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_DEFAULT; + module_ctx->module_jit_build_options = module_jit_build_options; module_ctx->module_jit_cache_disable = MODULE_DEFAULT; module_ctx->module_kernel_accel_max = MODULE_DEFAULT; module_ctx->module_kernel_accel_min = MODULE_DEFAULT; diff --git a/tools/benchmark_deep.pl b/tools/benchmark_deep.pl index 09e38928a..b75d371a0 100755 --- a/tools/benchmark_deep.pl +++ b/tools/benchmark_deep.pl @@ -548,7 +548,7 @@ sub get_module close (IN); - my $mask = (defined $benchmark_mask) ? $benchmark_mask : $default_mask; + my $mask = $default_mask; if ($pw_min != -1) { @@ -571,6 +571,8 @@ sub get_module } } + $mask = (defined $benchmark_mask) ? $benchmark_mask : $mask; + my $module = { "is_binary" => $is_binary,