From 7807eb6f2df6bced8742e44a746527e1dce0c4c3 Mon Sep 17 00:00:00 2001 From: jsteube Date: Wed, 30 Dec 2015 00:29:39 +0100 Subject: [PATCH] Fix -m 7800 for NV --- OpenCL/m07800_a0.cl | 12 ++++++------ OpenCL/m07800_a1.cl | 12 ++++++------ OpenCL/m07800_a3.cl | 12 ++++++------ src/shared.c | 4 ++++ 4 files changed, 22 insertions(+), 18 deletions(-) diff --git a/OpenCL/m07800_a0.cl b/OpenCL/m07800_a0.cl index 78e4df03e..79da3fa39 100644 --- a/OpenCL/m07800_a0.cl +++ b/OpenCL/m07800_a0.cl @@ -317,7 +317,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m07800_m04 (__glo * sha1 */ - u32 final[256]; + u32 final[32]; final[ 0] = swap32 (w0[0] | s0[0]); final[ 1] = swap32 (w0[1] | s0[1]); @@ -380,8 +380,8 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m07800_m04 (__glo digest[3] = SHA1M_D; digest[4] = SHA1M_E; - #pragma unroll 64 - for (int i = 0; i < 64; i++) final[i] = 0; + #pragma unroll 32 + for (int i = 0; i < 32; i++) final[i] = 0; final[0] = w0[0]; final[1] = w0[1]; @@ -597,7 +597,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m07800_s04 (__glo * sha1 */ - u32 final[256]; + u32 final[32]; final[ 0] = swap32 (w0[0] | s0[0]); final[ 1] = swap32 (w0[1] | s0[1]); @@ -660,8 +660,8 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m07800_s04 (__glo digest[3] = SHA1M_D; digest[4] = SHA1M_E; - #pragma unroll 64 - for (int i = 0; i < 64; i++) final[i] = 0; + #pragma unroll 32 + for (int i = 0; i < 32; i++) final[i] = 0; final[0] = w0[0]; final[1] = w0[1]; diff --git a/OpenCL/m07800_a1.cl b/OpenCL/m07800_a1.cl index c20d59b42..3c2f60b7c 100644 --- a/OpenCL/m07800_a1.cl +++ b/OpenCL/m07800_a1.cl @@ -369,7 +369,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m07800_m04 (__glo * sha1 */ - u32 final[256]; + u32 final[32]; final[ 0] = swap32 (w0[0] | s0[0]); final[ 1] = swap32 (w0[1] | s0[1]); @@ -432,8 +432,8 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m07800_m04 (__glo digest[3] = SHA1M_D; digest[4] = SHA1M_E; - #pragma unroll 64 - for (int i = 0; i < 64; i++) final[i] = 0; + #pragma unroll 32 + for (int i = 0; i < 32; i++) final[i] = 0; final[0] = w0[0]; final[1] = w0[1]; @@ -699,7 +699,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m07800_s04 (__glo * sha1 */ - u32 final[256]; + u32 final[32]; final[ 0] = swap32 (w0[0] | s0[0]); final[ 1] = swap32 (w0[1] | s0[1]); @@ -762,8 +762,8 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m07800_s04 (__glo digest[3] = SHA1M_D; digest[4] = SHA1M_E; - #pragma unroll 64 - for (int i = 0; i < 64; i++) final[i] = 0; + #pragma unroll 32 + for (int i = 0; i < 32; i++) final[i] = 0; final[0] = w0[0]; final[1] = w0[1]; diff --git a/OpenCL/m07800_a3.cl b/OpenCL/m07800_a3.cl index 06b9829ba..bc972d3bc 100644 --- a/OpenCL/m07800_a3.cl +++ b/OpenCL/m07800_a3.cl @@ -274,7 +274,7 @@ static void m07800m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_le * sha1 */ - u32 final[256]; + u32 final[32]; final[ 0] = swap32 (w0[0] | s0[0]); final[ 1] = swap32 (w0[1] | s0[1]); @@ -337,8 +337,8 @@ static void m07800m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_le digest[3] = SHA1M_D; digest[4] = SHA1M_E; - #pragma unroll 64 - for (int i = 0; i < 64; i++) final[i] = 0; + #pragma unroll + for (int i = 0; i < 32; i++) final[i] = 0; final[0] = w0[0]; final[1] = w0[1]; @@ -505,7 +505,7 @@ static void m07800s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_le * sha1 */ - u32 final[256]; + u32 final[32]; final[ 0] = swap32 (w0[0] | s0[0]); final[ 1] = swap32 (w0[1] | s0[1]); @@ -568,8 +568,8 @@ static void m07800s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_le digest[3] = SHA1M_D; digest[4] = SHA1M_E; - #pragma unroll 64 - for (int i = 0; i < 64; i++) final[i] = 0; + #pragma unroll 32 + for (int i = 0; i < 32; i++) final[i] = 0; final[0] = w0[0]; final[1] = w0[1]; diff --git a/src/shared.c b/src/shared.c index b6c82c688..265357abd 100644 --- a/src/shared.c +++ b/src/shared.c @@ -13296,6 +13296,10 @@ int sapg_parse_hash (char *input_buf, uint input_len, hash_t *hash_buf) } // SAP user names cannot be longer than 12 characters + // this is kinda buggy. if the username is in utf the length can be up to length 12*3 + // so far nobody complained so we stay with this because it helps in optimization + // final string can have a max size of 32 (password) + (10 * 5) = lengthMagicArray + 12 (max salt) + 1 (the 0x80) + if (user_len > 12) return (PARSER_SALT_LENGTH); // SAP user name cannot start with ! or ?