diff --git a/OpenCL/inc_common.cl b/OpenCL/inc_common.cl index bdd45b7b3..e66651488 100644 --- a/OpenCL/inc_common.cl +++ b/OpenCL/inc_common.cl @@ -138,69 +138,6 @@ inline float get_entropy (const u32 *buf, const int elems) * vector functions */ -inline void truncate_block (u32x w[4], const u32 len) -{ - switch (len) - { - case 0: w[0] &= 0; - w[1] &= 0; - w[2] &= 0; - w[3] &= 0; - break; - case 1: w[0] &= 0x000000FF; - w[1] &= 0; - w[2] &= 0; - w[3] &= 0; - break; - case 2: w[0] &= 0x0000FFFF; - w[1] &= 0; - w[2] &= 0; - w[3] &= 0; - break; - case 3: w[0] &= 0x00FFFFFF; - w[1] &= 0; - w[2] &= 0; - w[3] &= 0; - break; - case 4: w[1] &= 0; - w[2] &= 0; - w[3] &= 0; - break; - case 5: w[1] &= 0x000000FF; - w[2] &= 0; - w[3] &= 0; - break; - case 6: w[1] &= 0x0000FFFF; - w[2] &= 0; - w[3] &= 0; - break; - case 7: w[1] &= 0x00FFFFFF; - w[2] &= 0; - w[3] &= 0; - break; - case 8: w[2] &= 0; - w[3] &= 0; - break; - case 9: w[2] &= 0x000000FF; - w[3] &= 0; - break; - case 10: w[2] &= 0x0000FFFF; - w[3] &= 0; - break; - case 11: w[2] &= 0x00FFFFFF; - w[3] &= 0; - break; - case 12: w[3] &= 0; - break; - case 13: w[3] &= 0x000000FF; - break; - case 14: w[3] &= 0x0000FFFF; - break; - case 15: w[3] &= 0x00FFFFFF; - break; - } -} - inline void truncate_block_4x4_le (u32x w0[4], const u32 len) { switch (len) @@ -16235,69 +16172,6 @@ inline void append_0x80_8x4_S (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], u32 w } } -inline void truncate_block_S (u32 w[4], const u32 len) -{ - switch (len) - { - case 0: w[0] &= 0; - w[1] &= 0; - w[2] &= 0; - w[3] &= 0; - break; - case 1: w[0] &= 0x000000FF; - w[1] &= 0; - w[2] &= 0; - w[3] &= 0; - break; - case 2: w[0] &= 0x0000FFFF; - w[1] &= 0; - w[2] &= 0; - w[3] &= 0; - break; - case 3: w[0] &= 0x00FFFFFF; - w[1] &= 0; - w[2] &= 0; - w[3] &= 0; - break; - case 4: w[1] &= 0; - w[2] &= 0; - w[3] &= 0; - break; - case 5: w[1] &= 0x000000FF; - w[2] &= 0; - w[3] &= 0; - break; - case 6: w[1] &= 0x0000FFFF; - w[2] &= 0; - w[3] &= 0; - break; - case 7: w[1] &= 0x00FFFFFF; - w[2] &= 0; - w[3] &= 0; - break; - case 8: w[2] &= 0; - w[3] &= 0; - break; - case 9: w[2] &= 0x000000FF; - w[3] &= 0; - break; - case 10: w[2] &= 0x0000FFFF; - w[3] &= 0; - break; - case 11: w[2] &= 0x00FFFFFF; - w[3] &= 0; - break; - case 12: w[3] &= 0; - break; - case 13: w[3] &= 0x000000FF; - break; - case 14: w[3] &= 0x0000FFFF; - break; - case 15: w[3] &= 0x00FFFFFF; - break; - } -} - inline void make_utf16be_S (const u32 in[4], u32 out1[4], u32 out2[4]) { #ifdef IS_NV diff --git a/OpenCL/m06800.cl b/OpenCL/m06800.cl index 61037ba29..333ec14b1 100644 --- a/OpenCL/m06800.cl +++ b/OpenCL/m06800.cl @@ -1556,7 +1556,7 @@ __kernel void m06800_comp (__global pw_t *pws, __global const kernel_rule_t *rul out[2] = swap32 (out[2]); out[3] = swap32 (out[3]); - truncate_block (out, salt_len); + truncate_block_4x4_le (out, salt_len); if ((out[0] == salt_buf[0]) && (out[1] == salt_buf[1]) diff --git a/OpenCL/m07400.cl b/OpenCL/m07400.cl index ba2b3b8d6..e3a9fbda5 100644 --- a/OpenCL/m07400.cl +++ b/OpenCL/m07400.cl @@ -868,7 +868,7 @@ __kernel void m07400_init (__global pw_t *pws, __global const kernel_rule_t *rul alt_result_tmp[6] = 0; alt_result_tmp[7] = 0; - truncate_block (alt_result_tmp, pw_len); + truncate_block_4x4_le (alt_result_tmp, pw_len); /* Add the key string. */ @@ -978,7 +978,7 @@ __kernel void m07400_init (__global pw_t *pws, __global const kernel_rule_t *rul bswap8 (p_bytes); - truncate_block (p_bytes, pw_len); + truncate_block_4x4_le (p_bytes, pw_len); tmps[gid].p_bytes[0] = p_bytes[0]; tmps[gid].p_bytes[1] = p_bytes[1]; @@ -1021,7 +1021,7 @@ __kernel void m07400_init (__global pw_t *pws, __global const kernel_rule_t *rul bswap8 (s_bytes); - truncate_block (s_bytes, salt_len); + truncate_block_4x4_le (s_bytes, salt_len); tmps[gid].s_bytes[0] = s_bytes[0]; tmps[gid].s_bytes[1] = s_bytes[1]; diff --git a/OpenCL/m13100_a0.cl b/OpenCL/m13100_a0.cl index 9858c8be8..e719cd696 100644 --- a/OpenCL/m13100_a0.cl +++ b/OpenCL/m13100_a0.cl @@ -481,7 +481,8 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat { j = rc4_next_16 (rc4_key, i, j, edata2, w0); i += 16; edata2 += 4; - truncate_block (w0, edata2_left & 0xf); + truncate_block_4x4_le (w0, edata2_left & 0xf); + append_0x80_1x4 (w0, edata2_left & 0xf); w3[2] = (64 + edata2_len) * 8; @@ -494,7 +495,8 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat j = rc4_next_16 (rc4_key, i, j, edata2, w0); i += 16; edata2 += 4; j = rc4_next_16 (rc4_key, i, j, edata2, w1); i += 16; edata2 += 4; - truncate_block (w1, edata2_left & 0xf); + truncate_block_4x4_le (w1, edata2_left & 0xf); + append_0x80_1x4 (w1, edata2_left & 0xf); w3[2] = (64 + edata2_len) * 8; @@ -508,7 +510,8 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat j = rc4_next_16 (rc4_key, i, j, edata2, w1); i += 16; edata2 += 4; j = rc4_next_16 (rc4_key, i, j, edata2, w2); i += 16; edata2 += 4; - truncate_block (w2, edata2_left & 0xf); + truncate_block_4x4_le (w2, edata2_left & 0xf); + append_0x80_1x4 (w2, edata2_left & 0xf); w3[2] = (64 + edata2_len) * 8; @@ -523,7 +526,8 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat j = rc4_next_16 (rc4_key, i, j, edata2, w2); i += 16; edata2 += 4; j = rc4_next_16 (rc4_key, i, j, edata2, w3); i += 16; edata2 += 4; - truncate_block (w3, edata2_left & 0xf); + truncate_block_4x4_le (w3, edata2_left & 0xf); + append_0x80_1x4 (w3, edata2_left & 0xf); if (edata2_left < 56) diff --git a/OpenCL/m13100_a1.cl b/OpenCL/m13100_a1.cl index ffa9e43fa..0b7e6dc42 100644 --- a/OpenCL/m13100_a1.cl +++ b/OpenCL/m13100_a1.cl @@ -479,7 +479,8 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat { j = rc4_next_16 (rc4_key, i, j, edata2, w0); i += 16; edata2 += 4; - truncate_block (w0, edata2_left & 0xf); + truncate_block_4x4_le (w0, edata2_left & 0xf); + append_0x80_1x4 (w0, edata2_left & 0xf); w3[2] = (64 + edata2_len) * 8; @@ -492,7 +493,8 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat j = rc4_next_16 (rc4_key, i, j, edata2, w0); i += 16; edata2 += 4; j = rc4_next_16 (rc4_key, i, j, edata2, w1); i += 16; edata2 += 4; - truncate_block (w1, edata2_left & 0xf); + truncate_block_4x4_le (w1, edata2_left & 0xf); + append_0x80_1x4 (w1, edata2_left & 0xf); w3[2] = (64 + edata2_len) * 8; @@ -506,7 +508,8 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat j = rc4_next_16 (rc4_key, i, j, edata2, w1); i += 16; edata2 += 4; j = rc4_next_16 (rc4_key, i, j, edata2, w2); i += 16; edata2 += 4; - truncate_block (w2, edata2_left & 0xf); + truncate_block_4x4_le (w2, edata2_left & 0xf); + append_0x80_1x4 (w2, edata2_left & 0xf); w3[2] = (64 + edata2_len) * 8; @@ -521,7 +524,8 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat j = rc4_next_16 (rc4_key, i, j, edata2, w2); i += 16; edata2 += 4; j = rc4_next_16 (rc4_key, i, j, edata2, w3); i += 16; edata2 += 4; - truncate_block (w3, edata2_left & 0xf); + truncate_block_4x4_le (w3, edata2_left & 0xf); + append_0x80_1x4 (w3, edata2_left & 0xf); if (edata2_left < 56) diff --git a/OpenCL/m13100_a3.cl b/OpenCL/m13100_a3.cl index e66756fac..bcba78c30 100644 --- a/OpenCL/m13100_a3.cl +++ b/OpenCL/m13100_a3.cl @@ -479,7 +479,8 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat { j = rc4_next_16 (rc4_key, i, j, edata2, w0); i += 16; edata2 += 4; - truncate_block (w0, edata2_left & 0xf); + truncate_block_4x4_le (w0, edata2_left & 0xf); + append_0x80_1x4 (w0, edata2_left & 0xf); w3[2] = (64 + edata2_len) * 8; @@ -492,7 +493,8 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat j = rc4_next_16 (rc4_key, i, j, edata2, w0); i += 16; edata2 += 4; j = rc4_next_16 (rc4_key, i, j, edata2, w1); i += 16; edata2 += 4; - truncate_block (w1, edata2_left & 0xf); + truncate_block_4x4_le (w1, edata2_left & 0xf); + append_0x80_1x4 (w1, edata2_left & 0xf); w3[2] = (64 + edata2_len) * 8; @@ -506,7 +508,8 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat j = rc4_next_16 (rc4_key, i, j, edata2, w1); i += 16; edata2 += 4; j = rc4_next_16 (rc4_key, i, j, edata2, w2); i += 16; edata2 += 4; - truncate_block (w2, edata2_left & 0xf); + truncate_block_4x4_le (w2, edata2_left & 0xf); + append_0x80_1x4 (w2, edata2_left & 0xf); w3[2] = (64 + edata2_len) * 8; @@ -521,7 +524,8 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat j = rc4_next_16 (rc4_key, i, j, edata2, w2); i += 16; edata2 += 4; j = rc4_next_16 (rc4_key, i, j, edata2, w3); i += 16; edata2 += 4; - truncate_block (w3, edata2_left & 0xf); + truncate_block_4x4_le (w3, edata2_left & 0xf); + append_0x80_1x4 (w3, edata2_left & 0xf); if (edata2_left < 56) diff --git a/OpenCL/m13400.cl b/OpenCL/m13400.cl index 103639121..d87f6f8f1 100644 --- a/OpenCL/m13400.cl +++ b/OpenCL/m13400.cl @@ -939,7 +939,7 @@ __kernel void m13400_comp (__global pw_t *pws, __global const kernel_rule_t *rul // we need to clear the buffer of the padding data - truncate_block (out, 16 - pad_byte); + truncate_block_4x4_le (out, 16 - pad_byte); // it's also a good point to push our 0x80 @@ -1146,7 +1146,7 @@ __kernel void m13400_comp (__global pw_t *pws, __global const kernel_rule_t *rul // we need to clear the buffer of the padding data - truncate_block (out, 16 - pad_byte); + truncate_block_4x4_le (out, 16 - pad_byte); // it's also a good point to push our 0x80