X-Git-Url: https://www.flypig.org.uk/git/?a=blobdiff_plain;f=OpenCL%2Fm06300.cl;h=f0831079ebae76c17ed7a2ccc024a1571f858d8a;hb=161a6eb4bc643d8e636e96eda613f5137d30da59;hp=593d31105ae0c97ec02a5528c4625e4b704cb52e;hpb=50f39b3563fa429dbec441288bd9a21fc328c21e;p=hashcat.git diff --git a/OpenCL/m06300.cl b/OpenCL/m06300.cl index 593d311..f083107 100644 --- a/OpenCL/m06300.cl +++ b/OpenCL/m06300.cl @@ -5,27 +5,16 @@ #define _MD5_ -#include "include/constants.h" -#include "include/kernel_vendor.h" +#include "inc_vendor.cl" +#include "inc_hash_constants.h" +#include "inc_hash_functions.cl" +#include "inc_types.cl" +#include "inc_common.cl" -#define DGST_R0 0 -#define DGST_R1 1 -#define DGST_R2 2 -#define DGST_R3 3 +#define COMPARE_S "inc_comp_single.cl" +#define COMPARE_M "inc_comp_multi.cl" -#include "include/kernel_functions.c" -#include "types_ocl.c" -#include "common.c" - -#ifdef VECT_SIZE1 -#define COMPARE_M "check_multi_vect1_comp4.c" -#endif - -#ifdef VECT_SIZE4 -#define COMPARE_M "check_multi_vect4_comp4.c" -#endif - -static void md5_transform (const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], u32 digest[4]) +void md5_transform (const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], u32 digest[4]) { u32 a = digest[0]; u32 b = digest[1]; @@ -49,8 +38,6 @@ static void md5_transform (const u32 w0[4], const u32 w1[4], const u32 w2[4], co u32 we_t = w3[2]; u32 wf_t = 0; - u32 tmp2; - MD5_STEP (MD5_Fo, a, b, c, d, w0_t, MD5C00, MD5S00); MD5_STEP (MD5_Fo, d, a, b, c, w1_t, MD5C01, MD5S01); MD5_STEP (MD5_Fo, c, d, a, b, w2_t, MD5C02, MD5S02); @@ -85,22 +72,22 @@ static void md5_transform (const u32 w0[4], const u32 w1[4], const u32 w2[4], co MD5_STEP (MD5_Go, c, d, a, b, w7_t, MD5C1e, MD5S12); MD5_STEP (MD5_Go, b, c, d, a, wc_t, MD5C1f, MD5S13); - MD5_STEP (MD5_H1, a, b, c, d, w5_t, MD5C20, MD5S20); - MD5_STEP (MD5_H2, d, a, b, c, w8_t, MD5C21, MD5S21); - MD5_STEP (MD5_H1, c, d, a, b, wb_t, MD5C22, MD5S22); - MD5_STEP (MD5_H2, b, c, d, a, we_t, MD5C23, MD5S23); - MD5_STEP (MD5_H1, a, b, c, d, w1_t, MD5C24, MD5S20); - MD5_STEP (MD5_H2, d, a, b, c, w4_t, MD5C25, MD5S21); - MD5_STEP (MD5_H1, c, d, a, b, w7_t, MD5C26, MD5S22); - MD5_STEP (MD5_H2, b, c, d, a, wa_t, MD5C27, MD5S23); - MD5_STEP (MD5_H1, a, b, c, d, wd_t, MD5C28, MD5S20); - MD5_STEP (MD5_H2, d, a, b, c, w0_t, MD5C29, MD5S21); - MD5_STEP (MD5_H1, c, d, a, b, w3_t, MD5C2a, MD5S22); - MD5_STEP (MD5_H2, b, c, d, a, w6_t, MD5C2b, MD5S23); - MD5_STEP (MD5_H1, a, b, c, d, w9_t, MD5C2c, MD5S20); - MD5_STEP (MD5_H2, d, a, b, c, wc_t, MD5C2d, MD5S21); - MD5_STEP (MD5_H1, c, d, a, b, wf_t, MD5C2e, MD5S22); - MD5_STEP (MD5_H2, b, c, d, a, w2_t, MD5C2f, MD5S23); + MD5_STEP (MD5_H , a, b, c, d, w5_t, MD5C20, MD5S20); + MD5_STEP (MD5_H , d, a, b, c, w8_t, MD5C21, MD5S21); + MD5_STEP (MD5_H , c, d, a, b, wb_t, MD5C22, MD5S22); + MD5_STEP (MD5_H , b, c, d, a, we_t, MD5C23, MD5S23); + MD5_STEP (MD5_H , a, b, c, d, w1_t, MD5C24, MD5S20); + MD5_STEP (MD5_H , d, a, b, c, w4_t, MD5C25, MD5S21); + MD5_STEP (MD5_H , c, d, a, b, w7_t, MD5C26, MD5S22); + MD5_STEP (MD5_H , b, c, d, a, wa_t, MD5C27, MD5S23); + MD5_STEP (MD5_H , a, b, c, d, wd_t, MD5C28, MD5S20); + MD5_STEP (MD5_H , d, a, b, c, w0_t, MD5C29, MD5S21); + MD5_STEP (MD5_H , c, d, a, b, w3_t, MD5C2a, MD5S22); + MD5_STEP (MD5_H , b, c, d, a, w6_t, MD5C2b, MD5S23); + MD5_STEP (MD5_H , a, b, c, d, w9_t, MD5C2c, MD5S20); + MD5_STEP (MD5_H , d, a, b, c, wc_t, MD5C2d, MD5S21); + MD5_STEP (MD5_H , c, d, a, b, wf_t, MD5C2e, MD5S22); + MD5_STEP (MD5_H , b, c, d, a, w2_t, MD5C2f, MD5S23); MD5_STEP (MD5_I , a, b, c, d, w0_t, MD5C30, MD5S30); MD5_STEP (MD5_I , d, a, b, c, w7_t, MD5C31, MD5S31); @@ -125,1222 +112,323 @@ static void md5_transform (const u32 w0[4], const u32 w1[4], const u32 w2[4], co digest[3] += d; } -static void memcat16 (u32 block0[4], u32 block1[4], u32 block2[4], u32 block3[4], const u32 block_len, const u32 append[4]) +void memcat16 (u32 block0[4], u32 block1[4], u32 block2[4], u32 block3[4], const u32 block_len, const u32 append[4]) { - switch (block_len) - { - case 0: - block0[0] = append[0]; - block0[1] = append[1]; - block0[2] = append[2]; - block0[3] = append[3]; - break; - - case 1: - block0[0] = block0[0] | append[0] << 8; - block0[1] = append[0] >> 24 | append[1] << 8; - block0[2] = append[1] >> 24 | append[2] << 8; - block0[3] = append[2] >> 24 | append[3] << 8; - block1[0] = append[3] >> 24; - break; - - case 2: - block0[0] = block0[0] | append[0] << 16; - block0[1] = append[0] >> 16 | append[1] << 16; - block0[2] = append[1] >> 16 | append[2] << 16; - block0[3] = append[2] >> 16 | append[3] << 16; - block1[0] = append[3] >> 16; - break; - - case 3: - block0[0] = block0[0] | append[0] << 24; - block0[1] = append[0] >> 8 | append[1] << 24; - block0[2] = append[1] >> 8 | append[2] << 24; - block0[3] = append[2] >> 8 | append[3] << 24; - block1[0] = append[3] >> 8; - break; - - case 4: - block0[1] = append[0]; - block0[2] = append[1]; - block0[3] = append[2]; - block1[0] = append[3]; - break; - - case 5: - block0[1] = block0[1] | append[0] << 8; - block0[2] = append[0] >> 24 | append[1] << 8; - block0[3] = append[1] >> 24 | append[2] << 8; - block1[0] = append[2] >> 24 | append[3] << 8; - block1[1] = append[3] >> 24; - break; - - case 6: - block0[1] = block0[1] | append[0] << 16; - block0[2] = append[0] >> 16 | append[1] << 16; - block0[3] = append[1] >> 16 | append[2] << 16; - block1[0] = append[2] >> 16 | append[3] << 16; - block1[1] = append[3] >> 16; - break; - - case 7: - block0[1] = block0[1] | append[0] << 24; - block0[2] = append[0] >> 8 | append[1] << 24; - block0[3] = append[1] >> 8 | append[2] << 24; - block1[0] = append[2] >> 8 | append[3] << 24; - block1[1] = append[3] >> 8; - break; - - case 8: - block0[2] = append[0]; - block0[3] = append[1]; - block1[0] = append[2]; - block1[1] = append[3]; - break; - - case 9: - block0[2] = block0[2] | append[0] << 8; - block0[3] = append[0] >> 24 | append[1] << 8; - block1[0] = append[1] >> 24 | append[2] << 8; - block1[1] = append[2] >> 24 | append[3] << 8; - block1[2] = append[3] >> 24; - break; - - case 10: - block0[2] = block0[2] | append[0] << 16; - block0[3] = append[0] >> 16 | append[1] << 16; - block1[0] = append[1] >> 16 | append[2] << 16; - block1[1] = append[2] >> 16 | append[3] << 16; - block1[2] = append[3] >> 16; - break; - - case 11: - block0[2] = block0[2] | append[0] << 24; - block0[3] = append[0] >> 8 | append[1] << 24; - block1[0] = append[1] >> 8 | append[2] << 24; - block1[1] = append[2] >> 8 | append[3] << 24; - block1[2] = append[3] >> 8; - break; - - case 12: - block0[3] = append[0]; - block1[0] = append[1]; - block1[1] = append[2]; - block1[2] = append[3]; - break; - - case 13: - block0[3] = block0[3] | append[0] << 8; - block1[0] = append[0] >> 24 | append[1] << 8; - block1[1] = append[1] >> 24 | append[2] << 8; - block1[2] = append[2] >> 24 | append[3] << 8; - block1[3] = append[3] >> 24; - break; - - case 14: - block0[3] = block0[3] | append[0] << 16; - block1[0] = append[0] >> 16 | append[1] << 16; - block1[1] = append[1] >> 16 | append[2] << 16; - block1[2] = append[2] >> 16 | append[3] << 16; - block1[3] = append[3] >> 16; - break; - - case 15: - block0[3] = block0[3] | append[0] << 24; - block1[0] = append[0] >> 8 | append[1] << 24; - block1[1] = append[1] >> 8 | append[2] << 24; - block1[2] = append[2] >> 8 | append[3] << 24; - block1[3] = append[3] >> 8; - break; - - case 16: - block1[0] = append[0]; - block1[1] = append[1]; - block1[2] = append[2]; - block1[3] = append[3]; - break; - - case 17: - block1[0] = block1[0] | append[0] << 8; - block1[1] = append[0] >> 24 | append[1] << 8; - block1[2] = append[1] >> 24 | append[2] << 8; - block1[3] = append[2] >> 24 | append[3] << 8; - block2[0] = append[3] >> 24; - break; - - case 18: - block1[0] = block1[0] | append[0] << 16; - block1[1] = append[0] >> 16 | append[1] << 16; - block1[2] = append[1] >> 16 | append[2] << 16; - block1[3] = append[2] >> 16 | append[3] << 16; - block2[0] = append[3] >> 16; - break; - - case 19: - block1[0] = block1[0] | append[0] << 24; - block1[1] = append[0] >> 8 | append[1] << 24; - block1[2] = append[1] >> 8 | append[2] << 24; - block1[3] = append[2] >> 8 | append[3] << 24; - block2[0] = append[3] >> 8; - break; - - case 20: - block1[1] = append[0]; - block1[2] = append[1]; - block1[3] = append[2]; - block2[0] = append[3]; - break; - - case 21: - block1[1] = block1[1] | append[0] << 8; - block1[2] = append[0] >> 24 | append[1] << 8; - block1[3] = append[1] >> 24 | append[2] << 8; - block2[0] = append[2] >> 24 | append[3] << 8; - block2[1] = append[3] >> 24; - break; - - case 22: - block1[1] = block1[1] | append[0] << 16; - block1[2] = append[0] >> 16 | append[1] << 16; - block1[3] = append[1] >> 16 | append[2] << 16; - block2[0] = append[2] >> 16 | append[3] << 16; - block2[1] = append[3] >> 16; - break; - - case 23: - block1[1] = block1[1] | append[0] << 24; - block1[2] = append[0] >> 8 | append[1] << 24; - block1[3] = append[1] >> 8 | append[2] << 24; - block2[0] = append[2] >> 8 | append[3] << 24; - block2[1] = append[3] >> 8; - break; - - case 24: - block1[2] = append[0]; - block1[3] = append[1]; - block2[0] = append[2]; - block2[1] = append[3]; - break; - - case 25: - block1[2] = block1[2] | append[0] << 8; - block1[3] = append[0] >> 24 | append[1] << 8; - block2[0] = append[1] >> 24 | append[2] << 8; - block2[1] = append[2] >> 24 | append[3] << 8; - block2[2] = append[3] >> 24; - break; - - case 26: - block1[2] = block1[2] | append[0] << 16; - block1[3] = append[0] >> 16 | append[1] << 16; - block2[0] = append[1] >> 16 | append[2] << 16; - block2[1] = append[2] >> 16 | append[3] << 16; - block2[2] = append[3] >> 16; - break; - - case 27: - block1[2] = block1[2] | append[0] << 24; - block1[3] = append[0] >> 8 | append[1] << 24; - block2[0] = append[1] >> 8 | append[2] << 24; - block2[1] = append[2] >> 8 | append[3] << 24; - block2[2] = append[3] >> 8; - break; - - case 28: - block1[3] = append[0]; - block2[0] = append[1]; - block2[1] = append[2]; - block2[2] = append[3]; - break; - - case 29: - block1[3] = block1[3] | append[0] << 8; - block2[0] = append[0] >> 24 | append[1] << 8; - block2[1] = append[1] >> 24 | append[2] << 8; - block2[2] = append[2] >> 24 | append[3] << 8; - block2[3] = append[3] >> 24; - break; - - case 30: - block1[3] = block1[3] | append[0] << 16; - block2[0] = append[0] >> 16 | append[1] << 16; - block2[1] = append[1] >> 16 | append[2] << 16; - block2[2] = append[2] >> 16 | append[3] << 16; - block2[3] = append[3] >> 16; - break; - - case 31: - block1[3] = block1[3] | append[0] << 24; - block2[0] = append[0] >> 8 | append[1] << 24; - block2[1] = append[1] >> 8 | append[2] << 24; - block2[2] = append[2] >> 8 | append[3] << 24; - block2[3] = append[3] >> 8; - break; - - case 32: - block2[0] = append[0]; - block2[1] = append[1]; - block2[2] = append[2]; - block2[3] = append[3]; - break; - - case 33: - block2[0] = block2[0] | append[0] << 8; - block2[1] = append[0] >> 24 | append[1] << 8; - block2[2] = append[1] >> 24 | append[2] << 8; - block2[3] = append[2] >> 24 | append[3] << 8; - block3[0] = append[3] >> 24; - break; - - case 34: - block2[0] = block2[0] | append[0] << 16; - block2[1] = append[0] >> 16 | append[1] << 16; - block2[2] = append[1] >> 16 | append[2] << 16; - block2[3] = append[2] >> 16 | append[3] << 16; - block3[0] = append[3] >> 16; - break; - - case 35: - block2[0] = block2[0] | append[0] << 24; - block2[1] = append[0] >> 8 | append[1] << 24; - block2[2] = append[1] >> 8 | append[2] << 24; - block2[3] = append[2] >> 8 | append[3] << 24; - block3[0] = append[3] >> 8; - break; - - case 36: - block2[1] = append[0]; - block2[2] = append[1]; - block2[3] = append[2]; - block3[0] = append[3]; - break; - - case 37: - block2[1] = block2[1] | append[0] << 8; - block2[2] = append[0] >> 24 | append[1] << 8; - block2[3] = append[1] >> 24 | append[2] << 8; - block3[0] = append[2] >> 24 | append[3] << 8; - block3[1] = append[3] >> 24; - break; - - case 38: - block2[1] = block2[1] | append[0] << 16; - block2[2] = append[0] >> 16 | append[1] << 16; - block2[3] = append[1] >> 16 | append[2] << 16; - block3[0] = append[2] >> 16 | append[3] << 16; - block3[1] = append[3] >> 16; - break; - - case 39: - block2[1] = block2[1] | append[0] << 24; - block2[2] = append[0] >> 8 | append[1] << 24; - block2[3] = append[1] >> 8 | append[2] << 24; - block3[0] = append[2] >> 8 | append[3] << 24; - block3[1] = append[3] >> 8; - break; - - case 40: - block2[2] = append[0]; - block2[3] = append[1]; - block3[0] = append[2]; - block3[1] = append[3]; - break; + u32 tmp0; + u32 tmp1; + u32 tmp2; + u32 tmp3; + u32 tmp4; - case 41: - block2[2] = block2[2] | append[0] << 8; - block2[3] = append[0] >> 24 | append[1] << 8; - block3[0] = append[1] >> 24 | append[2] << 8; - block3[1] = append[2] >> 24 | append[3] << 8; - block3[2] = append[3] >> 24; - break; + #if defined IS_AMD || defined IS_GENERIC - case 42: - block2[2] = block2[2] | append[0] << 16; - block2[3] = append[0] >> 16 | append[1] << 16; - block3[0] = append[1] >> 16 | append[2] << 16; - block3[1] = append[2] >> 16 | append[3] << 16; - block3[2] = append[3] >> 16; - break; + const int offset_minus_4 = 4 - (block_len & 3); - case 43: - block2[2] = block2[2] | append[0] << 24; - block2[3] = append[0] >> 8 | append[1] << 24; - block3[0] = append[1] >> 8 | append[2] << 24; - block3[1] = append[2] >> 8 | append[3] << 24; - block3[2] = append[3] >> 8; - break; + tmp0 = amd_bytealign (append[0], 0, offset_minus_4); + tmp1 = amd_bytealign (append[1], append[0], offset_minus_4); + tmp2 = amd_bytealign (append[2], append[1], offset_minus_4); + tmp3 = amd_bytealign (append[3], append[2], offset_minus_4); + tmp4 = amd_bytealign ( 0, append[3], offset_minus_4); - case 44: - block2[3] = append[0]; - block3[0] = append[1]; - block3[1] = append[2]; - block3[2] = append[3]; - break; + const u32 mod = block_len & 3; - case 45: - block2[3] = block2[3] | append[0] << 8; - block3[0] = append[0] >> 24 | append[1] << 8; - block3[1] = append[1] >> 24 | append[2] << 8; - block3[2] = append[2] >> 24 | append[3] << 8; - block3[3] = append[3] >> 24; - break; - - case 46: - block2[3] = block2[3] | append[0] << 16; - block3[0] = append[0] >> 16 | append[1] << 16; - block3[1] = append[1] >> 16 | append[2] << 16; - block3[2] = append[2] >> 16 | append[3] << 16; - block3[3] = append[3] >> 16; - break; - - case 47: - block2[3] = block2[3] | append[0] << 24; - block3[0] = append[0] >> 8 | append[1] << 24; - block3[1] = append[1] >> 8 | append[2] << 24; - block3[2] = append[2] >> 8 | append[3] << 24; - block3[3] = append[3] >> 8; - break; - - case 48: - block3[0] = append[0]; - block3[1] = append[1]; - block3[2] = append[2]; - block3[3] = append[3]; - break; + if (mod == 0) + { + tmp0 = tmp1; + tmp1 = tmp2; + tmp2 = tmp3; + tmp3 = tmp4; + tmp4 = 0; + } - case 49: - block3[0] = block3[0] | append[0] << 8; - block3[1] = append[0] >> 24 | append[1] << 8; - block3[2] = append[1] >> 24 | append[2] << 8; - block3[3] = append[2] >> 24 | append[3] << 8; - break; + #endif - case 50: - block3[0] = block3[0] | append[0] << 16; - block3[1] = append[0] >> 16 | append[1] << 16; - block3[2] = append[1] >> 16 | append[2] << 16; - block3[3] = append[2] >> 16 | append[3] << 16; - break; + #ifdef IS_NV - case 51: - block3[0] = block3[0] | append[0] << 24; - block3[1] = append[0] >> 8 | append[1] << 24; - block3[2] = append[1] >> 8 | append[2] << 24; - block3[3] = append[2] >> 8 | append[3] << 24; - break; + const int offset_minus_4 = 4 - (block_len & 3); - case 52: - block3[1] = append[0]; - block3[2] = append[1]; - block3[3] = append[2]; - break; + const int selector = (0x76543210 >> (offset_minus_4 * 4)) & 0xffff; - case 53: - block3[1] = block3[1] | append[0] << 8; - block3[2] = append[0] >> 24 | append[1] << 8; - block3[3] = append[1] >> 24 | append[2] << 8; - break; + tmp0 = __byte_perm ( 0, append[0], selector); + tmp1 = __byte_perm (append[0], append[1], selector); + tmp2 = __byte_perm (append[1], append[2], selector); + tmp3 = __byte_perm (append[2], append[3], selector); + tmp4 = __byte_perm (append[3], 0, selector); - case 54: - block3[1] = block3[1] | append[0] << 16; - block3[2] = append[0] >> 16 | append[1] << 16; - block3[3] = append[1] >> 16 | append[2] << 16; - break; + #endif - case 55: - block3[1] = block3[1] | append[0] << 24; - block3[2] = append[0] >> 8 | append[1] << 24; - block3[3] = append[1] >> 8 | append[2] << 24; - break; + const u32 div = block_len / 4; - case 56: - block3[2] = append[0]; - block3[3] = append[1]; - break; + switch (div) + { + case 0: block0[0] |= tmp0; + block0[1] = tmp1; + block0[2] = tmp2; + block0[3] = tmp3; + block1[0] = tmp4; + break; + case 1: block0[1] |= tmp0; + block0[2] = tmp1; + block0[3] = tmp2; + block1[0] = tmp3; + block1[1] = tmp4; + break; + case 2: block0[2] |= tmp0; + block0[3] = tmp1; + block1[0] = tmp2; + block1[1] = tmp3; + block1[2] = tmp4; + break; + case 3: block0[3] |= tmp0; + block1[0] = tmp1; + block1[1] = tmp2; + block1[2] = tmp3; + block1[3] = tmp4; + break; + case 4: block1[0] |= tmp0; + block1[1] = tmp1; + block1[2] = tmp2; + block1[3] = tmp3; + block2[0] = tmp4; + break; + case 5: block1[1] |= tmp0; + block1[2] = tmp1; + block1[3] = tmp2; + block2[0] = tmp3; + block2[1] = tmp4; + break; + case 6: block1[2] |= tmp0; + block1[3] = tmp1; + block2[0] = tmp2; + block2[1] = tmp3; + block2[2] = tmp4; + break; + case 7: block1[3] |= tmp0; + block2[0] = tmp1; + block2[1] = tmp2; + block2[2] = tmp3; + block2[3] = tmp4; + break; + case 8: block2[0] |= tmp0; + block2[1] = tmp1; + block2[2] = tmp2; + block2[3] = tmp3; + block3[0] = tmp4; + break; + case 9: block2[1] |= tmp0; + block2[2] = tmp1; + block2[3] = tmp2; + block3[0] = tmp3; + block3[1] = tmp4; + break; } } -static void memcat16_x80 (u32 block0[4], u32 block1[4], u32 block2[4], u32 block3[4], const u32 block_len, const u32 append[4]) +void memcat16_x80 (u32 block0[4], u32 block1[4], u32 block2[4], u32 block3[4], const u32 block_len, const u32 append[4]) { - switch (block_len) - { - case 0: - block0[0] = append[0]; - block0[1] = append[1]; - block0[2] = append[2]; - block0[3] = append[3]; - block1[0] = 0x80; - break; - - case 1: - block0[0] = block0[0] | append[0] << 8; - block0[1] = append[0] >> 24 | append[1] << 8; - block0[2] = append[1] >> 24 | append[2] << 8; - block0[3] = append[2] >> 24 | append[3] << 8; - block1[0] = append[3] >> 24 | 0x80u << 8; - break; - - case 2: - block0[0] = block0[0] | append[0] << 16; - block0[1] = append[0] >> 16 | append[1] << 16; - block0[2] = append[1] >> 16 | append[2] << 16; - block0[3] = append[2] >> 16 | append[3] << 16; - block1[0] = append[3] >> 16 | 0x80u << 16; - break; - - case 3: - block0[0] = block0[0] | append[0] << 24; - block0[1] = append[0] >> 8 | append[1] << 24; - block0[2] = append[1] >> 8 | append[2] << 24; - block0[3] = append[2] >> 8 | append[3] << 24; - block1[0] = append[3] >> 8 | 0x80u << 24; - break; - - case 4: - block0[1] = append[0]; - block0[2] = append[1]; - block0[3] = append[2]; - block1[0] = append[3]; - block1[1] = 0x80; - break; - - case 5: - block0[1] = block0[1] | append[0] << 8; - block0[2] = append[0] >> 24 | append[1] << 8; - block0[3] = append[1] >> 24 | append[2] << 8; - block1[0] = append[2] >> 24 | append[3] << 8; - block1[1] = append[3] >> 24 | 0x80u << 8; - break; - - case 6: - block0[1] = block0[1] | append[0] << 16; - block0[2] = append[0] >> 16 | append[1] << 16; - block0[3] = append[1] >> 16 | append[2] << 16; - block1[0] = append[2] >> 16 | append[3] << 16; - block1[1] = append[3] >> 16 | 0x80u << 16; - break; - - case 7: - block0[1] = block0[1] | append[0] << 24; - block0[2] = append[0] >> 8 | append[1] << 24; - block0[3] = append[1] >> 8 | append[2] << 24; - block1[0] = append[2] >> 8 | append[3] << 24; - block1[1] = append[3] >> 8 | 0x80u << 24; - break; - - case 8: - block0[2] = append[0]; - block0[3] = append[1]; - block1[0] = append[2]; - block1[1] = append[3]; - block1[2] = 0x80; - break; - - case 9: - block0[2] = block0[2] | append[0] << 8; - block0[3] = append[0] >> 24 | append[1] << 8; - block1[0] = append[1] >> 24 | append[2] << 8; - block1[1] = append[2] >> 24 | append[3] << 8; - block1[2] = append[3] >> 24 | 0x80u << 8; - break; - - case 10: - block0[2] = block0[2] | append[0] << 16; - block0[3] = append[0] >> 16 | append[1] << 16; - block1[0] = append[1] >> 16 | append[2] << 16; - block1[1] = append[2] >> 16 | append[3] << 16; - block1[2] = append[3] >> 16 | 0x80u << 16; - break; - - case 11: - block0[2] = block0[2] | append[0] << 24; - block0[3] = append[0] >> 8 | append[1] << 24; - block1[0] = append[1] >> 8 | append[2] << 24; - block1[1] = append[2] >> 8 | append[3] << 24; - block1[2] = append[3] >> 8 | 0x80u << 24; - break; - - case 12: - block0[3] = append[0]; - block1[0] = append[1]; - block1[1] = append[2]; - block1[2] = append[3]; - block1[3] = 0x80; - break; - - case 13: - block0[3] = block0[3] | append[0] << 8; - block1[0] = append[0] >> 24 | append[1] << 8; - block1[1] = append[1] >> 24 | append[2] << 8; - block1[2] = append[2] >> 24 | append[3] << 8; - block1[3] = append[3] >> 24 | 0x80u << 8; - break; - - case 14: - block0[3] = block0[3] | append[0] << 16; - block1[0] = append[0] >> 16 | append[1] << 16; - block1[1] = append[1] >> 16 | append[2] << 16; - block1[2] = append[2] >> 16 | append[3] << 16; - block1[3] = append[3] >> 16 | 0x80u << 16; - break; - - case 15: - block0[3] = block0[3] | append[0] << 24; - block1[0] = append[0] >> 8 | append[1] << 24; - block1[1] = append[1] >> 8 | append[2] << 24; - block1[2] = append[2] >> 8 | append[3] << 24; - block1[3] = append[3] >> 8 | 0x80u << 24; - break; - - case 16: - block1[0] = append[0]; - block1[1] = append[1]; - block1[2] = append[2]; - block1[3] = append[3]; - block2[0] = 0x80; - break; - - case 17: - block1[0] = block1[0] | append[0] << 8; - block1[1] = append[0] >> 24 | append[1] << 8; - block1[2] = append[1] >> 24 | append[2] << 8; - block1[3] = append[2] >> 24 | append[3] << 8; - block2[0] = append[3] >> 24 | 0x80u << 8; - break; - - case 18: - block1[0] = block1[0] | append[0] << 16; - block1[1] = append[0] >> 16 | append[1] << 16; - block1[2] = append[1] >> 16 | append[2] << 16; - block1[3] = append[2] >> 16 | append[3] << 16; - block2[0] = append[3] >> 16 | 0x80u << 16; - break; - - case 19: - block1[0] = block1[0] | append[0] << 24; - block1[1] = append[0] >> 8 | append[1] << 24; - block1[2] = append[1] >> 8 | append[2] << 24; - block1[3] = append[2] >> 8 | append[3] << 24; - block2[0] = append[3] >> 8 | 0x80u << 24; - break; - - case 20: - block1[1] = append[0]; - block1[2] = append[1]; - block1[3] = append[2]; - block2[0] = append[3]; - block2[1] = 0x80; - break; - - case 21: - block1[1] = block1[1] | append[0] << 8; - block1[2] = append[0] >> 24 | append[1] << 8; - block1[3] = append[1] >> 24 | append[2] << 8; - block2[0] = append[2] >> 24 | append[3] << 8; - block2[1] = append[3] >> 24 | 0x80u << 8; - break; - - case 22: - block1[1] = block1[1] | append[0] << 16; - block1[2] = append[0] >> 16 | append[1] << 16; - block1[3] = append[1] >> 16 | append[2] << 16; - block2[0] = append[2] >> 16 | append[3] << 16; - block2[1] = append[3] >> 16 | 0x80u << 16; - break; - - case 23: - block1[1] = block1[1] | append[0] << 24; - block1[2] = append[0] >> 8 | append[1] << 24; - block1[3] = append[1] >> 8 | append[2] << 24; - block2[0] = append[2] >> 8 | append[3] << 24; - block2[1] = append[3] >> 8 | 0x80u << 24; - break; - - case 24: - block1[2] = append[0]; - block1[3] = append[1]; - block2[0] = append[2]; - block2[1] = append[3]; - block2[2] = 0x80; - break; - - case 25: - block1[2] = block1[2] | append[0] << 8; - block1[3] = append[0] >> 24 | append[1] << 8; - block2[0] = append[1] >> 24 | append[2] << 8; - block2[1] = append[2] >> 24 | append[3] << 8; - block2[2] = append[3] >> 24 | 0x80u << 8; - break; - - case 26: - block1[2] = block1[2] | append[0] << 16; - block1[3] = append[0] >> 16 | append[1] << 16; - block2[0] = append[1] >> 16 | append[2] << 16; - block2[1] = append[2] >> 16 | append[3] << 16; - block2[2] = append[3] >> 16 | 0x80u << 16; - break; - - case 27: - block1[2] = block1[2] | append[0] << 24; - block1[3] = append[0] >> 8 | append[1] << 24; - block2[0] = append[1] >> 8 | append[2] << 24; - block2[1] = append[2] >> 8 | append[3] << 24; - block2[2] = append[3] >> 8 | 0x80u << 24; - break; - - case 28: - block1[3] = append[0]; - block2[0] = append[1]; - block2[1] = append[2]; - block2[2] = append[3]; - block2[3] = 0x80; - break; - - case 29: - block1[3] = block1[3] | append[0] << 8; - block2[0] = append[0] >> 24 | append[1] << 8; - block2[1] = append[1] >> 24 | append[2] << 8; - block2[2] = append[2] >> 24 | append[3] << 8; - block2[3] = append[3] >> 24 | 0x80u << 8; - break; - - case 30: - block1[3] = block1[3] | append[0] << 16; - block2[0] = append[0] >> 16 | append[1] << 16; - block2[1] = append[1] >> 16 | append[2] << 16; - block2[2] = append[2] >> 16 | append[3] << 16; - block2[3] = append[3] >> 16 | 0x80u << 16; - break; - - case 31: - block1[3] = block1[3] | append[0] << 24; - block2[0] = append[0] >> 8 | append[1] << 24; - block2[1] = append[1] >> 8 | append[2] << 24; - block2[2] = append[2] >> 8 | append[3] << 24; - block2[3] = append[3] >> 8 | 0x80u << 24; - break; - - case 32: - block2[0] = append[0]; - block2[1] = append[1]; - block2[2] = append[2]; - block2[3] = append[3]; - block3[0] = 0x80; - break; - - case 33: - block2[0] = block2[0] | append[0] << 8; - block2[1] = append[0] >> 24 | append[1] << 8; - block2[2] = append[1] >> 24 | append[2] << 8; - block2[3] = append[2] >> 24 | append[3] << 8; - block3[0] = append[3] >> 24 | 0x80u << 8; - break; - - case 34: - block2[0] = block2[0] | append[0] << 16; - block2[1] = append[0] >> 16 | append[1] << 16; - block2[2] = append[1] >> 16 | append[2] << 16; - block2[3] = append[2] >> 16 | append[3] << 16; - block3[0] = append[3] >> 16 | 0x80u << 16; - break; - - case 35: - block2[0] = block2[0] | append[0] << 24; - block2[1] = append[0] >> 8 | append[1] << 24; - block2[2] = append[1] >> 8 | append[2] << 24; - block2[3] = append[2] >> 8 | append[3] << 24; - block3[0] = append[3] >> 8 | 0x80u << 24; - break; - - case 36: - block2[1] = append[0]; - block2[2] = append[1]; - block2[3] = append[2]; - block3[0] = append[3]; - block3[1] = 0x80; - break; - - case 37: - block2[1] = block2[1] | append[0] << 8; - block2[2] = append[0] >> 24 | append[1] << 8; - block2[3] = append[1] >> 24 | append[2] << 8; - block3[0] = append[2] >> 24 | append[3] << 8; - block3[1] = append[3] >> 24 | 0x80u << 8; - break; - - case 38: - block2[1] = block2[1] | append[0] << 16; - block2[2] = append[0] >> 16 | append[1] << 16; - block2[3] = append[1] >> 16 | append[2] << 16; - block3[0] = append[2] >> 16 | append[3] << 16; - block3[1] = append[3] >> 16 | 0x80u << 16; - break; - - case 39: - block2[1] = block2[1] | append[0] << 24; - block2[2] = append[0] >> 8 | append[1] << 24; - block2[3] = append[1] >> 8 | append[2] << 24; - block3[0] = append[2] >> 8 | append[3] << 24; - block3[1] = append[3] >> 8 | 0x80u << 24; - break; - - case 40: - block2[2] = append[0]; - block2[3] = append[1]; - block3[0] = append[2]; - block3[1] = append[3]; - block3[2] = 0x80; - break; - - case 41: - block2[2] = block2[2] | append[0] << 8; - block2[3] = append[0] >> 24 | append[1] << 8; - block3[0] = append[1] >> 24 | append[2] << 8; - block3[1] = append[2] >> 24 | append[3] << 8; - block3[2] = append[3] >> 24 | 0x80u << 8; - break; + u32 tmp0; + u32 tmp1; + u32 tmp2; + u32 tmp3; + u32 tmp4; - case 42: - block2[2] = block2[2] | append[0] << 16; - block2[3] = append[0] >> 16 | append[1] << 16; - block3[0] = append[1] >> 16 | append[2] << 16; - block3[1] = append[2] >> 16 | append[3] << 16; - block3[2] = append[3] >> 16 | 0x80u << 16; - break; + #if defined IS_AMD || defined IS_GENERIC - case 43: - block2[2] = block2[2] | append[0] << 24; - block2[3] = append[0] >> 8 | append[1] << 24; - block3[0] = append[1] >> 8 | append[2] << 24; - block3[1] = append[2] >> 8 | append[3] << 24; - block3[2] = append[3] >> 8 | 0x80u << 24; - break; + const int offset_minus_4 = 4 - (block_len & 3); - case 44: - block2[3] = append[0]; - block3[0] = append[1]; - block3[1] = append[2]; - block3[2] = append[3]; - block3[3] = 0x80; - break; + tmp0 = amd_bytealign (append[0], 0, offset_minus_4); + tmp1 = amd_bytealign (append[1], append[0], offset_minus_4); + tmp2 = amd_bytealign (append[2], append[1], offset_minus_4); + tmp3 = amd_bytealign (append[3], append[2], offset_minus_4); + tmp4 = amd_bytealign ( 0x80, append[3], offset_minus_4); - case 45: - block2[3] = block2[3] | append[0] << 8; - block3[0] = append[0] >> 24 | append[1] << 8; - block3[1] = append[1] >> 24 | append[2] << 8; - block3[2] = append[2] >> 24 | append[3] << 8; - block3[3] = append[3] >> 24 | 0x80u << 8; - break; - - case 46: - block2[3] = block2[3] | append[0] << 16; - block3[0] = append[0] >> 16 | append[1] << 16; - block3[1] = append[1] >> 16 | append[2] << 16; - block3[2] = append[2] >> 16 | append[3] << 16; - block3[3] = append[3] >> 16 | 0x80u << 16; - break; - - case 47: - block2[3] = block2[3] | append[0] << 24; - block3[0] = append[0] >> 8 | append[1] << 24; - block3[1] = append[1] >> 8 | append[2] << 24; - block3[2] = append[2] >> 8 | append[3] << 24; - block3[3] = append[3] >> 8 | 0x80u << 24; - break; + const u32 mod = block_len & 3; - case 48: - block3[0] = append[0]; - block3[1] = append[1]; - block3[2] = append[2]; - block3[3] = append[3]; - break; + if (mod == 0) + { + tmp0 = tmp1; + tmp1 = tmp2; + tmp2 = tmp3; + tmp3 = tmp4; + tmp4 = 0x80; + } - case 49: - block3[0] = block3[0] | append[0] << 8; - block3[1] = append[0] >> 24 | append[1] << 8; - block3[2] = append[1] >> 24 | append[2] << 8; - block3[3] = append[2] >> 24 | append[3] << 8; - break; + #endif - case 50: - block3[0] = block3[0] | append[0] << 16; - block3[1] = append[0] >> 16 | append[1] << 16; - block3[2] = append[1] >> 16 | append[2] << 16; - block3[3] = append[2] >> 16 | append[3] << 16; - break; + #ifdef IS_NV - case 51: - block3[0] = block3[0] | append[0] << 24; - block3[1] = append[0] >> 8 | append[1] << 24; - block3[2] = append[1] >> 8 | append[2] << 24; - block3[3] = append[2] >> 8 | append[3] << 24; - break; + const int offset_minus_4 = 4 - (block_len & 3); - case 52: - block3[1] = append[0]; - block3[2] = append[1]; - block3[3] = append[2]; - break; + const int selector = (0x76543210 >> (offset_minus_4 * 4)) & 0xffff; - case 53: - block3[1] = block3[1] | append[0] << 8; - block3[2] = append[0] >> 24 | append[1] << 8; - block3[3] = append[1] >> 24 | append[2] << 8; - break; + tmp0 = __byte_perm ( 0, append[0], selector); + tmp1 = __byte_perm (append[0], append[1], selector); + tmp2 = __byte_perm (append[1], append[2], selector); + tmp3 = __byte_perm (append[2], append[3], selector); + tmp4 = __byte_perm (append[3], 0x80, selector); - case 54: - block3[1] = block3[1] | append[0] << 16; - block3[2] = append[0] >> 16 | append[1] << 16; - block3[3] = append[1] >> 16 | append[2] << 16; - break; + #endif - case 55: - block3[1] = block3[1] | append[0] << 24; - block3[2] = append[0] >> 8 | append[1] << 24; - block3[3] = append[1] >> 8 | append[2] << 24; - break; + const u32 div = block_len / 4; - case 56: - block3[2] = append[0]; - block3[3] = append[1]; - break; + switch (div) + { + case 0: block0[0] |= tmp0; + block0[1] = tmp1; + block0[2] = tmp2; + block0[3] = tmp3; + block1[0] = tmp4; + break; + case 1: block0[1] |= tmp0; + block0[2] = tmp1; + block0[3] = tmp2; + block1[0] = tmp3; + block1[1] = tmp4; + break; + case 2: block0[2] |= tmp0; + block0[3] = tmp1; + block1[0] = tmp2; + block1[1] = tmp3; + block1[2] = tmp4; + break; + case 3: block0[3] |= tmp0; + block1[0] = tmp1; + block1[1] = tmp2; + block1[2] = tmp3; + block1[3] = tmp4; + break; + case 4: block1[0] |= tmp0; + block1[1] = tmp1; + block1[2] = tmp2; + block1[3] = tmp3; + block2[0] = tmp4; + break; + case 5: block1[1] |= tmp0; + block1[2] = tmp1; + block1[3] = tmp2; + block2[0] = tmp3; + block2[1] = tmp4; + break; + case 6: block1[2] |= tmp0; + block1[3] = tmp1; + block2[0] = tmp2; + block2[1] = tmp3; + block2[2] = tmp4; + break; + case 7: block1[3] |= tmp0; + block2[0] = tmp1; + block2[1] = tmp2; + block2[2] = tmp3; + block2[3] = tmp4; + break; + case 8: block2[0] |= tmp0; + block2[1] = tmp1; + block2[2] = tmp2; + block2[3] = tmp3; + block3[0] = tmp4; + break; + case 9: block2[1] |= tmp0; + block2[2] = tmp1; + block2[3] = tmp2; + block3[0] = tmp3; + block3[1] = tmp4; + break; } } -static void memcat8 (u32 block0[4], u32 block1[4], u32 block2[4], u32 block3[4], const u32 block_len, const u32 append[2]) +void memcat8 (u32 block0[4], u32 block1[4], u32 block2[4], u32 block3[4], const u32 block_len, const u32 append[2]) { - switch (block_len) - { - case 0: - block0[0] = append[0]; - block0[1] = append[1]; - break; - - case 1: - block0[0] = block0[0] | append[0] << 8; - block0[1] = append[0] >> 24 | append[1] << 8; - block0[2] = append[1] >> 24; - break; - - case 2: - block0[0] = block0[0] | append[0] << 16; - block0[1] = append[0] >> 16 | append[1] << 16; - block0[2] = append[1] >> 16; - break; - - case 3: - block0[0] = block0[0] | append[0] << 24; - block0[1] = append[0] >> 8 | append[1] << 24; - block0[2] = append[1] >> 8; - break; - - case 4: - block0[1] = append[0]; - block0[2] = append[1]; - break; - - case 5: - block0[1] = block0[1] | append[0] << 8; - block0[2] = append[0] >> 24 | append[1] << 8; - block0[3] = append[1] >> 24; - break; - - case 6: - block0[1] = block0[1] | append[0] << 16; - block0[2] = append[0] >> 16 | append[1] << 16; - block0[3] = append[1] >> 16; - break; - - case 7: - block0[1] = block0[1] | append[0] << 24; - block0[2] = append[0] >> 8 | append[1] << 24; - block0[3] = append[1] >> 8; - break; - - case 8: - block0[2] = append[0]; - block0[3] = append[1]; - break; - - case 9: - block0[2] = block0[2] | append[0] << 8; - block0[3] = append[0] >> 24 | append[1] << 8; - block1[0] = append[1] >> 24; - break; - - case 10: - block0[2] = block0[2] | append[0] << 16; - block0[3] = append[0] >> 16 | append[1] << 16; - block1[0] = append[1] >> 16; - break; - - case 11: - block0[2] = block0[2] | append[0] << 24; - block0[3] = append[0] >> 8 | append[1] << 24; - block1[0] = append[1] >> 8; - break; - - case 12: - block0[3] = append[0]; - block1[0] = append[1]; - break; - - case 13: - block0[3] = block0[3] | append[0] << 8; - block1[0] = append[0] >> 24 | append[1] << 8; - block1[1] = append[1] >> 24; - break; - - case 14: - block0[3] = block0[3] | append[0] << 16; - block1[0] = append[0] >> 16 | append[1] << 16; - block1[1] = append[1] >> 16; - break; - - case 15: - block0[3] = block0[3] | append[0] << 24; - block1[0] = append[0] >> 8 | append[1] << 24; - block1[1] = append[1] >> 8; - break; - - case 16: - block1[0] = append[0]; - block1[1] = append[1]; - break; - - case 17: - block1[0] = block1[0] | append[0] << 8; - block1[1] = append[0] >> 24 | append[1] << 8; - block1[2] = append[1] >> 24; - break; - - case 18: - block1[0] = block1[0] | append[0] << 16; - block1[1] = append[0] >> 16 | append[1] << 16; - block1[2] = append[1] >> 16; - break; - - case 19: - block1[0] = block1[0] | append[0] << 24; - block1[1] = append[0] >> 8 | append[1] << 24; - block1[2] = append[1] >> 8; - break; - - case 20: - block1[1] = append[0]; - block1[2] = append[1]; - break; - - case 21: - block1[1] = block1[1] | append[0] << 8; - block1[2] = append[0] >> 24 | append[1] << 8; - block1[3] = append[1] >> 24; - break; - - case 22: - block1[1] = block1[1] | append[0] << 16; - block1[2] = append[0] >> 16 | append[1] << 16; - block1[3] = append[1] >> 16; - break; - - case 23: - block1[1] = block1[1] | append[0] << 24; - block1[2] = append[0] >> 8 | append[1] << 24; - block1[3] = append[1] >> 8; - break; - - case 24: - block1[2] = append[0]; - block1[3] = append[1]; - break; - - case 25: - block1[2] = block1[2] | append[0] << 8; - block1[3] = append[0] >> 24 | append[1] << 8; - block2[0] = append[1] >> 24; - break; - - case 26: - block1[2] = block1[2] | append[0] << 16; - block1[3] = append[0] >> 16 | append[1] << 16; - block2[0] = append[1] >> 16; - break; - - case 27: - block1[2] = block1[2] | append[0] << 24; - block1[3] = append[0] >> 8 | append[1] << 24; - block2[0] = append[1] >> 8; - break; - - case 28: - block1[3] = append[0]; - block2[0] = append[1]; - break; - - case 29: - block1[3] = block1[3] | append[0] << 8; - block2[0] = append[0] >> 24 | append[1] << 8; - block2[1] = append[1] >> 24; - break; - - case 30: - block1[3] = block1[3] | append[0] << 16; - block2[0] = append[0] >> 16 | append[1] << 16; - block2[1] = append[1] >> 16; - break; - - case 31: - block1[3] = block1[3] | append[0] << 24; - block2[0] = append[0] >> 8 | append[1] << 24; - block2[1] = append[1] >> 8; - break; - - case 32: - block2[0] = append[0]; - block2[1] = append[1]; - break; - - case 33: - block2[0] = block2[0] | append[0] << 8; - block2[1] = append[0] >> 24 | append[1] << 8; - block2[2] = append[1] >> 24; - break; - - case 34: - block2[0] = block2[0] | append[0] << 16; - block2[1] = append[0] >> 16 | append[1] << 16; - block2[2] = append[1] >> 16; - break; - - case 35: - block2[0] = block2[0] | append[0] << 24; - block2[1] = append[0] >> 8 | append[1] << 24; - block2[2] = append[1] >> 8; - break; - - case 36: - block2[1] = append[0]; - block2[2] = append[1]; - break; - - case 37: - block2[1] = block2[1] | append[0] << 8; - block2[2] = append[0] >> 24 | append[1] << 8; - block2[3] = append[1] >> 24; - break; - - case 38: - block2[1] = block2[1] | append[0] << 16; - block2[2] = append[0] >> 16 | append[1] << 16; - block2[3] = append[1] >> 16; - break; - - case 39: - block2[1] = block2[1] | append[0] << 24; - block2[2] = append[0] >> 8 | append[1] << 24; - block2[3] = append[1] >> 8; - break; - - case 40: - block2[2] = append[0]; - block2[3] = append[1]; - break; - - case 41: - block2[2] = block2[2] | append[0] << 8; - block2[3] = append[0] >> 24 | append[1] << 8; - block3[0] = append[1] >> 24; - break; - - case 42: - block2[2] = block2[2] | append[0] << 16; - block2[3] = append[0] >> 16 | append[1] << 16; - block3[0] = append[1] >> 16; - break; - - case 43: - block2[2] = block2[2] | append[0] << 24; - block2[3] = append[0] >> 8 | append[1] << 24; - block3[0] = append[1] >> 8; - break; + u32 tmp0; + u32 tmp1; + u32 tmp2; - case 44: - block2[3] = append[0]; - block3[0] = append[1]; - break; + #if defined IS_AMD || defined IS_GENERIC - case 45: - block2[3] = block2[3] | append[0] << 8; - block3[0] = append[0] >> 24 | append[1] << 8; - block3[1] = append[1] >> 24; - break; + const int offset_minus_4 = 4 - (block_len & 3); - case 46: - block2[3] = block2[3] | append[0] << 16; - block3[0] = append[0] >> 16 | append[1] << 16; - block3[1] = append[1] >> 16; - break; + tmp0 = amd_bytealign (append[0], 0, offset_minus_4); + tmp1 = amd_bytealign (append[1], append[0], offset_minus_4); + tmp2 = amd_bytealign ( 0, append[1], offset_minus_4); - case 47: - block2[3] = block2[3] | append[0] << 24; - block3[0] = append[0] >> 8 | append[1] << 24; - block3[1] = append[1] >> 8; - break; + const u32 mod = block_len & 3; - case 48: - block3[0] = append[0]; - block3[1] = append[1]; - break; + if (mod == 0) + { + tmp0 = tmp1; + tmp1 = tmp2; + tmp2 = 0; + } - case 49: - block3[0] = block3[0] | append[0] << 8; - block3[1] = append[0] >> 24 | append[1] << 8; - block3[2] = append[1] >> 24; - break; + #endif - case 50: - block3[0] = block3[0] | append[0] << 16; - block3[1] = append[0] >> 16 | append[1] << 16; - block3[2] = append[1] >> 16; - break; + #ifdef IS_NV - case 51: - block3[0] = block3[0] | append[0] << 24; - block3[1] = append[0] >> 8 | append[1] << 24; - block3[2] = append[1] >> 8; - break; + const int offset_minus_4 = 4 - (block_len & 3); - case 52: - block3[1] = append[0]; - block3[2] = append[1]; - break; + const int selector = (0x76543210 >> (offset_minus_4 * 4)) & 0xffff; - case 53: - block3[1] = block3[1] | append[0] << 8; - block3[2] = append[0] >> 24 | append[1] << 8; - block3[3] = append[1] >> 24; - break; + tmp0 = __byte_perm ( 0, append[0], selector); + tmp1 = __byte_perm (append[0], append[1], selector); + tmp2 = __byte_perm (append[1], 0, selector); - case 54: - block3[1] = block3[1] | append[0] << 16; - block3[2] = append[0] >> 16 | append[1] << 16; - block3[3] = append[1] >> 16; - break; + #endif - case 55: - block3[1] = block3[1] | append[0] << 24; - block3[2] = append[0] >> 8 | append[1] << 24; - block3[3] = append[1] >> 8; - break; + const u32 div = block_len / 4; - case 56: - block3[2] = append[0]; - block3[3] = append[1]; - break; + switch (div) + { + case 0: block0[0] |= tmp0; + block0[1] = tmp1; + block0[2] = tmp2; + break; + case 1: block0[1] |= tmp0; + block0[2] = tmp1; + block0[3] = tmp2; + break; + case 2: block0[2] |= tmp0; + block0[3] = tmp1; + block1[0] = tmp2; + break; + case 3: block0[3] |= tmp0; + block1[0] = tmp1; + block1[1] = tmp2; + break; + case 4: block1[0] |= tmp0; + block1[1] = tmp1; + block1[2] = tmp2; + break; + case 5: block1[1] |= tmp0; + block1[2] = tmp1; + block1[3] = tmp2; + break; + case 6: block1[2] |= tmp0; + block1[3] = tmp1; + block2[0] = tmp2; + break; + case 7: block1[3] |= tmp0; + block2[0] = tmp1; + block2[1] = tmp2; + break; + case 8: block2[0] |= tmp0; + block2[1] = tmp1; + block2[2] = tmp2; + break; + case 9: block2[1] |= tmp0; + block2[2] = tmp1; + block2[3] = tmp2; + break; + case 10: block2[2] |= tmp0; + block2[3] = tmp1; + block3[0] = tmp2; + break; + case 11: block2[3] |= tmp0; + block3[0] = tmp1; + block3[1] = tmp2; + break; } } -static void append_1st (u32 block0[4], u32 block1[4], u32 block2[4], u32 block3[4], const u32 block_len, const u32 append) +void append_1st (u32 block0[4], u32 block1[4], u32 block2[4], u32 block3[4], const u32 block_len, const u32 append) { switch (block_len) { @@ -1574,7 +662,7 @@ static void append_1st (u32 block0[4], u32 block1[4], u32 block2[4], u32 block3[ } } -__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06300_init (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global md5crypt_tmp_t *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) +__kernel void m06300_init (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global md5crypt_tmp_t *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *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 u32 gid_max) { /** * base @@ -1732,7 +820,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06300_init (__gl tmps[gid].digest_buf[3] = digest[3]; } -__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06300_loop (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global md5crypt_tmp_t *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) +__kernel void m06300_loop (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global md5crypt_tmp_t *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *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 u32 gid_max) { /** * base @@ -1920,7 +1008,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06300_loop (__gl tmps[gid].digest_buf[3] = digest[3]; } -__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06300_comp (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global md5crypt_tmp_t *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) +__kernel void m06300_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global md5crypt_tmp_t *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *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 u32 gid_max) { /** * modifier