From 6053f473eb286887f8680f817d42f3aed5e2f5d0 Mon Sep 17 00:00:00 2001 From: Royce Williams Date: Thu, 1 Nov 2018 11:17:02 -0800 Subject: [PATCH] trailing whitespace --- OpenCL/m07000_a3-optimized.cl | 8 ++++---- OpenCL/m15400_a0-optimized.cl | 34 ++++++++++++++++----------------- OpenCL/m15400_a1-optimized.cl | 36 +++++++++++++++++------------------ OpenCL/m15400_a3-optimized.cl | 34 ++++++++++++++++----------------- OpenCL/m16400_a1-optimized.cl | 2 +- OpenCL/m18200_a1-optimized.cl | 2 +- docs/limits.txt | 10 +++++----- tools/test.pl | 2 +- 8 files changed, 64 insertions(+), 64 deletions(-) diff --git a/OpenCL/m07000_a3-optimized.cl b/OpenCL/m07000_a3-optimized.cl index bee348d51..e86c48942 100644 --- a/OpenCL/m07000_a3-optimized.cl +++ b/OpenCL/m07000_a3-optimized.cl @@ -132,8 +132,8 @@ DECLSPEC void m07000m (u32 *w0, u32 *w1, u32 *w2, u32 *w3, const u32 pw_len, __g t2[2] = w1[3]; t2[3] = w2[0]; t3[0] = w2[1]; - t3[1] = w2[2]; - t3[2] = w2[3]; + t3[1] = w2[2]; + t3[2] = w2[3]; t3[3] = w3[0]; t0[0] |= salt_buf0[0]; @@ -417,8 +417,8 @@ DECLSPEC void m07000s (u32 *w0, u32 *w1, u32 *w2, u32 *w3, const u32 pw_len, __g t2[2] = w1[3]; t2[3] = w2[0]; t3[0] = w2[1]; - t3[1] = w2[2]; - t3[2] = w2[3]; + t3[1] = w2[2]; + t3[2] = w2[3]; t3[3] = w3[0]; t0[0] |= salt_buf0[0]; diff --git a/OpenCL/m15400_a0-optimized.cl b/OpenCL/m15400_a0-optimized.cl index 9477aed4a..2f40eac18 100644 --- a/OpenCL/m15400_a0-optimized.cl +++ b/OpenCL/m15400_a0-optimized.cl @@ -38,12 +38,12 @@ DECLSPEC void chacha20_transform (const u32x *w0, const u32x *w1, const u32 *pos */ u32x ctx[16]; - + ctx[ 0] = CHACHA_CONST_00; ctx[ 1] = CHACHA_CONST_01; ctx[ 2] = CHACHA_CONST_02; ctx[ 3] = CHACHA_CONST_03; - ctx[ 4] = w0[0]; + ctx[ 4] = w0[0]; ctx[ 5] = w0[1]; ctx[ 6] = w0[2]; ctx[ 7] = w0[3]; @@ -80,7 +80,7 @@ DECLSPEC void chacha20_transform (const u32x *w0, const u32x *w1, const u32 *pos x[15] = ctx[15]; #pragma unroll - for (u8 i = 0; i < 10; i++) + for (u8 i = 0; i < 10; i++) { /* Column round */ QR(0, 4, 8, 12); @@ -115,11 +115,11 @@ DECLSPEC void chacha20_transform (const u32x *w0, const u32x *w1, const u32 *pos if (offset > 56) { /** - * Generate a second 64 byte keystream + * Generate a second 64 byte keystream */ - + ctx[12]++; - + if (all(ctx[12] == 0)) ctx[13]++; x[16] = ctx[ 0]; @@ -189,17 +189,17 @@ DECLSPEC void chacha20_transform (const u32x *w0, const u32x *w1, const u32 *pos digest[1] ^= x[index + 1] << (32 - remain * 8); digest[0] ^= x[index + 1] >> ( 0 + remain * 8); - digest[0] ^= x[index + 2] << (32 - remain * 8); + digest[0] ^= x[index + 2] << (32 - remain * 8); } else { digest[1] ^= x[index + 0]; - digest[0] ^= x[index + 1]; + digest[0] ^= x[index + 1]; } -} +} __kernel void m15400_m04 (__global pw_t *pws, __constant const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const chacha20_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) -{ +{ /** * modifier */ @@ -222,7 +222,7 @@ __kernel void m15400_m04 (__global pw_t *pws, __constant const kernel_rule_t *ru const u32 pw_len = pws[gid].pw_len; /** - * Salt prep + * Salt prep */ u32 iv[2] = { 0 }; @@ -262,7 +262,7 @@ __kernel void m15400_m04 (__global pw_t *pws, __constant const kernel_rule_t *ru const u32x r3 = digest[3]; COMPARE_M_SIMD(r0, r1, r2, r3); - } + } } __kernel void m15400_m08 (__global pw_t *pws, __constant const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const chacha20_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) @@ -274,7 +274,7 @@ __kernel void m15400_m16 (__global pw_t *pws, __constant const kernel_rule_t *ru } __kernel void m15400_s04 (__global pw_t *pws, __constant const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const chacha20_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) -{ +{ /** * modifier */ @@ -318,7 +318,7 @@ __kernel void m15400_s04 (__global pw_t *pws, __constant const kernel_rule_t *ru plain[0] = esalt_bufs[digests_offset].plain[0]; plain[1] = esalt_bufs[digests_offset].plain[1]; - + /** * digest */ @@ -330,7 +330,7 @@ __kernel void m15400_s04 (__global pw_t *pws, __constant const kernel_rule_t *ru digests_buf[digests_offset].digest_buf[DGST_R2], digests_buf[digests_offset].digest_buf[DGST_R3] }; - + /** * loop */ @@ -339,7 +339,7 @@ __kernel void m15400_s04 (__global pw_t *pws, __constant const kernel_rule_t *ru { u32x w0[4] = { 0 }; u32x w1[4] = { 0 }; - + const u32x out_len = apply_rules_vect(pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1); u32x digest[4] = { 0 }; @@ -352,7 +352,7 @@ __kernel void m15400_s04 (__global pw_t *pws, __constant const kernel_rule_t *ru const u32x r3 = digest[3]; COMPARE_S_SIMD(r0, r1, r2, r3); - } + } } __kernel void m15400_s08 (__global pw_t *pws, __constant const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const chacha20_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) diff --git a/OpenCL/m15400_a1-optimized.cl b/OpenCL/m15400_a1-optimized.cl index e0db643d7..7eecd306e 100644 --- a/OpenCL/m15400_a1-optimized.cl +++ b/OpenCL/m15400_a1-optimized.cl @@ -38,12 +38,12 @@ DECLSPEC void chacha20_transform (const u32x *w0, const u32x *w1, const u32 *pos */ u32x ctx[16]; - + ctx[ 0] = CHACHA_CONST_00; ctx[ 1] = CHACHA_CONST_01; ctx[ 2] = CHACHA_CONST_02; ctx[ 3] = CHACHA_CONST_03; - ctx[ 4] = w0[0]; + ctx[ 4] = w0[0]; ctx[ 5] = w0[1]; ctx[ 6] = w0[2]; ctx[ 7] = w0[3]; @@ -80,7 +80,7 @@ DECLSPEC void chacha20_transform (const u32x *w0, const u32x *w1, const u32 *pos x[15] = ctx[15]; #pragma unroll - for (u8 i = 0; i < 10; i++) + for (u8 i = 0; i < 10; i++) { /* Column round */ QR(0, 4, 8, 12); @@ -115,11 +115,11 @@ DECLSPEC void chacha20_transform (const u32x *w0, const u32x *w1, const u32 *pos if (offset > 56) { /** - * Generate a second 64 byte keystream + * Generate a second 64 byte keystream */ - + ctx[12]++; - + if (all(ctx[12] == 0)) ctx[13]++; x[16] = ctx[ 0]; @@ -189,17 +189,17 @@ DECLSPEC void chacha20_transform (const u32x *w0, const u32x *w1, const u32 *pos digest[1] ^= x[index + 1] << (32 - remain * 8); digest[0] ^= x[index + 1] >> ( 0 + remain * 8); - digest[0] ^= x[index + 2] << (32 - remain * 8); + digest[0] ^= x[index + 2] << (32 - remain * 8); } else { digest[1] ^= x[index + 0]; - digest[0] ^= x[index + 1]; + digest[0] ^= x[index + 1]; } -} +} __kernel void m15400_m04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const chacha20_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) -{ +{ /** * modifier */ @@ -222,14 +222,14 @@ __kernel void m15400_m04 (__global pw_t *pws, __global const kernel_rule_t *rule const u32 pw_l_len = pws[gid].pw_len; /** - * Salt prep + * Salt prep */ u32 iv[2] = { 0 }; u32 plain[2] = { 0 }; - u32 position[2] = { 0 }; + u32 position[2] = { 0 }; u32 offset = 0; - + position[0] = esalt_bufs[digests_offset].position[0]; position[1] = esalt_bufs[digests_offset].position[1]; @@ -326,7 +326,7 @@ __kernel void m15400_m16 (__global pw_t *pws, __global const kernel_rule_t *rule } __kernel void m15400_s04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const chacha20_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) -{ +{ /** * modifier */ @@ -352,14 +352,14 @@ __kernel void m15400_s04 (__global pw_t *pws, __global const kernel_rule_t *rule const u32 pw_l_len = pws[gid].pw_len; /** - * Salt prep + * Salt prep */ u32 iv[2] = { 0 }; u32 plain[2] = { 0 }; u32 position[2] = { 0 }; u32 offset = 0; - + position[0] = esalt_bufs[digests_offset].position[0]; position[1] = esalt_bufs[digests_offset].position[1]; @@ -382,7 +382,7 @@ __kernel void m15400_s04 (__global pw_t *pws, __global const kernel_rule_t *rule digests_buf[digests_offset].digest_buf[DGST_R2], digests_buf[digests_offset].digest_buf[DGST_R3] }; - + /** * loop */ @@ -456,7 +456,7 @@ __kernel void m15400_s04 (__global pw_t *pws, __global const kernel_rule_t *rule const u32x r3 = digest[3]; COMPARE_S_SIMD(r0, r1, r2, r3); - } + } } __kernel void m15400_s08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const chacha20_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) diff --git a/OpenCL/m15400_a3-optimized.cl b/OpenCL/m15400_a3-optimized.cl index 152051210..b2eaf4bbb 100644 --- a/OpenCL/m15400_a3-optimized.cl +++ b/OpenCL/m15400_a3-optimized.cl @@ -36,12 +36,12 @@ DECLSPEC void chacha20_transform (const u32x *w0, const u32x *w1, const u32 *pos */ u32x ctx[16]; - + ctx[ 0] = CHACHA_CONST_00; ctx[ 1] = CHACHA_CONST_01; ctx[ 2] = CHACHA_CONST_02; ctx[ 3] = CHACHA_CONST_03; - ctx[ 4] = w0[0]; + ctx[ 4] = w0[0]; ctx[ 5] = w0[1]; ctx[ 6] = w0[2]; ctx[ 7] = w0[3]; @@ -78,7 +78,7 @@ DECLSPEC void chacha20_transform (const u32x *w0, const u32x *w1, const u32 *pos x[15] = ctx[15]; #pragma unroll - for (u8 i = 0; i < 10; i++) + for (u8 i = 0; i < 10; i++) { /* Column round */ QR(0, 4, 8, 12); @@ -113,11 +113,11 @@ DECLSPEC void chacha20_transform (const u32x *w0, const u32x *w1, const u32 *pos if (offset > 56) { /** - * Generate a second 64 byte keystream + * Generate a second 64 byte keystream */ - + ctx[12]++; - + if (all(ctx[12] == 0)) ctx[13]++; x[16] = ctx[ 0]; @@ -187,17 +187,17 @@ DECLSPEC void chacha20_transform (const u32x *w0, const u32x *w1, const u32 *pos digest[1] ^= x[index + 1] << (32 - remain * 8); digest[0] ^= x[index + 1] >> ( 0 + remain * 8); - digest[0] ^= x[index + 2] << (32 - remain * 8); + digest[0] ^= x[index + 2] << (32 - remain * 8); } else { digest[1] ^= x[index + 0]; - digest[0] ^= x[index + 1]; + digest[0] ^= x[index + 1]; } -} +} __kernel void m15400_m04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const chacha20_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) -{ +{ /** * modifier */ @@ -220,7 +220,7 @@ __kernel void m15400_m04 (__global pw_t *pws, __global const kernel_rule_t *rule u32x out_len = pws[gid].pw_len; /** - * Salt prep + * Salt prep */ u32 iv[2] = { 0 }; @@ -232,7 +232,7 @@ __kernel void m15400_m04 (__global pw_t *pws, __global const kernel_rule_t *rule position[1] = esalt_bufs[digests_offset].position[1]; offset = esalt_bufs[digests_offset].offset; - + iv[0] = esalt_bufs[digests_offset].iv[0]; iv[1] = esalt_bufs[digests_offset].iv[1]; @@ -273,7 +273,7 @@ __kernel void m15400_m04 (__global pw_t *pws, __global const kernel_rule_t *rule COMPARE_M_SIMD(r0, r1, r2, r3); } -} +} __kernel void m15400_m08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const chacha20_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { @@ -284,7 +284,7 @@ __kernel void m15400_m16 (__global pw_t *pws, __global const kernel_rule_t *rule } __kernel void m15400_s04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const chacha20_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) -{ +{ /** * modifier */ @@ -307,7 +307,7 @@ __kernel void m15400_s04 (__global pw_t *pws, __global const kernel_rule_t *rule u32 out_len = pws[gid].pw_len; /** - * Salt prep + * Salt prep */ u32 iv[2] = { 0 }; @@ -348,7 +348,7 @@ __kernel void m15400_s04 (__global pw_t *pws, __global const kernel_rule_t *rule { const u32x w0r = words_buf_r[il_pos / VECT_SIZE]; const u32x w0x = w0l | w0r; - + u32x w0_t[4]; u32x w1_t[4]; @@ -372,7 +372,7 @@ __kernel void m15400_s04 (__global pw_t *pws, __global const kernel_rule_t *rule COMPARE_S_SIMD(r0, r1, r2, r3); } -} +} __kernel void m15400_s08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const chacha20_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) { diff --git a/OpenCL/m16400_a1-optimized.cl b/OpenCL/m16400_a1-optimized.cl index e7201890c..63184aeff 100644 --- a/OpenCL/m16400_a1-optimized.cl +++ b/OpenCL/m16400_a1-optimized.cl @@ -322,7 +322,7 @@ __kernel void m16400_s04 (__global pw_t *pws, __global const kernel_rule_t *rule w3[1] = (wordl3[1] | wordr3[1]) ^ 0x5c5c5c5c; w3[2] = 0x5c5c5c5c; w3[3] = 0x5c5c5c5c; - + /** * md5 */ diff --git a/OpenCL/m18200_a1-optimized.cl b/OpenCL/m18200_a1-optimized.cl index d7e836638..8089f09b7 100644 --- a/OpenCL/m18200_a1-optimized.cl +++ b/OpenCL/m18200_a1-optimized.cl @@ -242,7 +242,7 @@ DECLSPEC int decrypt_and_check (__local RC4_KEY *rc4_key, u32 *data, __global co else if length > 256: length is on 3 bytes, the first byte is 0x82, and the fourth byte is 0x30 (class=SEQUENCE) */ - + rc4_next_16 (rc4_key, 0, 0, edata2 + 0, out0); if (((out0[2] & 0x00ff80ff) != 0x00300079) && diff --git a/docs/limits.txt b/docs/limits.txt index 7c97ff5a6..884a781be 100644 --- a/docs/limits.txt +++ b/docs/limits.txt @@ -18,7 +18,7 @@ Dedicated hash modes allow unlimited salt length support. UTF-16 is mostly seen on Windows. UTF-8 (as mostly used on Linux and macOS) are fine. -Important: That does not mean UTF-16 file content, which is fully supported. +Important: That does not mean UTF-16 file content, which is fully supported. It only means the filename itself. @@ -44,9 +44,9 @@ For example, on a Vega64: 64 * 512 * 1024 * 1024 * 20 = 687,194,767,360 bytes ## Hashcat GPU memory usage may be limited by maximum allocation sizes of OpenCL drivers ## -Most hashcat hash modes only use a single OpenCL allocation. +Most hashcat hash modes only use a single OpenCL allocation. -The size of this allocation is limited by GPU drivers / OpenCL runtimes. +The size of this allocation is limited by GPU drivers / OpenCL runtimes. Only a few modes (like scrypt) make more than one allocation. @@ -54,13 +54,13 @@ Only a few modes (like scrypt) make more than one allocation. ## The maximum number of functions per rule is limited to 31 ## -This makes the size of one rule 128 byte. +This makes the size of one rule 128 byte. On the other hand, there is a 25% OpenCL single allocation memory limit. A typical GPU of today has 8GB = 2GB/128 = 16M rules max -If hashcat supported more functions per rule, it would be limited to fewer rules. +If hashcat supported more functions per rule, it would be limited to fewer rules. This is a trade-off game. diff --git a/tools/test.pl b/tools/test.pl index 6a5674393..c709ad063 100755 --- a/tools/test.pl +++ b/tools/test.pl @@ -10319,7 +10319,7 @@ sub rnd elsif ($mode == 18200) { $salt_buf = get_random_kerberos5_as_rep_salt (); - } + } else { my @salt_arr;