diff --git a/OpenCL/m24000-pure.cl b/OpenCL/m24000-pure.cl new file mode 100644 index 000000000..37e219e2c --- /dev/null +++ b/OpenCL/m24000-pure.cl @@ -0,0 +1,917 @@ +/** + * Author......: See docs/credits.txt + * License.....: MIT + */ + +#ifdef KERNEL_STATIC +#include M2S(INCLUDE_PATH/inc_vendor.h) +#include M2S(INCLUDE_PATH/inc_types.h) +#include M2S(INCLUDE_PATH/inc_platform.cl) +#include M2S(INCLUDE_PATH/inc_common.cl) +#include M2S(INCLUDE_PATH/inc_hash_sha256.cl) +#include M2S(INCLUDE_PATH/inc_cipher_aes.cl) +#include M2S(INCLUDE_PATH/inc_cipher_twofish.cl) +#include M2S(INCLUDE_PATH/inc_cipher_serpent.cl) +#include M2S(INCLUDE_PATH/inc_cipher_camellia.cl) +#endif + +typedef struct +{ + #ifndef SCRYPT_TMP_ELEM + #define SCRYPT_TMP_ELEM 1 + #endif + + uint4 P[SCRYPT_TMP_ELEM]; + +} scrypt_tmp_t; + +typedef struct bestcrypt_scrypt +{ + u32 salt_buf[24]; + u32 ciphertext[96]; + u32 version; + +} bestcrypt_scrypt_t; + +#if defined IS_CUDA || defined IS_HIP + +inline __device__ uint4 operator & (const uint4 a, const u32 b) { return make_uint4 ((a.x & b ), (a.y & b ), (a.z & b ), (a.w & b )); } +inline __device__ uint4 operator << (const uint4 a, const u32 b) { return make_uint4 ((a.x << b ), (a.y << b ), (a.z << b ), (a.w << b )); } +inline __device__ uint4 operator >> (const uint4 a, const u32 b) { return make_uint4 ((a.x >> b ), (a.y >> b ), (a.z >> b ), (a.w >> b )); } +inline __device__ uint4 operator + (const uint4 a, const uint4 b) { return make_uint4 ((a.x + b.x), (a.y + b.y), (a.z + b.z), (a.w + b.w)); } +inline __device__ uint4 operator ^ (const uint4 a, const uint4 b) { return make_uint4 ((a.x ^ b.x), (a.y ^ b.y), (a.z ^ b.z), (a.w ^ b.w)); } +inline __device__ uint4 operator | (const uint4 a, const uint4 b) { return make_uint4 ((a.x | b.x), (a.y | b.y), (a.z | b.z), (a.w | b.w)); } +inline __device__ void operator ^= ( uint4 &a, const uint4 b) { a.x ^= b.x; a.y ^= b.y; a.z ^= b.z; a.w ^= b.w; } + +inline __device__ uint4 rotate (const uint4 a, const int n) +{ + return ((a << n) | ((a >> (32 - n)))); +} + +#endif + +DECLSPEC uint4 hc_swap32_4 (uint4 v) +{ + return (rotate ((v & 0x00FF00FF), 24u) | rotate ((v & 0xFF00FF00), 8u)); +} + +#define GET_SCRYPT_CNT(r,p) (2 * (r) * 16 * (p)) +#define GET_SMIX_CNT(r,N) (2 * (r) * 16 * (N)) +#define GET_STATE_CNT(r) (2 * (r) * 16) + +#define SCRYPT_CNT GET_SCRYPT_CNT (SCRYPT_R, SCRYPT_P) +#define SCRYPT_CNT4 (SCRYPT_CNT / 4) +#define STATE_CNT GET_STATE_CNT (SCRYPT_R) +#define STATE_CNT4 (STATE_CNT / 4) + +#define ADD_ROTATE_XOR(r,i1,i2,s) (r) ^= rotate ((i1) + (i2), (s)); + +#if defined IS_CUDA || defined IS_HIP + +#define SALSA20_2R() \ +{ \ + ADD_ROTATE_XOR (X1, X0, X3, 7); \ + ADD_ROTATE_XOR (X2, X1, X0, 9); \ + ADD_ROTATE_XOR (X3, X2, X1, 13); \ + ADD_ROTATE_XOR (X0, X3, X2, 18); \ + \ + X1 = make_uint4 (X1.w, X1.x, X1.y, X1.z); \ + X2 = make_uint4 (X2.z, X2.w, X2.x, X2.y); \ + X3 = make_uint4 (X3.y, X3.z, X3.w, X3.x); \ + \ + ADD_ROTATE_XOR (X3, X0, X1, 7); \ + ADD_ROTATE_XOR (X2, X3, X0, 9); \ + ADD_ROTATE_XOR (X1, X2, X3, 13); \ + ADD_ROTATE_XOR (X0, X1, X2, 18); \ + \ + X1 = make_uint4 (X1.y, X1.z, X1.w, X1.x); \ + X2 = make_uint4 (X2.z, X2.w, X2.x, X2.y); \ + X3 = make_uint4 (X3.w, X3.x, X3.y, X3.z); \ +} +#elif defined IS_METAL +#define SALSA20_2R() \ +{ \ + ADD_ROTATE_XOR (X1, X0, X3, 7); \ + ADD_ROTATE_XOR (X2, X1, X0, 9); \ + ADD_ROTATE_XOR (X3, X2, X1, 13); \ + ADD_ROTATE_XOR (X0, X3, X2, 18); \ + \ + X1 = X1.wxyz; \ + X2 = X2.zwxy; \ + X3 = X3.yzwx; \ + \ + ADD_ROTATE_XOR (X3, X0, X1, 7); \ + ADD_ROTATE_XOR (X2, X3, X0, 9); \ + ADD_ROTATE_XOR (X1, X2, X3, 13); \ + ADD_ROTATE_XOR (X0, X1, X2, 18); \ + \ + X1 = X1.yzwx; \ + X2 = X2.zwxy; \ + X3 = X3.wxyz; \ +} +#else +#define SALSA20_2R() \ +{ \ + ADD_ROTATE_XOR (X1, X0, X3, 7); \ + ADD_ROTATE_XOR (X2, X1, X0, 9); \ + ADD_ROTATE_XOR (X3, X2, X1, 13); \ + ADD_ROTATE_XOR (X0, X3, X2, 18); \ + \ + X1 = X1.s3012; \ + X2 = X2.s2301; \ + X3 = X3.s1230; \ + \ + ADD_ROTATE_XOR (X3, X0, X1, 7); \ + ADD_ROTATE_XOR (X2, X3, X0, 9); \ + ADD_ROTATE_XOR (X1, X2, X3, 13); \ + ADD_ROTATE_XOR (X0, X1, X2, 18); \ + \ + X1 = X1.s1230; \ + X2 = X2.s2301; \ + X3 = X3.s3012; \ +} +#endif + +#define SALSA20_8_XOR() \ +{ \ + R0 = R0 ^ Y0; \ + R1 = R1 ^ Y1; \ + R2 = R2 ^ Y2; \ + R3 = R3 ^ Y3; \ + \ + uint4 X0 = R0; \ + uint4 X1 = R1; \ + uint4 X2 = R2; \ + uint4 X3 = R3; \ + \ + SALSA20_2R (); \ + SALSA20_2R (); \ + SALSA20_2R (); \ + SALSA20_2R (); \ + \ + R0 = R0 + X0; \ + R1 = R1 + X1; \ + R2 = R2 + X2; \ + R3 = R3 + X3; \ +} + +DECLSPEC void salsa_r (PRIVATE_AS uint4 *TI) +{ + uint4 R0 = TI[STATE_CNT4 - 4]; + uint4 R1 = TI[STATE_CNT4 - 3]; + uint4 R2 = TI[STATE_CNT4 - 2]; + uint4 R3 = TI[STATE_CNT4 - 1]; + + uint4 TO[STATE_CNT4]; + + int idx_y = 0; + int idx_r1 = 0; + int idx_r2 = SCRYPT_R * 4; + + for (int i = 0; i < SCRYPT_R; i++) + { + uint4 Y0; + uint4 Y1; + uint4 Y2; + uint4 Y3; + + Y0 = TI[idx_y++]; + Y1 = TI[idx_y++]; + Y2 = TI[idx_y++]; + Y3 = TI[idx_y++]; + + SALSA20_8_XOR (); + + TO[idx_r1++] = R0; + TO[idx_r1++] = R1; + TO[idx_r1++] = R2; + TO[idx_r1++] = R3; + + Y0 = TI[idx_y++]; + Y1 = TI[idx_y++]; + Y2 = TI[idx_y++]; + Y3 = TI[idx_y++]; + + SALSA20_8_XOR (); + + TO[idx_r2++] = R0; + TO[idx_r2++] = R1; + TO[idx_r2++] = R2; + TO[idx_r2++] = R3; + } + + #pragma unroll + for (int i = 0; i < STATE_CNT4; i++) + { + TI[i] = TO[i]; + } +} + +DECLSPEC void scrypt_smix (PRIVATE_AS uint4 *X, PRIVATE_AS uint4 *T, GLOBAL_AS uint4 *V0, GLOBAL_AS uint4 *V1, GLOBAL_AS uint4 *V2, GLOBAL_AS uint4 *V3, const u64 gid) +{ + #define Coord(xd4,y,z) (((xd4) * ySIZE * zSIZE) + ((y) * zSIZE) + (z)) + #define CO Coord(xd4,y,z) + + const u32 ySIZE = SCRYPT_N / SCRYPT_TMTO; + const u32 zSIZE = STATE_CNT4; + + const u32 x = (u32) gid; + + const u32 xd4 = x / 4; + const u32 xm4 = x & 3; + + GLOBAL_AS uint4 *V; + + switch (xm4) + { + case 0: V = V0; break; + case 1: V = V1; break; + case 2: V = V2; break; + case 3: V = V3; break; + } + + #ifdef _unroll + #pragma unroll + #endif + for (u32 i = 0; i < STATE_CNT4; i += 4) + { + #if defined IS_CUDA || defined IS_HIP + T[0] = make_uint4 (X[i + 0].x, X[i + 1].y, X[i + 2].z, X[i + 3].w); + T[1] = make_uint4 (X[i + 1].x, X[i + 2].y, X[i + 3].z, X[i + 0].w); + T[2] = make_uint4 (X[i + 2].x, X[i + 3].y, X[i + 0].z, X[i + 1].w); + T[3] = make_uint4 (X[i + 3].x, X[i + 0].y, X[i + 1].z, X[i + 2].w); + #elif defined IS_METAL + T[0] = uint4 (X[i + 0].x, X[i + 1].y, X[i + 2].z, X[i + 3].w); + T[1] = uint4 (X[i + 1].x, X[i + 2].y, X[i + 3].z, X[i + 0].w); + T[2] = uint4 (X[i + 2].x, X[i + 3].y, X[i + 0].z, X[i + 1].w); + T[3] = uint4 (X[i + 3].x, X[i + 0].y, X[i + 1].z, X[i + 2].w); + #else + T[0] = (uint4) (X[i + 0].x, X[i + 1].y, X[i + 2].z, X[i + 3].w); + T[1] = (uint4) (X[i + 1].x, X[i + 2].y, X[i + 3].z, X[i + 0].w); + T[2] = (uint4) (X[i + 2].x, X[i + 3].y, X[i + 0].z, X[i + 1].w); + T[3] = (uint4) (X[i + 3].x, X[i + 0].y, X[i + 1].z, X[i + 2].w); + #endif + + X[i + 0] = T[0]; + X[i + 1] = T[1]; + X[i + 2] = T[2]; + X[i + 3] = T[3]; + } + + for (u32 y = 0; y < ySIZE; y++) + { + for (u32 z = 0; z < zSIZE; z++) V[CO] = X[z]; + + for (u32 i = 0; i < SCRYPT_TMTO; i++) salsa_r (X); + } + + for (u32 i = 0; i < SCRYPT_N; i++) + { + const u32 k = X[zSIZE - 4].x & (SCRYPT_N - 1); + + const u32 y = k / SCRYPT_TMTO; + + const u32 km = k - (y * SCRYPT_TMTO); + + for (u32 z = 0; z < zSIZE; z++) T[z] = V[CO]; + + for (u32 i = 0; i < km; i++) salsa_r (T); + + for (u32 z = 0; z < zSIZE; z++) X[z] ^= T[z]; + + salsa_r (X); + } + + #ifdef _unroll + #pragma unroll + #endif + for (u32 i = 0; i < STATE_CNT4; i += 4) + { + #if defined IS_CUDA || defined IS_HIP + T[0] = make_uint4 (X[i + 0].x, X[i + 3].y, X[i + 2].z, X[i + 1].w); + T[1] = make_uint4 (X[i + 1].x, X[i + 0].y, X[i + 3].z, X[i + 2].w); + T[2] = make_uint4 (X[i + 2].x, X[i + 1].y, X[i + 0].z, X[i + 3].w); + T[3] = make_uint4 (X[i + 3].x, X[i + 2].y, X[i + 1].z, X[i + 0].w); + #elif defined IS_METAL + T[0] = uint4 (X[i + 0].x, X[i + 3].y, X[i + 2].z, X[i + 1].w); + T[1] = uint4 (X[i + 1].x, X[i + 0].y, X[i + 3].z, X[i + 2].w); + T[2] = uint4 (X[i + 2].x, X[i + 1].y, X[i + 0].z, X[i + 3].w); + T[3] = uint4 (X[i + 3].x, X[i + 2].y, X[i + 1].z, X[i + 0].w); + #else + T[0] = (uint4) (X[i + 0].x, X[i + 3].y, X[i + 2].z, X[i + 1].w); + T[1] = (uint4) (X[i + 1].x, X[i + 0].y, X[i + 3].z, X[i + 2].w); + T[2] = (uint4) (X[i + 2].x, X[i + 1].y, X[i + 0].z, X[i + 3].w); + T[3] = (uint4) (X[i + 3].x, X[i + 2].y, X[i + 1].z, X[i + 0].w); + #endif + + X[i + 0] = T[0]; + X[i + 1] = T[1]; + X[i + 2] = T[2]; + X[i + 3] = T[3]; + } +} + +#ifndef KECCAK_ROUNDS +#define KECCAK_ROUNDS 24 +#endif + +#define Theta1(s) (st[0 + s] ^ st[5 + s] ^ st[10 + s] ^ st[15 + s] ^ st[20 + s]) + +#define Theta2(s) \ +{ \ + st[ 0 + s] ^= t; \ + st[ 5 + s] ^= t; \ + st[10 + s] ^= t; \ + st[15 + s] ^= t; \ + st[20 + s] ^= t; \ +} + +#define Rho_Pi(s) \ +{ \ + u32 j = keccakf_piln[s]; \ + u32 k = keccakf_rotc[s]; \ + bc0 = st[j]; \ + st[j] = hc_rotl64_S (t, k); \ + t = bc0; \ +} + +#define Chi(s) \ +{ \ + bc0 = st[0 + s]; \ + bc1 = st[1 + s]; \ + bc2 = st[2 + s]; \ + bc3 = st[3 + s]; \ + bc4 = st[4 + s]; \ + st[0 + s] ^= ~bc1 & bc2; \ + st[1 + s] ^= ~bc2 & bc3; \ + st[2 + s] ^= ~bc3 & bc4; \ + st[3 + s] ^= ~bc4 & bc0; \ + st[4 + s] ^= ~bc0 & bc1; \ +} + +CONSTANT_VK u64a keccakf_rndc[24] = +{ + 0x0000000000000001, 0x0000000000008082, 0x800000000000808a, + 0x8000000080008000, 0x000000000000808b, 0x0000000080000001, + 0x8000000080008081, 0x8000000000008009, 0x000000000000008a, + 0x0000000000000088, 0x0000000080008009, 0x000000008000000a, + 0x000000008000808b, 0x800000000000008b, 0x8000000000008089, + 0x8000000000008003, 0x8000000000008002, 0x8000000000000080, + 0x000000000000800a, 0x800000008000000a, 0x8000000080008081, + 0x8000000000008080, 0x0000000080000001, 0x8000000080008008 +}; + +DECLSPEC void keccak_transform_S (PRIVATE_AS u64 *st) +{ + const u8 keccakf_rotc[24] = + { + 1, 3, 6, 10, 15, 21, 28, 36, 45, 55, 2, 14, + 27, 41, 56, 8, 25, 43, 62, 18, 39, 61, 20, 44 + }; + + const u8 keccakf_piln[24] = + { + 10, 7, 11, 17, 18, 3, 5, 16, 8, 21, 24, 4, + 15, 23, 19, 13, 12, 2, 20, 14, 22, 9, 6, 1 + }; + + /** + * Keccak + */ + + int round; + + for (round = 0; round < KECCAK_ROUNDS; round++) + { + // Theta + + u64 bc0 = Theta1 (0); + u64 bc1 = Theta1 (1); + u64 bc2 = Theta1 (2); + u64 bc3 = Theta1 (3); + u64 bc4 = Theta1 (4); + + u64 t; + + t = bc4 ^ hc_rotl64_S (bc1, 1); Theta2 (0); + t = bc0 ^ hc_rotl64_S (bc2, 1); Theta2 (1); + t = bc1 ^ hc_rotl64_S (bc3, 1); Theta2 (2); + t = bc2 ^ hc_rotl64_S (bc4, 1); Theta2 (3); + t = bc3 ^ hc_rotl64_S (bc0, 1); Theta2 (4); + + // Rho Pi + + t = st[1]; + + Rho_Pi (0); + Rho_Pi (1); + Rho_Pi (2); + Rho_Pi (3); + Rho_Pi (4); + Rho_Pi (5); + Rho_Pi (6); + Rho_Pi (7); + Rho_Pi (8); + Rho_Pi (9); + Rho_Pi (10); + Rho_Pi (11); + Rho_Pi (12); + Rho_Pi (13); + Rho_Pi (14); + Rho_Pi (15); + Rho_Pi (16); + Rho_Pi (17); + Rho_Pi (18); + Rho_Pi (19); + Rho_Pi (20); + Rho_Pi (21); + Rho_Pi (22); + Rho_Pi (23); + + // Chi + + Chi (0); + Chi (5); + Chi (10); + Chi (15); + Chi (20); + + // Iota + + st[0] ^= keccakf_rndc[round]; + } +} + +KERNEL_FQ void m24000_init (KERN_ATTR_TMPS_ESALT (scrypt_tmp_t, bestcrypt_scrypt_t)) +{ + /** + * base + */ + + const u64 gid = get_global_id (0); + + if (gid >= GID_CNT) return; + + sha256_hmac_ctx_t sha256_hmac_ctx; + + sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len); + + sha256_hmac_update_global_swap (&sha256_hmac_ctx, salt_bufs[SALT_POS_HOST].salt_buf, salt_bufs[SALT_POS_HOST].salt_len); + + for (u32 i = 0, j = 1, k = 0; i < SCRYPT_CNT; i += 8, j += 1, k += 2) + { + sha256_hmac_ctx_t sha256_hmac_ctx2 = sha256_hmac_ctx; + + u32 w0[4]; + u32 w1[4]; + u32 w2[4]; + u32 w3[4]; + + w0[0] = j; + w0[1] = 0; + w0[2] = 0; + w0[3] = 0; + w1[0] = 0; + w1[1] = 0; + w1[2] = 0; + w1[3] = 0; + w2[0] = 0; + w2[1] = 0; + w2[2] = 0; + w2[3] = 0; + w3[0] = 0; + w3[1] = 0; + w3[2] = 0; + w3[3] = 0; + + sha256_hmac_update_64 (&sha256_hmac_ctx2, w0, w1, w2, w3, 4); + + sha256_hmac_final (&sha256_hmac_ctx2); + + u32 digest[8]; + + digest[0] = sha256_hmac_ctx2.opad.h[0]; + digest[1] = sha256_hmac_ctx2.opad.h[1]; + digest[2] = sha256_hmac_ctx2.opad.h[2]; + digest[3] = sha256_hmac_ctx2.opad.h[3]; + digest[4] = sha256_hmac_ctx2.opad.h[4]; + digest[5] = sha256_hmac_ctx2.opad.h[5]; + digest[6] = sha256_hmac_ctx2.opad.h[6]; + digest[7] = sha256_hmac_ctx2.opad.h[7]; + + #if defined IS_CUDA || defined IS_HIP + const uint4 tmp0 = make_uint4 (digest[0], digest[1], digest[2], digest[3]); + const uint4 tmp1 = make_uint4 (digest[4], digest[5], digest[6], digest[7]); + #elif defined IS_METAL + const uint4 tmp0 = uint4 (digest[0], digest[1], digest[2], digest[3]); + const uint4 tmp1 = uint4 (digest[4], digest[5], digest[6], digest[7]); + #else + const uint4 tmp0 = (uint4) (digest[0], digest[1], digest[2], digest[3]); + const uint4 tmp1 = (uint4) (digest[4], digest[5], digest[6], digest[7]); + #endif + + tmps[gid].P[k + 0] = tmp0; + tmps[gid].P[k + 1] = tmp1; + } +} + +KERNEL_FQ void m24000_loop (KERN_ATTR_TMPS_ESALT (scrypt_tmp_t, bestcrypt_scrypt_t)) +{ + const u64 gid = get_global_id (0); + + if (gid >= GID_CNT) return; + + GLOBAL_AS uint4 *d_scrypt0_buf = (GLOBAL_AS uint4 *) d_extra0_buf; + GLOBAL_AS uint4 *d_scrypt1_buf = (GLOBAL_AS uint4 *) d_extra1_buf; + GLOBAL_AS uint4 *d_scrypt2_buf = (GLOBAL_AS uint4 *) d_extra2_buf; + GLOBAL_AS uint4 *d_scrypt3_buf = (GLOBAL_AS uint4 *) d_extra3_buf; + + uint4 X[STATE_CNT4]; + uint4 T[STATE_CNT4]; + + #ifdef _unroll + #pragma unroll + #endif + for (int z = 0; z < STATE_CNT4; z++) X[z] = hc_swap32_4 (tmps[gid].P[z]); + + scrypt_smix (X, T, d_scrypt0_buf, d_scrypt1_buf, d_scrypt2_buf, d_scrypt3_buf, gid); + + #ifdef _unroll + #pragma unroll + #endif + for (int z = 0; z < STATE_CNT4; z++) tmps[gid].P[z] = hc_swap32_4 (X[z]); + + #if SCRYPT_P >= 1 + for (int i = STATE_CNT4; i < SCRYPT_CNT4; i += STATE_CNT4) + { + for (int z = 0; z < STATE_CNT4; z++) X[z] = hc_swap32_4 (tmps[gid].P[i + z]); + + scrypt_smix (X, T, d_scrypt0_buf, d_scrypt1_buf, d_scrypt2_buf, d_scrypt3_buf, gid); + + for (int z = 0; z < STATE_CNT4; z++) tmps[gid].P[i + z] = hc_swap32_4 (X[z]); + } + #endif +} + +KERNEL_FQ void m24000_comp (KERN_ATTR_TMPS_ESALT (scrypt_tmp_t, bestcrypt_scrypt_t)) +{ + /** + * base + */ + + const u64 gid = get_global_id (0); + const u64 lid = get_local_id (0); + const u64 lsz = get_local_size (0); + + /** + * aes shared + */ + + #ifdef REAL_SHM + + LOCAL_VK u32 s_td0[256]; + LOCAL_VK u32 s_td1[256]; + LOCAL_VK u32 s_td2[256]; + LOCAL_VK u32 s_td3[256]; + LOCAL_VK u32 s_td4[256]; + + LOCAL_VK u32 s_te0[256]; + LOCAL_VK u32 s_te1[256]; + LOCAL_VK u32 s_te2[256]; + LOCAL_VK u32 s_te3[256]; + LOCAL_VK u32 s_te4[256]; + + for (u32 i = lid; i < 256; i += lsz) + { + s_td0[i] = td0[i]; + s_td1[i] = td1[i]; + s_td2[i] = td2[i]; + s_td3[i] = td3[i]; + s_td4[i] = td4[i]; + + s_te0[i] = te0[i]; + s_te1[i] = te1[i]; + s_te2[i] = te2[i]; + s_te3[i] = te3[i]; + s_te4[i] = te4[i]; + } + + SYNC_THREADS (); + + #else + + CONSTANT_AS u32a *s_td0 = td0; + CONSTANT_AS u32a *s_td1 = td1; + CONSTANT_AS u32a *s_td2 = td2; + CONSTANT_AS u32a *s_td3 = td3; + CONSTANT_AS u32a *s_td4 = td4; + + CONSTANT_AS u32a *s_te0 = te0; + CONSTANT_AS u32a *s_te1 = te1; + CONSTANT_AS u32a *s_te2 = te2; + CONSTANT_AS u32a *s_te3 = te3; + CONSTANT_AS u32a *s_te4 = te4; + #endif + + /** + * AES part + */ + + if (gid >= GID_CNT) return; + + /** + * 2nd pbkdf2, creates B + */ + + u32 w0[4]; + u32 w1[4]; + u32 w2[4]; + u32 w3[4]; + + sha256_hmac_ctx_t ctx; + + sha256_hmac_init_global_swap (&ctx, pws[gid].i, pws[gid].pw_len); + + for (u32 l = 0; l < SCRYPT_CNT4; l += 4) + { + uint4 tmp; + + tmp = tmps[gid].P[l + 0]; + + w0[0] = tmp.x; + w0[1] = tmp.y; + w0[2] = tmp.z; + w0[3] = tmp.w; + + tmp = tmps[gid].P[l + 1]; + + w1[0] = tmp.x; + w1[1] = tmp.y; + w1[2] = tmp.z; + w1[3] = tmp.w; + + tmp = tmps[gid].P[l + 2]; + + w2[0] = tmp.x; + w2[1] = tmp.y; + w2[2] = tmp.z; + w2[3] = tmp.w; + + tmp = tmps[gid].P[l + 3]; + + w3[0] = tmp.x; + w3[1] = tmp.y; + w3[2] = tmp.z; + w3[3] = tmp.w; + + sha256_hmac_update_64 (&ctx, w0, w1, w2, w3, 64); + } + + w0[0] = 1; + w0[1] = 0; + w0[2] = 0; + w0[3] = 0; + w1[0] = 0; + w1[1] = 0; + w1[2] = 0; + w1[3] = 0; + w2[0] = 0; + w2[1] = 0; + w2[2] = 0; + w2[3] = 0; + w3[0] = 0; + w3[1] = 0; + w3[2] = 0; + w3[3] = 0; + + sha256_hmac_update_64 (&ctx, w0, w1, w2, w3, 4); + + sha256_hmac_final (&ctx); + + u32 version = esalt_bufs[DIGESTS_OFFSET_HOST].version; + + u32 iv[4] = { 0 }; + + u32 res[20]; // full would be 24 x u32 (96 bytes) + + if (version == 0x38) //0x38 is char for '8' which is the crypto type passed in position 3 of hash ( $08$ ) + { + + #define KEYLEN 60 + + u32 ks[KEYLEN]; + + AES256_set_decrypt_key (ks, ctx.opad.h, s_te0, s_te1, s_te2, s_te3, s_td0, s_td1, s_td2, s_td3); + + for (u32 i = 0; i < 20; i += 4) // 96 bytes output would contain the full 32 byte checksum + { + u32 data[4]; + + data[0] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[i + 0]; + data[1] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[i + 1]; + data[2] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[i + 2]; + data[3] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[i + 3]; + + u32 out[4]; + + aes256_decrypt (ks, data, out, s_td0, s_td1, s_td2, s_td3, s_td4); + + res[i + 0] = hc_swap32_S (out[0] ^ iv[0]); + res[i + 1] = hc_swap32_S (out[1] ^ iv[1]); + res[i + 2] = hc_swap32_S (out[2] ^ iv[2]); + res[i + 3] = hc_swap32_S (out[3] ^ iv[3]); + + iv[0] = data[0]; + iv[1] = data[1]; + iv[2] = data[2]; + iv[3] = data[3]; + } + } + + + if (version == 0x39) //0x39 is char for '9' which is the crypto type passed in position 3 of hash ( $09$ ) + { + u32 sk[4]; + u32 lk[40]; + + ctx.opad.h[0] = hc_swap32_S (ctx.opad.h[0]); + ctx.opad.h[1] = hc_swap32_S (ctx.opad.h[1]); + ctx.opad.h[2] = hc_swap32_S (ctx.opad.h[2]); + ctx.opad.h[3] = hc_swap32_S (ctx.opad.h[3]); + ctx.opad.h[4] = hc_swap32_S (ctx.opad.h[4]); + ctx.opad.h[5] = hc_swap32_S (ctx.opad.h[5]); + ctx.opad.h[6] = hc_swap32_S (ctx.opad.h[6]); + ctx.opad.h[7] = hc_swap32_S (ctx.opad.h[7]); + + twofish256_set_key (sk, lk, ctx.opad.h); + + for (u32 i = 0; i < 20; i += 4) // 96 bytes output would contain the full 32 byte checksum + { + u32 data[4]; + + data[0] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[i + 0]; + data[1] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[i + 1]; + data[2] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[i + 2]; + data[3] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[i + 3]; + + + u32 out[4]; + + twofish256_decrypt (sk, lk, data, out); + + res[i + 0] = hc_swap32_S (out[0] ^ iv[0]); + res[i + 1] = hc_swap32_S (out[1] ^ iv[1]); + res[i + 2] = hc_swap32_S (out[2] ^ iv[2]); + res[i + 3] = hc_swap32_S (out[3] ^ iv[3]); + + iv[0] = data[0]; + iv[1] = data[1]; + iv[2] = data[2]; + iv[3] = data[3]; + } + } + + if (version == 0x61) //0x61 is char for 'a' which is the crypto type passed in position 3 of hash ( $0a$ ) + { + u32 ks_serpent[140]; + + ctx.opad.h[0] = hc_swap32_S (ctx.opad.h[0]); + ctx.opad.h[1] = hc_swap32_S (ctx.opad.h[1]); + ctx.opad.h[2] = hc_swap32_S (ctx.opad.h[2]); + ctx.opad.h[3] = hc_swap32_S (ctx.opad.h[3]); + ctx.opad.h[4] = hc_swap32_S (ctx.opad.h[4]); + ctx.opad.h[5] = hc_swap32_S (ctx.opad.h[5]); + ctx.opad.h[6] = hc_swap32_S (ctx.opad.h[6]); + ctx.opad.h[7] = hc_swap32_S (ctx.opad.h[7]); + + serpent256_set_key (ks_serpent, ctx.opad.h); + + for (u32 i = 0; i < 20; i += 4) // 96 bytes output would contain the full 32 byte checksum + { + u32 data[4]; + + data[0] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[i + 0]; + data[1] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[i + 1]; + data[2] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[i + 2]; + data[3] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[i + 3]; + + + u32 out[4]; + + serpent256_decrypt (ks_serpent, data, out); + + res[i + 0] = hc_swap32_S (out[0] ^ iv[0]); + res[i + 1] = hc_swap32_S (out[1] ^ iv[1]); + res[i + 2] = hc_swap32_S (out[2] ^ iv[2]); + res[i + 3] = hc_swap32_S (out[3] ^ iv[3]); + + iv[0] = data[0]; + iv[1] = data[1]; + iv[2] = data[2]; + iv[3] = data[3]; + } + } + + if (version == 0x66) //0x66 is char for 'f' which is the crypto type passed in position 3 of hash ( $0f$ ) + { + u32 ks_camellia[68]; + + ctx.opad.h[0] = hc_swap32_S (ctx.opad.h[0]); + ctx.opad.h[1] = hc_swap32_S (ctx.opad.h[1]); + ctx.opad.h[2] = hc_swap32_S (ctx.opad.h[2]); + ctx.opad.h[3] = hc_swap32_S (ctx.opad.h[3]); + ctx.opad.h[4] = hc_swap32_S (ctx.opad.h[4]); + ctx.opad.h[5] = hc_swap32_S (ctx.opad.h[5]); + ctx.opad.h[6] = hc_swap32_S (ctx.opad.h[6]); + ctx.opad.h[7] = hc_swap32_S (ctx.opad.h[7]); + + camellia256_set_key (ks_camellia, ctx.opad.h); + + for (u32 i = 0; i < 20; i += 4) // 96 bytes output would contain the full 32 byte checksum + { + u32 data[4]; + + data[0] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[i + 0]; + data[1] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[i + 1]; + data[2] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[i + 2]; + data[3] = esalt_bufs[DIGESTS_OFFSET_HOST].ciphertext[i + 3]; + + + u32 out[4]; + + camellia256_decrypt (ks_camellia, data, out); + + res[i + 0] = hc_swap32_S (out[0] ^ iv[0]); + res[i + 1] = hc_swap32_S (out[1] ^ iv[1]); + res[i + 2] = hc_swap32_S (out[2] ^ iv[2]); + res[i + 3] = hc_swap32_S (out[3] ^ iv[3]); + + iv[0] = data[0]; + iv[1] = data[1]; + iv[2] = data[2]; + iv[3] = data[3]; + } + } + + u32 digest[8]; + + digest[0] = SHA256M_A; + digest[1] = SHA256M_B; + digest[2] = SHA256M_C; + digest[3] = SHA256M_D; + digest[4] = SHA256M_E; + digest[5] = SHA256M_F; + digest[6] = SHA256M_G; + digest[7] = SHA256M_H; + + w0[0] = res[ 0]; + w0[1] = res[ 1]; + w0[2] = res[ 2]; + w0[3] = res[ 3]; + w1[0] = res[ 4]; + w1[1] = res[ 5]; + w1[2] = res[ 6]; + w1[3] = res[ 7]; + w2[0] = res[ 8]; + w2[1] = res[ 9]; + w2[2] = res[10]; + w2[3] = res[11]; + w3[0] = res[12]; + w3[1] = res[13]; + w3[2] = res[14]; + w3[3] = res[15]; + + sha256_transform (w0, w1, w2, w3, digest); + + w0[0] = 0x80000000; + w0[1] = 0; + w0[2] = 0; + w0[3] = 0; + w1[0] = 0; + w1[1] = 0; + w1[2] = 0; + w1[3] = 0; + w2[0] = 0; + w2[1] = 0; + w2[2] = 0; + w2[3] = 0; + w3[0] = 0; + w3[1] = 0; + w3[2] = 0; + w3[3] = 64 * 8; + + sha256_transform (w0, w1, w2, w3, digest); + + if ((digest[0] == res[16]) && + (digest[1] == res[17]) && + (digest[2] == res[18]) && + (digest[3] == res[19])) + { + if (hc_atomic_inc (&hashes_shown[DIGESTS_OFFSET_HOST]) == 0) + { + mark_hash (plains_buf, d_return_buf, SALT_POS_HOST, DIGESTS_CNT, 0, DIGESTS_OFFSET_HOST + 0, gid, 0, 0, 0); + } + + return; + } +} diff --git a/docs/changes.txt b/docs/changes.txt index 88c5b9216..ce74e562d 100644 --- a/docs/changes.txt +++ b/docs/changes.txt @@ -8,6 +8,7 @@ - Added hash-mode: Adobe AEM (SSPR, SHA-256 with Salt) - Added hash-mode: Adobe AEM (SSPR, SHA-512 with Salt) - Added hash-mode: Anope IRC Services (enc_sha256) +- Added hash-mode: BestCrypt v4 Volume Encryption - Added hash-mode: Bisq .wallet (scrypt) - Added hash-mode: Bitcoin raw private key (P2PKH) - Added hash-mode: Bitcoin raw private key (P2SH(P2WPKH)) diff --git a/docs/readme.txt b/docs/readme.txt index 886de7dd8..38d1bb594 100644 --- a/docs/readme.txt +++ b/docs/readme.txt @@ -325,6 +325,7 @@ NVIDIA GPUs require "NVIDIA Driver" (440.64 or later) and "CUDA Toolkit" (9.0 or - LUKS v1 - VeraCrypt - BestCrypt v3 Volume Encryption +- BestCrypt v4 Volume Encryption - FileVault 2 - VirtualBox (PBKDF2-HMAC-SHA256 & AES-128-XTS) - VirtualBox (PBKDF2-HMAC-SHA256 & AES-256-XTS) diff --git a/src/modules/module_24000.c b/src/modules/module_24000.c new file mode 100644 index 000000000..a3a69e488 --- /dev/null +++ b/src/modules/module_24000.c @@ -0,0 +1,533 @@ +/** + * Author......: See docs/credits.txt + * License.....: MIT + */ + +#include "common.h" +#include "types.h" +#include "modules.h" +#include "bitops.h" +#include "convert.h" +#include "shared.h" + +static const u32 ATTACK_EXEC = ATTACK_EXEC_OUTSIDE_KERNEL; +static const u32 DGST_POS0 = 0; +static const u32 DGST_POS1 = 1; +static const u32 DGST_POS2 = 2; +static const u32 DGST_POS3 = 3; +static const u32 DGST_SIZE = DGST_SIZE_4_8; +static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; +static const char *HASH_NAME = "BestCrypt v4 Volume Encryption"; +static const u64 KERN_TYPE = 24000; +static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE; +static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE + | OPTS_TYPE_PT_GENERATE_LE + | OPTS_TYPE_MP_MULTI_DISABLE + | OPTS_TYPE_SELF_TEST_DISABLE + | OPTS_TYPE_ST_HEX; +static const u32 SALT_TYPE = SALT_TYPE_EMBEDDED; +static const char *ST_PASS = "hashcat"; +static const char *ST_HASH = "$bcve$4$08$323631353538333233323034363039393534383233393530$9f7892b8324b1d8cd36b5f2f8705b407131261620a89370db8369046646f5f82b96780453948db53b04928ae0cc47066f13454b34e31b58ea44ce943bcba14fcbd87f17205a31a896df182629ceea164d87e9e29127e8d865ca0bee52f832723"; + +u32 module_attack_exec (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return ATTACK_EXEC; } +u32 module_dgst_pos0 (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return DGST_POS0; } +u32 module_dgst_pos1 (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return DGST_POS1; } +u32 module_dgst_pos2 (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return DGST_POS2; } +u32 module_dgst_pos3 (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return DGST_POS3; } +u32 module_dgst_size (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return DGST_SIZE; } +u32 module_hash_category (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return HASH_CATEGORY; } +const char *module_hash_name (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return HASH_NAME; } +u64 module_kern_type (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return KERN_TYPE; } +u32 module_opti_type (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return OPTI_TYPE; } +u64 module_opts_type (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return OPTS_TYPE; } +u32 module_salt_type (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return SALT_TYPE; } +const char *module_st_hash (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return ST_HASH; } +const char *module_st_pass (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return ST_PASS; } + +typedef struct bestcrypt_scrypt +{ + u32 salt_buf[24]; + u32 ciphertext[96]; + u32 version; + +} bestcrypt_scrypt_t; + +// limit scrypt accel otherwise we hurt ourself when calculating the scrypt tmto +// 16 is actually a bit low, we may need to change this depending on user response + +static const char *SIGNATURE_BESTCRYPT_SCRYPT = "$bcve$"; +static const u32 SCRYPT_MAX_ACCEL = 256; +static const u32 SCRYPT_MAX_THREADS = 4; + +static const u64 SCRYPT_N = 32768; +static const u64 SCRYPT_R = 16; +static const u64 SCRYPT_P = 1; + +u32 module_kernel_accel_min (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) +{ + const u32 kernel_accel_min = 1; + + return kernel_accel_min; +} + +u32 module_kernel_accel_max (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) +{ + const u32 kernel_accel_max = (user_options->kernel_accel_chgd == true) ? user_options->kernel_accel : SCRYPT_MAX_ACCEL; + + return kernel_accel_max; +} + +u32 module_kernel_loops_min (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) +{ + const u32 kernel_loops_min = 1; + + return kernel_loops_min; +} + +u32 module_kernel_loops_max (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) +{ + const u32 kernel_loops_max = 1; + + return kernel_loops_max; +} + +u32 module_kernel_threads_min (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) +{ + const u32 kernel_threads_min = 1; + + return kernel_threads_min; +} + +u32 module_kernel_threads_max (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) +{ + const u32 kernel_threads_max = (user_options->kernel_threads_chgd == true) ? user_options->kernel_threads : SCRYPT_MAX_THREADS; + + return kernel_threads_max; +} + +u64 module_esalt_size (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) +{ + const u64 esalt_size = (const u64) sizeof (bestcrypt_scrypt_t); + + return esalt_size; +} + +u32 module_pw_max (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) +{ + // this overrides the reductions of PW_MAX in case optimized kernel is selected + // IOW, even in optimized kernel mode it support length 256 + + const u32 pw_max = PW_MAX; + + return pw_max; +} + +u64 module_extra_buffer_size (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra, MAYBE_UNUSED const hashes_t *hashes, MAYBE_UNUSED const hc_device_param_t *device_param) +{ + // we need to set the self-test hash settings to pass the self-test + // the decoder for the self-test is called after this function + + const u32 scrypt_N = (hashes->salts_buf[0].scrypt_N) ? hashes->salts_buf[0].scrypt_N : SCRYPT_N; + const u32 scrypt_r = (hashes->salts_buf[0].scrypt_r) ? hashes->salts_buf[0].scrypt_r : SCRYPT_R; + + const u64 kernel_power_max = ((OPTS_TYPE & OPTS_TYPE_MP_MULTI_DISABLE) ? 1 : device_param->device_processors) * device_param->kernel_threads_max * device_param->kernel_accel_max; + + u32 tmto_start = 1; + u32 tmto_stop = 6; + + if (user_options->scrypt_tmto) + { + tmto_start = user_options->scrypt_tmto; + tmto_stop = user_options->scrypt_tmto; + } + + // size_pws + + const u64 size_pws = kernel_power_max * sizeof (pw_t); + + const u64 size_pws_amp = size_pws; + + // size_pws_comp + + const u64 size_pws_comp = kernel_power_max * (sizeof (u32) * 64); + + // size_pws_idx + + const u64 size_pws_idx = (kernel_power_max + 1) * sizeof (pw_idx_t); + + // size_tmps + + const u64 size_tmps = kernel_power_max * hashconfig->tmp_size; + + // size_hooks + + const u64 size_hooks = kernel_power_max * hashconfig->hook_size; + +/* + u64 size_pws_pre = 4; + u64 size_pws_base = 4; + + if (user_options->slow_candidates == true) + { + // size_pws_pre + + size_pws_pre = kernel_power_max * sizeof (pw_pre_t); + + // size_pws_base + + size_pws_base = kernel_power_max * sizeof (pw_pre_t); + } +*/ + + // sometimes device_available_mem and device_maxmem_alloc reported back from the opencl runtime are a bit inaccurate. + // let's add some extra space just to be sure. + // now depends on the kernel-accel value (where scrypt and similar benefits), but also hard minimum 64mb and maximum 1024mb limit +/* + u64 EXTRA_SPACE = (1024ULL * 1024ULL) * device_param->kernel_accel_max; + + EXTRA_SPACE = MAX (EXTRA_SPACE, ( 64ULL * 1024ULL * 1024ULL)); + EXTRA_SPACE = MIN (EXTRA_SPACE, (1024ULL * 1024ULL * 1024ULL)); +*/ + const u64 scrypt_extra_space + = device_param->size_bfs + + device_param->size_combs + + device_param->size_digests + + device_param->size_esalts + + device_param->size_markov_css + + device_param->size_plains + + device_param->size_results + + device_param->size_root_css + + device_param->size_rules + + device_param->size_rules_c + + device_param->size_salts + + device_param->size_shown + + device_param->size_tm + + device_param->size_st_digests + + device_param->size_st_salts + + device_param->size_st_esalts + + size_pws + + size_pws_amp + + size_pws_comp + + size_pws_idx + + size_tmps + + size_hooks; +// + size_pws_pre +// + size_pws_base; +/* + + EXTRA_SPACE; +*/ + bool not_enough_memory = true; + + u64 size_scrypt = 0; + + u32 tmto; + + for (tmto = tmto_start; tmto <= tmto_stop; tmto++) + { + size_scrypt = (128ULL * scrypt_r) * scrypt_N; + + size_scrypt /= 1ull << tmto; + + size_scrypt *= kernel_power_max; + + if ((size_scrypt / 4) > device_param->device_maxmem_alloc) continue; + + if ((size_scrypt + scrypt_extra_space) > device_param->device_available_mem) continue; + + not_enough_memory = false; + + break; + } + + if (not_enough_memory == true) return -1; + + return size_scrypt; +} + +u64 module_tmp_size (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) +{ + const u64 tmp_size = 0; // we'll add some later + + return tmp_size; +} + +u64 module_extra_tmp_size (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra, MAYBE_UNUSED const hashes_t *hashes) +{ + const u64 scrypt_N = (hashes->salts_buf[0].scrypt_N) ? hashes->salts_buf[0].scrypt_N : SCRYPT_N; + const u64 scrypt_r = (hashes->salts_buf[0].scrypt_r) ? hashes->salts_buf[0].scrypt_r : SCRYPT_R; + const u64 scrypt_p = (hashes->salts_buf[0].scrypt_p) ? hashes->salts_buf[0].scrypt_p : SCRYPT_P; + + // we need to check that all hashes have the same scrypt settings + + for (u32 i = 1; i < hashes->salts_cnt; i++) + { + if ((hashes->salts_buf[i].scrypt_N != scrypt_N) + || (hashes->salts_buf[i].scrypt_r != scrypt_r) + || (hashes->salts_buf[i].scrypt_p != scrypt_p)) + { + return -1; + } + } + + const u64 tmp_size = 128ULL * scrypt_r * scrypt_p; + + return tmp_size; +} + +bool module_unstable_warning (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra, MAYBE_UNUSED const hc_device_param_t *device_param) +{ + if (device_param->opencl_platform_vendor_id == VENDOR_ID_APPLE) + { + // Invalid extra buffer size. + if ((device_param->is_metal == false) && (device_param->opencl_device_type & CL_DEVICE_TYPE_GPU)) + { + return true; + } +/* + if ((device_param->opencl_device_vendor_id == VENDOR_ID_INTEL_SDK) && (device_param->opencl_device_type & CL_DEVICE_TYPE_GPU)) + { + return true; + } +*/ + } + + // amdgpu-pro-18.50-708488-ubuntu-18.04: Segmentation fault + if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == false)) + { + return true; + } + + return false; +} + +bool module_warmup_disable (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) +{ + return true; +} + +char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra, MAYBE_UNUSED const hashes_t *hashes, MAYBE_UNUSED const hc_device_param_t *device_param) +{ + const u64 scrypt_N = (hashes->salts_buf[0].scrypt_N) ? hashes->salts_buf[0].scrypt_N : SCRYPT_N; + const u64 scrypt_r = (hashes->salts_buf[0].scrypt_r) ? hashes->salts_buf[0].scrypt_r : SCRYPT_R; + const u64 scrypt_p = (hashes->salts_buf[0].scrypt_p) ? hashes->salts_buf[0].scrypt_p : SCRYPT_P; + + const u64 extra_buffer_size = device_param->extra_buffer_size; + + const u64 kernel_power_max = ((OPTS_TYPE & OPTS_TYPE_MP_MULTI_DISABLE) ? 1 : device_param->device_processors) * device_param->kernel_threads_max * device_param->kernel_accel_max; + + const u64 size_scrypt = 128ULL * scrypt_r * scrypt_N; + + const u64 scrypt_tmto_final = (kernel_power_max * size_scrypt) / extra_buffer_size; + + const u64 tmp_size = 128ULL * scrypt_r * scrypt_p; + + char *jit_build_options = NULL; + + hc_asprintf (&jit_build_options, "-D SCRYPT_N=%u -D SCRYPT_R=%u -D SCRYPT_P=%u -D SCRYPT_TMTO=%" PRIu64 " -D SCRYPT_TMP_ELEM=%" PRIu64, + hashes->salts_buf[0].scrypt_N, + hashes->salts_buf[0].scrypt_r, + hashes->salts_buf[0].scrypt_p, + scrypt_tmto_final, + tmp_size / 16); + + return jit_build_options; +} + +int module_hash_decode (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED void *digest_buf, MAYBE_UNUSED salt_t *salt, MAYBE_UNUSED void *esalt_buf, MAYBE_UNUSED void *hook_salt_buf, MAYBE_UNUSED hashinfo_t *hash_info, const char *line_buf, MAYBE_UNUSED const int line_len) +{ + u32 *digest = (u32 *) digest_buf; + + bestcrypt_scrypt_t *bestcrypt_scrypt = (bestcrypt_scrypt_t *) esalt_buf; + + hc_token_t token; + + memset (&token, 0, sizeof (hc_token_t)); + + token.token_cnt = 5; + + token.signatures_cnt = 1; + token.signatures_buf[0] = SIGNATURE_BESTCRYPT_SCRYPT; + + token.len[0] = 6; + token.attr[0] = TOKEN_ATTR_FIXED_LENGTH + | TOKEN_ATTR_VERIFY_SIGNATURE; + + token.sep[1] = '$'; + token.len[1] = 1; + token.attr[1] = TOKEN_ATTR_FIXED_LENGTH + | TOKEN_ATTR_VERIFY_DIGIT; + + token.sep[2] = '$'; + token.len[2] = 2; + token.attr[2] = TOKEN_ATTR_FIXED_LENGTH + | TOKEN_ATTR_VERIFY_HEX; + + token.sep[3] = '$'; + token.len[3] = 48; + token.attr[3] = TOKEN_ATTR_FIXED_LENGTH + | TOKEN_ATTR_VERIFY_HEX; + + token.len[4] = 192; + token.attr[4] = TOKEN_ATTR_FIXED_LENGTH + | TOKEN_ATTR_VERIFY_HEX; + + const int rc_tokenizer = input_tokenizer ((const u8 *) line_buf, line_len, &token); + + if (rc_tokenizer != PARSER_OK) return (rc_tokenizer); + + const u8 *format_type_pos = token.buf[1]; + + if (format_type_pos[0] != '4') return (PARSER_SALT_VALUE); + + const u8 *crypto_type_pos = token.buf[2]; + + bestcrypt_scrypt->version = crypto_type_pos[1]; + + // scrypt settings + + const u32 scrypt_N = SCRYPT_N; + const u32 scrypt_r = SCRYPT_R; + const u32 scrypt_p = SCRYPT_P; + + salt->scrypt_N = scrypt_N; + salt->scrypt_r = scrypt_r; + salt->scrypt_p = scrypt_p; + + // salt + + const u8 *salt_pos = token.buf[3]; + const int salt_len = token.len[3]; + + const bool parse_rc = generic_salt_decode (hashconfig, salt_pos, salt_len, (u8 *) salt->salt_buf, (int *) &salt->salt_len); + + if (parse_rc == false) return (PARSER_SALT_LENGTH); + + bestcrypt_scrypt->salt_buf[0] = salt->salt_buf[0]; + bestcrypt_scrypt->salt_buf[1] = salt->salt_buf[1]; + + bestcrypt_scrypt->salt_buf[0] = byte_swap_32 (salt->salt_buf[0]); + bestcrypt_scrypt->salt_buf[1] = byte_swap_32 (salt->salt_buf[1]); + + salt->salt_iter = 1; + +// salt->salt_iter = salt->scrypt_N; + salt->salt_repeats = salt->scrypt_p - 1; + + // ciphertext + + const u8 *ciphertext_pos = token.buf[4]; + const int ciphertext_len = token.len[4]; + + hex_decode (ciphertext_pos, ciphertext_len, (u8 *) bestcrypt_scrypt->ciphertext); + + digest[0] = bestcrypt_scrypt->ciphertext[16]; + digest[1] = bestcrypt_scrypt->ciphertext[17]; + digest[2] = bestcrypt_scrypt->ciphertext[18]; + digest[3] = bestcrypt_scrypt->ciphertext[19]; + + return (PARSER_OK); +} + +int module_hash_encode (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const void *digest_buf, MAYBE_UNUSED const salt_t *salt, MAYBE_UNUSED const void *esalt_buf, MAYBE_UNUSED const void *hook_salt_buf, MAYBE_UNUSED const hashinfo_t *hash_info, char *line_buf, MAYBE_UNUSED const int line_size) +{ + char tmp_salt[SALT_MAX * 2]; + + const int salt_len = generic_salt_encode (hashconfig, (const u8 *) salt->salt_buf, (const int) salt->salt_len, (u8 *) tmp_salt); + + tmp_salt[salt_len] = 0; + + bestcrypt_scrypt_t *bestcrypt_scrypt = (bestcrypt_scrypt_t *) esalt_buf; + + u8 data_hex[193] = { 0 }; + + hex_encode ((u8 *) bestcrypt_scrypt->ciphertext, 96, data_hex); + + const int line_len = snprintf (line_buf, line_size, "%s4$0%c$%s$%s", + SIGNATURE_BESTCRYPT_SCRYPT, + bestcrypt_scrypt->version, + (char *) tmp_salt, + data_hex + ); + + return line_len; +} + +void module_init (module_ctx_t *module_ctx) +{ + module_ctx->module_context_size = MODULE_CONTEXT_SIZE_CURRENT; + module_ctx->module_interface_version = MODULE_INTERFACE_VERSION_CURRENT; + + module_ctx->module_attack_exec = module_attack_exec; + module_ctx->module_benchmark_esalt = MODULE_DEFAULT; + module_ctx->module_benchmark_hook_salt = MODULE_DEFAULT; + module_ctx->module_benchmark_mask = MODULE_DEFAULT; + module_ctx->module_benchmark_charset = MODULE_DEFAULT; + module_ctx->module_benchmark_salt = MODULE_DEFAULT; + module_ctx->module_build_plain_postprocess = MODULE_DEFAULT; + module_ctx->module_deep_comp_kernel = MODULE_DEFAULT; + module_ctx->module_deprecated_notice = MODULE_DEFAULT; + module_ctx->module_dgst_pos0 = module_dgst_pos0; + module_ctx->module_dgst_pos1 = module_dgst_pos1; + module_ctx->module_dgst_pos2 = module_dgst_pos2; + module_ctx->module_dgst_pos3 = module_dgst_pos3; + module_ctx->module_dgst_size = module_dgst_size; + module_ctx->module_dictstat_disable = MODULE_DEFAULT; + module_ctx->module_esalt_size = module_esalt_size; + module_ctx->module_extra_buffer_size = module_extra_buffer_size; + module_ctx->module_extra_tmp_size = module_extra_tmp_size; + module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT; + module_ctx->module_forced_outfile_format = MODULE_DEFAULT; + module_ctx->module_hash_binary_count = MODULE_DEFAULT; + module_ctx->module_hash_binary_parse = MODULE_DEFAULT; + module_ctx->module_hash_binary_save = MODULE_DEFAULT; + module_ctx->module_hash_decode_postprocess = MODULE_DEFAULT; + module_ctx->module_hash_decode_potfile = MODULE_DEFAULT; + module_ctx->module_hash_decode_zero_hash = MODULE_DEFAULT; + module_ctx->module_hash_decode = module_hash_decode; + module_ctx->module_hash_encode_status = MODULE_DEFAULT; + module_ctx->module_hash_encode_potfile = MODULE_DEFAULT; + module_ctx->module_hash_encode = module_hash_encode; + module_ctx->module_hash_init_selftest = MODULE_DEFAULT; + module_ctx->module_hash_mode = MODULE_DEFAULT; + module_ctx->module_hash_category = module_hash_category; + module_ctx->module_hash_name = module_hash_name; + module_ctx->module_hashes_count_min = MODULE_DEFAULT; + module_ctx->module_hashes_count_max = MODULE_DEFAULT; + module_ctx->module_hlfmt_disable = MODULE_DEFAULT; + module_ctx->module_hook_extra_param_size = MODULE_DEFAULT; + module_ctx->module_hook_extra_param_init = MODULE_DEFAULT; + module_ctx->module_hook_extra_param_term = MODULE_DEFAULT; + module_ctx->module_hook12 = MODULE_DEFAULT; + module_ctx->module_hook23 = MODULE_DEFAULT; + module_ctx->module_hook_salt_size = MODULE_DEFAULT; + module_ctx->module_hook_size = MODULE_DEFAULT; + module_ctx->module_jit_build_options = module_jit_build_options; + module_ctx->module_jit_cache_disable = MODULE_DEFAULT; + module_ctx->module_kernel_accel_max = module_kernel_accel_max; + module_ctx->module_kernel_accel_min = module_kernel_accel_min; + module_ctx->module_kernel_loops_max = module_kernel_loops_max; + module_ctx->module_kernel_loops_min = module_kernel_loops_min; + module_ctx->module_kernel_threads_max = module_kernel_threads_max; + module_ctx->module_kernel_threads_min = module_kernel_threads_min; + module_ctx->module_kern_type = module_kern_type; + module_ctx->module_kern_type_dynamic = MODULE_DEFAULT; + module_ctx->module_opti_type = module_opti_type; + module_ctx->module_opts_type = module_opts_type; + module_ctx->module_outfile_check_disable = MODULE_DEFAULT; + module_ctx->module_outfile_check_nocomp = MODULE_DEFAULT; + module_ctx->module_potfile_custom_check = MODULE_DEFAULT; + module_ctx->module_potfile_disable = MODULE_DEFAULT; + module_ctx->module_potfile_keep_all_hashes = MODULE_DEFAULT; + module_ctx->module_pwdump_column = MODULE_DEFAULT; + module_ctx->module_pw_max = module_pw_max; + module_ctx->module_pw_min = MODULE_DEFAULT; + module_ctx->module_salt_max = MODULE_DEFAULT; + module_ctx->module_salt_min = MODULE_DEFAULT; + module_ctx->module_salt_type = module_salt_type; + module_ctx->module_separator = MODULE_DEFAULT; + module_ctx->module_st_hash = module_st_hash; + module_ctx->module_st_pass = module_st_pass; + module_ctx->module_tmp_size = module_tmp_size; + module_ctx->module_unstable_warning = module_unstable_warning; + module_ctx->module_warmup_disable = module_warmup_disable; +} diff --git a/tools/install_modules.sh b/tools/install_modules.sh index 5167a9f31..683d20000 100755 --- a/tools/install_modules.sh +++ b/tools/install_modules.sh @@ -19,7 +19,9 @@ cpan install Authen::Passphrase::LANManager \ Compress::Zlib \ Convert::EBCDIC \ Crypt::AuthEnc::GCM \ + Crypt::Camellia \ Crypt::CBC \ + Crypt::Cipher::Serpent \ Crypt::DES \ Crypt::DES_EDE3 \ Crypt::Digest::BLAKE2s_256 \ diff --git a/tools/test_modules/m24000.pm b/tools/test_modules/m24000.pm new file mode 100644 index 000000000..7f84b373c --- /dev/null +++ b/tools/test_modules/m24000.pm @@ -0,0 +1,144 @@ +#!/usr/bin/env perl + +## +## Author......: See docs/credits.txt +## License.....: MIT +## + +use strict; +use warnings; + +use Crypt::ScryptKDF qw (scrypt_raw); +use Digest::SHA qw (sha256); +use Crypt::CBC; +use Crypt::Rijndael; +use Crypt::Twofish; +use Crypt::Cipher::Serpent; # Crypt::Serpent doesn't work for me :( +use Crypt::Camellia; + +sub module_constraints { [[0, 256], [24, 24], [-1, -1], [-1, -1], [-1, -1]] } + +my $SCRYPT_N = 32768; +my $SCRYPT_R = 16; +my $SCRYPT_P = 1; + +sub module_generate_hash +{ + my $word = shift; + my $salt = shift; + my $data = shift; + my $type = shift; + + # most heavy part of the algorithm: + + my $key = scrypt_raw ($word, $salt, $SCRYPT_N, $SCRYPT_R, $SCRYPT_P, 32); + + my %crypto_types = ( + 8 => 'Crypt::Rijndael', # '08' => AES + 9 => 'Crypt::Twofish', # '09' => Twofish + 10 => 'Crypt::Cipher::Serpent', # '0a' => Serpent + 15 => 'Crypt::Camellia' # '0f' => Camellia + ); + + my @crypto_type_conv = (8, 9, 10, 15); + + if (! defined ($type)) + { + my $rand_type_num = random_number (0, 3); + + $type = $crypto_type_conv[$rand_type_num]; + } + + my $crypto_algo = $crypto_types{$type}; + + my $crypt = Crypt::CBC->new ({ + cipher => $crypto_algo, + key => $key, + iv => "\x00" x 16, + keysize => 32, + literal_key => 1, + header => "none", + padding => "none" + }); + + if (defined ($data)) # decrypt + { + my $plain_text = $crypt->decrypt ($data); + + my $part1 = substr ($plain_text, 0, 64); + my $part2 = substr ($plain_text, 64, 32); + + my $hash = sha256 ($part1); + + if ($hash ne $part2) # wrong => fake the data + { + $data = "\x00" x length ($data); # 64 + 32 = 96 + } + } + else # encrypt + { + $data = random_bytes (64); + + my $hash = sha256 ($data); + + $data = $crypt->encrypt ($data . $hash); + } + + return sprintf ("\$bcve\$4\$%02x\$%s\$%s", $type, unpack ("H*", $salt), unpack ("H*", $data)); +} + +sub module_verify_hash +{ + my $line = shift; + + my $idx1 = index ($line, ':'); + + return if ($idx1 < 1); + + my $hash = substr ($line, 0, $idx1); + my $word = substr ($line, $idx1 + 1); + + return if (substr ($hash, 0, 8) ne "\$bcve\$4\$"); + + $idx1 = index ($hash, '$', 8); + + return if ($idx1 < 1); + + # crypto type + + my $crypto_type = substr ($hash, 8, $idx1 - 8); + + return unless ($crypto_type eq "08") || + ($crypto_type eq "09") || + ($crypto_type eq "0a") || + ($crypto_type eq "0f"); + + $crypto_type = hex ($crypto_type); + + # salt + + my $idx2 = index ($hash, '$', $idx1 + 1); + + my $salt = substr ($hash, $idx1 + 1, $idx2 - $idx1 - 1); + + return unless ($salt =~ m/^[0-9a-fA-F]+$/); + + # data + + my $data = substr ($hash, $idx2 + 1); + + return unless ($data =~ m/^[0-9a-fA-F]+$/); + + # convert to hex: + + $salt = pack ("H*", $salt); + $data = pack ("H*", $data); + + my $word_packed = pack_if_HEX_notation ($word); + + my $new_hash = module_generate_hash ($word_packed, $salt, $data, $crypto_type); + + return ($new_hash, $word); +} + +1; diff --git a/tunings/Module_24000.hctune b/tunings/Module_24000.hctune new file mode 100644 index 000000000..11fb49268 --- /dev/null +++ b/tunings/Module_24000.hctune @@ -0,0 +1,5 @@ +#Device Attack Hash Vector Kernel Kernel +#Name Mode Type Width Accel Loops + +DEVICE_TYPE_CPU * 24000 1 N A +DEVICE_TYPE_GPU * 24000 1 1 A