From 4ed8977e487bf6b0bc66b9e9bd89cfbe241649fe Mon Sep 17 00:00:00 2001 From: jsteube Date: Mon, 27 Jun 2016 21:28:48 +0200 Subject: [PATCH] Dynamic tmp_t length for scrypt --- OpenCL/inc_types.cl | 8 +++++--- OpenCL/m08900.cl | 17 +++++------------ include/types.h | 11 +++-------- src/hashcat.c | 46 ++++++++++++++++++++++++++++++--------------- 4 files changed, 44 insertions(+), 38 deletions(-) diff --git a/OpenCL/inc_types.cl b/OpenCL/inc_types.cl index f0f4ad5..a51fc6f 100644 --- a/OpenCL/inc_types.cl +++ b/OpenCL/inc_types.cl @@ -689,8 +689,6 @@ typedef struct u32 scrypt_N; u32 scrypt_r; u32 scrypt_p; - u32 scrypt_tmto; - u32 scrypt_phy; } salt_t; @@ -1370,6 +1368,10 @@ typedef struct typedef struct { - uint4 P[64]; + #ifndef SCRYPT_TMP_ELEM + #define SCRYPT_TMP_ELEM 1 + #endif + + uint4 P[SCRYPT_TMP_ELEM]; } scrypt_tmp_t; diff --git a/OpenCL/m08900.cl b/OpenCL/m08900.cl index 6b50d19..2a8a2e0 100644 --- a/OpenCL/m08900.cl +++ b/OpenCL/m08900.cl @@ -668,20 +668,15 @@ void salsa_r (uint4 *TI) } } -void scrypt_smix (uint4 *X, uint4 *T, const u32 phy, __global uint4 *V) +void scrypt_smix (uint4 *X, uint4 *T, __global uint4 *V) { - #define Coord(x,y,z) (((x) * zSIZE) + ((y) * zSIZE * xSIZE) + (z)) + #define Coord(x,y,z) (((x) * ySIZE * zSIZE) + ((y) * zSIZE) + (z)) #define CO Coord(x,y,z) - const u32 xSIZE = phy; const u32 ySIZE = SCRYPT_N / SCRYPT_TMTO; const u32 zSIZE = STATE_CNT4; - const u32 lid = get_local_id (0); - const u32 lsz = get_local_size (0); - const u32 rid = get_group_id (0); - - const u32 x = (rid * lsz) + lid; + const u32 x = get_global_id (0); #ifdef _unroll #pragma unroll @@ -887,8 +882,6 @@ __kernel void m08900_loop (__global pw_t *pws, __global kernel_rule_t *rules_buf if (gid >= gid_max) return; - const u32 scrypt_phy = salt_bufs[salt_pos].scrypt_phy; - uint4 X[STATE_CNT4]; uint4 T[STATE_CNT4]; @@ -897,7 +890,7 @@ __kernel void m08900_loop (__global pw_t *pws, __global kernel_rule_t *rules_buf #endif for (int z = 0; z < STATE_CNT4; z++) X[z] = swap32_4 (tmps[gid].P[z]); - scrypt_smix (X, T, scrypt_phy, d_scryptV_buf); + scrypt_smix (X, T, d_scryptV_buf); #ifdef _unroll #pragma unroll @@ -909,7 +902,7 @@ __kernel void m08900_loop (__global pw_t *pws, __global kernel_rule_t *rules_buf { for (int z = 0; z < STATE_CNT4; z++) X[z] = swap32_4 (tmps[gid].P[i + z]); - scrypt_smix (X, T, scrypt_phy, d_scryptV_buf); + scrypt_smix (X, T, d_scryptV_buf); for (int z = 0; z < STATE_CNT4; z++) tmps[gid].P[i + z] = swap32_4 (X[z]); } diff --git a/include/types.h b/include/types.h index f1eca83..b9f860f 100644 --- a/include/types.h +++ b/include/types.h @@ -32,8 +32,6 @@ typedef struct uint scrypt_N; uint scrypt_r; uint scrypt_p; - uint scrypt_tmto; - uint scrypt_phy; } salt_t; @@ -295,12 +293,6 @@ typedef struct } win8phone_t; -typedef struct -{ - uint P[256]; - -} scrypt_tmp_t; - typedef struct { uint digest[4]; @@ -1210,6 +1202,9 @@ typedef struct void *esalts_buf; + uint scrypt_tmp_size; + uint scrypt_tmto_final; + /** * logging */ diff --git a/src/hashcat.c b/src/hashcat.c index 96efc1b..7c01f7a 100644 --- a/src/hashcat.c +++ b/src/hashcat.c @@ -15011,9 +15011,7 @@ int main (int argc, char **argv) if (hash_mode == 3000) kernel_threads = 64; // DES if (hash_mode == 3200) kernel_threads = 8; // Blowfish if (hash_mode == 7500) kernel_threads = 64; // RC4 - if (hash_mode == 8900) kernel_threads = 32; // scrypt if (hash_mode == 9000) kernel_threads = 8; // Blowfish - if (hash_mode == 9300) kernel_threads = 32; // scrypt if (hash_mode == 9700) kernel_threads = 64; // RC4 if (hash_mode == 9710) kernel_threads = 64; // RC4 if (hash_mode == 9800) kernel_threads = 64; // RC4 @@ -15063,6 +15061,24 @@ int main (int argc, char **argv) if ((hash_mode == 8900) || (hash_mode == 9300)) { + // we need to check that all hashes have the same scrypt settings + + const u32 scrypt_N = data.salts_buf[0].scrypt_N; + const u32 scrypt_r = data.salts_buf[0].scrypt_r; + const u32 scrypt_p = data.salts_buf[0].scrypt_p; + + for (uint i = 1; i < salts_cnt; i++) + { + if ((data.salts_buf[i].scrypt_N != scrypt_N) + || (data.salts_buf[i].scrypt_r != scrypt_r) + || (data.salts_buf[i].scrypt_p != scrypt_p)) + { + log_error ("ERROR: Mixed scrypt settings not supported"); + + return -1; + } + } + uint tmto_start = 0; uint tmto_stop = 10; @@ -15100,15 +15116,16 @@ int main (int argc, char **argv) } } + data.scrypt_tmp_size = (128 * scrypt_r); + device_param->kernel_accel_min = 1; device_param->kernel_accel_max = 8; - for (uint tmto = tmto_start; tmto < tmto_stop; tmto++) - { - // TODO: in theory the following calculation needs to be done per salt, not global - // we assume all hashes have the same scrypt settings + uint tmto; - size_scryptV = (128 * data.salts_buf[0].scrypt_r) * data.salts_buf[0].scrypt_N; + for (tmto = tmto_start; tmto < tmto_stop; tmto++) + { + size_scryptV = (128 * scrypt_r) * scrypt_N; size_scryptV /= 1 << tmto; @@ -15123,21 +15140,20 @@ int main (int argc, char **argv) for (uint salts_pos = 0; salts_pos < data.salts_cnt; salts_pos++) { - data.salts_buf[salts_pos].scrypt_tmto = tmto; - data.salts_buf[salts_pos].scrypt_phy = device_param->device_processors * device_param->kernel_threads * device_param->kernel_accel_max; + data.scrypt_tmto_final = tmto; } break; } - if (data.salts_buf[0].scrypt_phy == 0) + if (tmto == tmto_stop) { log_error ("ERROR: Can't allocate enough device memory"); return -1; } - if (quiet == 0) log_info ("SCRYPT tmto optimizer value set to: %u, mem: %u\n", data.salts_buf[0].scrypt_tmto, size_scryptV); + if (quiet == 0) log_info ("SCRYPT tmto optimizer value set to: %u, mem: %u\n", data.scrypt_tmto_final, size_scryptV); } /** @@ -15265,11 +15281,11 @@ int main (int argc, char **argv) case 7900: size_tmps = kernel_power_max * sizeof (drupal7_tmp_t); break; case 8200: size_tmps = kernel_power_max * sizeof (pbkdf2_sha512_tmp_t); break; case 8800: size_tmps = kernel_power_max * sizeof (androidfde_tmp_t); break; - case 8900: size_tmps = kernel_power_max * sizeof (scrypt_tmp_t); break; + case 8900: size_tmps = kernel_power_max * data.scrypt_tmp_size; break; case 9000: size_tmps = kernel_power_max * sizeof (pwsafe2_tmp_t); break; case 9100: size_tmps = kernel_power_max * sizeof (lotus8_tmp_t); break; case 9200: size_tmps = kernel_power_max * sizeof (pbkdf2_sha256_tmp_t); break; - case 9300: size_tmps = kernel_power_max * sizeof (scrypt_tmp_t); break; + case 9300: size_tmps = kernel_power_max * data.scrypt_tmp_size; break; case 9400: size_tmps = kernel_power_max * sizeof (office2007_tmp_t); break; case 9500: size_tmps = kernel_power_max * sizeof (office2010_tmp_t); break; case 9600: size_tmps = kernel_power_max * sizeof (office2013_tmp_t); break; @@ -15640,11 +15656,11 @@ int main (int argc, char **argv) if (force_jit_compilation == 1500) { - snprintf (build_opts_update, sizeof (build_opts_update) - 1, "%s -DDESCRYPT_SALT=%d", build_opts, data.salts_buf[0].salt_buf[0]); + snprintf (build_opts_update, sizeof (build_opts_update) - 1, "%s -DDESCRYPT_SALT=%u", build_opts, data.salts_buf[0].salt_buf[0]); } else if (force_jit_compilation == 8900) { - snprintf (build_opts_update, sizeof (build_opts_update) - 1, "%s -DSCRYPT_N=%d -DSCRYPT_R=%d -DSCRYPT_P=%d -DSCRYPT_TMTO=%d", build_opts, data.salts_buf[0].scrypt_N, data.salts_buf[0].scrypt_r, data.salts_buf[0].scrypt_p, 1 << data.salts_buf[0].scrypt_tmto); + snprintf (build_opts_update, sizeof (build_opts_update) - 1, "%s -DSCRYPT_N=%u -DSCRYPT_R=%u -DSCRYPT_P=%u -DSCRYPT_TMTO=%u -DSCRYPT_TMP_ELEM=%u", build_opts, data.salts_buf[0].scrypt_N, data.salts_buf[0].scrypt_r, data.salts_buf[0].scrypt_p, 1 << data.scrypt_tmto_final, data.scrypt_tmp_size / 16); } else { -- 2.25.1