1
0
mirror of https://github.com/hashcat/hashcat.git synced 2024-12-22 22:58:30 +00:00

Merge pull request #1793 from Naufragous/veracrypt-kuznyechik

Add Kuznyechik support for VeraCrypt kernels
This commit is contained in:
Jens Steube 2018-11-22 16:31:02 +01:00 committed by GitHub
commit 38e97bd89a
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23
31 changed files with 812 additions and 15 deletions

View File

@ -0,0 +1,303 @@
/* *
* This is an OpenCL implementation of the encryption algorithm: *
* *
* GOST R 34.12-2015 Kuznyechik by A.S.Kuzmin and A.A.Nechaev *
* *
* Author of the original C implementation: *
* *
* Markku-Juhani O. Saarinen <mjos@iki.fi> *
* https://github.com/mjosaarinen/kuznechik *
* *
* Adapted for GPU use with hashcat by Ruslan Yushaev. *
* *
*/
__constant u32a k_sbox[256] =
{
0xfc, 0xee, 0xdd, 0x11, 0xcf, 0x6e, 0x31, 0x16,
0xfb, 0xc4, 0xfa, 0xda, 0x23, 0xc5, 0x04, 0x4d,
0xe9, 0x77, 0xf0, 0xdb, 0x93, 0x2e, 0x99, 0xba,
0x17, 0x36, 0xf1, 0xbb, 0x14, 0xcd, 0x5f, 0xc1,
0xf9, 0x18, 0x65, 0x5a, 0xe2, 0x5c, 0xef, 0x21,
0x81, 0x1c, 0x3c, 0x42, 0x8b, 0x01, 0x8e, 0x4f,
0x05, 0x84, 0x02, 0xae, 0xe3, 0x6a, 0x8f, 0xa0,
0x06, 0x0b, 0xed, 0x98, 0x7f, 0xd4, 0xd3, 0x1f,
0xeb, 0x34, 0x2c, 0x51, 0xea, 0xc8, 0x48, 0xab,
0xf2, 0x2a, 0x68, 0xa2, 0xfd, 0x3a, 0xce, 0xcc,
0xb5, 0x70, 0x0e, 0x56, 0x08, 0x0c, 0x76, 0x12,
0xbf, 0x72, 0x13, 0x47, 0x9c, 0xb7, 0x5d, 0x87,
0x15, 0xa1, 0x96, 0x29, 0x10, 0x7b, 0x9a, 0xc7,
0xf3, 0x91, 0x78, 0x6f, 0x9d, 0x9e, 0xb2, 0xb1,
0x32, 0x75, 0x19, 0x3d, 0xff, 0x35, 0x8a, 0x7e,
0x6d, 0x54, 0xc6, 0x80, 0xc3, 0xbd, 0x0d, 0x57,
0xdf, 0xf5, 0x24, 0xa9, 0x3e, 0xa8, 0x43, 0xc9,
0xd7, 0x79, 0xd6, 0xf6, 0x7c, 0x22, 0xb9, 0x03,
0xe0, 0x0f, 0xec, 0xde, 0x7a, 0x94, 0xb0, 0xbc,
0xdc, 0xe8, 0x28, 0x50, 0x4e, 0x33, 0x0a, 0x4a,
0xa7, 0x97, 0x60, 0x73, 0x1e, 0x00, 0x62, 0x44,
0x1a, 0xb8, 0x38, 0x82, 0x64, 0x9f, 0x26, 0x41,
0xad, 0x45, 0x46, 0x92, 0x27, 0x5e, 0x55, 0x2f,
0x8c, 0xa3, 0xa5, 0x7d, 0x69, 0xd5, 0x95, 0x3b,
0x07, 0x58, 0xb3, 0x40, 0x86, 0xac, 0x1d, 0xf7,
0x30, 0x37, 0x6b, 0xe4, 0x88, 0xd9, 0xe7, 0x89,
0xe1, 0x1b, 0x83, 0x49, 0x4c, 0x3f, 0xf8, 0xfe,
0x8d, 0x53, 0xaa, 0x90, 0xca, 0xd8, 0x85, 0x61,
0x20, 0x71, 0x67, 0xa4, 0x2d, 0x2b, 0x09, 0x5b,
0xcb, 0x9b, 0x25, 0xd0, 0xbe, 0xe5, 0x6c, 0x52,
0x59, 0xa6, 0x74, 0xd2, 0xe6, 0xf4, 0xb4, 0xc0,
0xd1, 0x66, 0xaf, 0xc2, 0x39, 0x4b, 0x63, 0xb6
};
__constant u32a k_sbox_inv[256] =
{
0xa5, 0x2d, 0x32, 0x8f, 0x0e, 0x30, 0x38, 0xc0,
0x54, 0xe6, 0x9e, 0x39, 0x55, 0x7e, 0x52, 0x91,
0x64, 0x03, 0x57, 0x5a, 0x1c, 0x60, 0x07, 0x18,
0x21, 0x72, 0xa8, 0xd1, 0x29, 0xc6, 0xa4, 0x3f,
0xe0, 0x27, 0x8d, 0x0c, 0x82, 0xea, 0xae, 0xb4,
0x9a, 0x63, 0x49, 0xe5, 0x42, 0xe4, 0x15, 0xb7,
0xc8, 0x06, 0x70, 0x9d, 0x41, 0x75, 0x19, 0xc9,
0xaa, 0xfc, 0x4d, 0xbf, 0x2a, 0x73, 0x84, 0xd5,
0xc3, 0xaf, 0x2b, 0x86, 0xa7, 0xb1, 0xb2, 0x5b,
0x46, 0xd3, 0x9f, 0xfd, 0xd4, 0x0f, 0x9c, 0x2f,
0x9b, 0x43, 0xef, 0xd9, 0x79, 0xb6, 0x53, 0x7f,
0xc1, 0xf0, 0x23, 0xe7, 0x25, 0x5e, 0xb5, 0x1e,
0xa2, 0xdf, 0xa6, 0xfe, 0xac, 0x22, 0xf9, 0xe2,
0x4a, 0xbc, 0x35, 0xca, 0xee, 0x78, 0x05, 0x6b,
0x51, 0xe1, 0x59, 0xa3, 0xf2, 0x71, 0x56, 0x11,
0x6a, 0x89, 0x94, 0x65, 0x8c, 0xbb, 0x77, 0x3c,
0x7b, 0x28, 0xab, 0xd2, 0x31, 0xde, 0xc4, 0x5f,
0xcc, 0xcf, 0x76, 0x2c, 0xb8, 0xd8, 0x2e, 0x36,
0xdb, 0x69, 0xb3, 0x14, 0x95, 0xbe, 0x62, 0xa1,
0x3b, 0x16, 0x66, 0xe9, 0x5c, 0x6c, 0x6d, 0xad,
0x37, 0x61, 0x4b, 0xb9, 0xe3, 0xba, 0xf1, 0xa0,
0x85, 0x83, 0xda, 0x47, 0xc5, 0xb0, 0x33, 0xfa,
0x96, 0x6f, 0x6e, 0xc2, 0xf6, 0x50, 0xff, 0x5d,
0xa9, 0x8e, 0x17, 0x1b, 0x97, 0x7d, 0xec, 0x58,
0xf7, 0x1f, 0xfb, 0x7c, 0x09, 0x0d, 0x7a, 0x67,
0x45, 0x87, 0xdc, 0xe8, 0x4f, 0x1d, 0x4e, 0x04,
0xeb, 0xf8, 0xf3, 0x3e, 0x3d, 0xbd, 0x8a, 0x88,
0xdd, 0xcd, 0x0b, 0x13, 0x98, 0x02, 0x93, 0x80,
0x90, 0xd0, 0x24, 0x34, 0xcb, 0xed, 0xf4, 0xce,
0x99, 0x10, 0x44, 0x40, 0x92, 0x3a, 0x01, 0x26,
0x12, 0x1a, 0x48, 0x68, 0xf5, 0x81, 0x8b, 0xc7,
0xd6, 0x20, 0x0a, 0x08, 0x00, 0x4c, 0xd7, 0x74
};
__constant int k_vec[16] =
{
0x94, 0x20, 0x85, 0x10, 0xc2, 0xc0, 0x01, 0xfb,
0x01, 0xc0, 0xc2, 0x10, 0x85, 0x20, 0x94, 0x01
};
#define extract_byte(x,n) (((x) >> (8 * (n))) & 0xff)
#define k_lookup(w,sbox) \
for (int i = 0; i < 4; i++) \
w[i] = (sbox[extract_byte(w[i],0)] << 0) \
| (sbox[extract_byte(w[i],1)] << 8) \
| (sbox[extract_byte(w[i],2)] << 16) \
| (sbox[extract_byte(w[i],3)] << 24)
#define k_vec_xor(n) \
for (int y = k_vec[(n)]; y > 0; y >>= 1) \
{ \
z ^= x * (y & 1); \
x = ((x << 1) ^ ((x >> 7) * 0xc3)) & 0xff; \
}
DECLSPEC void kuznyechik_linear (u32 *w)
{
// used inside k_vec_xor macro
u32 x;
u32 z;
for (int i = 0; i < 16; i++)
{
z = 0;
x = extract_byte (w[3], 3); k_vec_xor (15);
x = extract_byte (w[3], 2); k_vec_xor (14);
x = extract_byte (w[3], 1); k_vec_xor (13);
x = extract_byte (w[3], 0); k_vec_xor (12);
x = extract_byte (w[2], 3); k_vec_xor (11);
x = extract_byte (w[2], 2); k_vec_xor (10);
x = extract_byte (w[2], 1); k_vec_xor ( 9);
x = extract_byte (w[2], 0); k_vec_xor ( 8);
x = extract_byte (w[1], 3); k_vec_xor ( 7);
x = extract_byte (w[1], 2); k_vec_xor ( 6);
x = extract_byte (w[1], 1); k_vec_xor ( 5);
x = extract_byte (w[1], 0); k_vec_xor ( 4);
x = extract_byte (w[0], 3); k_vec_xor ( 3);
x = extract_byte (w[0], 2); k_vec_xor ( 2);
x = extract_byte (w[0], 1); k_vec_xor ( 1);
x = extract_byte (w[0], 0); k_vec_xor ( 0);
// right-shift data block, prepend calculated byte
w[3] = (w[3] << 8) | (w[2] >> 24);
w[2] = (w[2] << 8) | (w[1] >> 24);
w[1] = (w[1] << 8) | (w[0] >> 24);
w[0] = (w[0] << 8) | z;
}
}
DECLSPEC void kuznyechik_linear_inv (u32 *w)
{
// used inside k_vec_xor macro
u32 x;
u32 z;
for (int i = 0; i < 16; i++)
{
z = extract_byte (w[0], 0);
//left-shift data block
w[0] = (w[0] >> 8) | (w[1] << 24);
w[1] = (w[1] >> 8) | (w[2] << 24);
w[2] = (w[2] >> 8) | (w[3] << 24);
w[3] = (w[3] >> 8);
x = extract_byte (w[0], 0); k_vec_xor ( 0);
x = extract_byte (w[0], 1); k_vec_xor ( 1);
x = extract_byte (w[0], 2); k_vec_xor ( 2);
x = extract_byte (w[0], 3); k_vec_xor ( 3);
x = extract_byte (w[1], 0); k_vec_xor ( 4);
x = extract_byte (w[1], 1); k_vec_xor ( 5);
x = extract_byte (w[1], 2); k_vec_xor ( 6);
x = extract_byte (w[1], 3); k_vec_xor ( 7);
x = extract_byte (w[2], 0); k_vec_xor ( 8);
x = extract_byte (w[2], 1); k_vec_xor ( 9);
x = extract_byte (w[2], 2); k_vec_xor (10);
x = extract_byte (w[2], 3); k_vec_xor (11);
x = extract_byte (w[3], 0); k_vec_xor (12);
x = extract_byte (w[3], 1); k_vec_xor (13);
x = extract_byte (w[3], 2); k_vec_xor (14);
//append calculated byte
w[3] |= (z << 24);
}
}
DECLSPEC void kuznyechik_set_key (u32 *ks, const u32 *ukey)
{
u32 counter[4];
u32 x[4];
u32 y[4];
u32 z[4];
x[0] = ukey[0];
x[1] = ukey[1];
x[2] = ukey[2];
x[3] = ukey[3];
y[0] = ukey[4];
y[1] = ukey[5];
y[2] = ukey[6];
y[3] = ukey[7];
ks[0] = ukey[0];
ks[1] = ukey[1];
ks[2] = ukey[2];
ks[3] = ukey[3];
ks[4] = ukey[4];
ks[5] = ukey[5];
ks[6] = ukey[6];
ks[7] = ukey[7];
for (int i = 1; i <= 32; i++)
{
counter[0] = 0;
counter[1] = 0;
counter[2] = 0;
counter[3] = (i << 24);
kuznyechik_linear (counter);
z[0] = x[0] ^ counter[0];
z[1] = x[1] ^ counter[1];
z[2] = x[2] ^ counter[2];
z[3] = x[3] ^ counter[3];
k_lookup (z, k_sbox);
kuznyechik_linear (z);
z[0] ^= y[0];
z[1] ^= y[1];
z[2] ^= y[2];
z[3] ^= y[3];
y[0] = x[0];
y[1] = x[1];
y[2] = x[2];
y[3] = x[3];
x[0] = z[0];
x[1] = z[1];
x[2] = z[2];
x[3] = z[3];
if ((i & 7) == 0)
{
ks[i + 0] = x[0];
ks[i + 1] = x[1];
ks[i + 2] = x[2];
ks[i + 3] = x[3];
ks[i + 4] = y[0];
ks[i + 5] = y[1];
ks[i + 6] = y[2];
ks[i + 7] = y[3];
}
}
}
DECLSPEC void kuznyechik_encrypt (const u32 *ks, const u32 *in, u32 *out)
{
out[0] = in[0];
out[1] = in[1];
out[2] = in[2];
out[3] = in[3];
for (int i = 0; i < 9; i++)
{
out[0] ^= ks[4 * i + 0];
out[1] ^= ks[4 * i + 1];
out[2] ^= ks[4 * i + 2];
out[3] ^= ks[4 * i + 3];
k_lookup (out, k_sbox);
kuznyechik_linear (out);
}
out[0] ^= ks[4 * 9 + 0];
out[1] ^= ks[4 * 9 + 1];
out[2] ^= ks[4 * 9 + 2];
out[3] ^= ks[4 * 9 + 3];
}
DECLSPEC void kuznyechik_decrypt (const u32 *ks, const u32 *in, u32 *out)
{
out[0] = in[0];
out[1] = in[1];
out[2] = in[2];
out[3] = in[3];
out[0] ^= ks[4 * 9 + 0];
out[1] ^= ks[4 * 9 + 1];
out[2] ^= ks[4 * 9 + 2];
out[3] ^= ks[4 * 9 + 3];
for (int i = 8; i >= 0; i--)
{
kuznyechik_linear_inv (out);
k_lookup (out, k_sbox_inv);
out[0] ^= ks[4 * i + 0];
out[1] ^= ks[4 * i + 1];
out[2] ^= ks[4 * i + 2];
out[3] ^= ks[4 * i + 3];
}
}

229
OpenCL/inc_veracrypt_xts.cl Normal file
View File

@ -0,0 +1,229 @@
DECLSPEC void kuznyechik_decrypt_xts_first (const u32 *ukey1, const u32 *ukey2, const u32 *in, u32 *out, u32 *S, u32 *T, u32 *ks)
{
out[0] = in[0];
out[1] = in[1];
out[2] = in[2];
out[3] = in[3];
kuznyechik_set_key (ks, ukey2);
kuznyechik_encrypt (ks, S, T);
out[0] ^= T[0];
out[1] ^= T[1];
out[2] ^= T[2];
out[3] ^= T[3];
kuznyechik_set_key (ks, ukey1);
kuznyechik_decrypt (ks, out, out);
out[0] ^= T[0];
out[1] ^= T[1];
out[2] ^= T[2];
out[3] ^= T[3];
}
DECLSPEC void kuznyechik_decrypt_xts_next (const u32 *in, u32 *out, u32 *T, u32 *ks)
{
out[0] = in[0];
out[1] = in[1];
out[2] = in[2];
out[3] = in[3];
xts_mul2 (T, T);
out[0] ^= T[0];
out[1] ^= T[1];
out[2] ^= T[2];
out[3] ^= T[3];
kuznyechik_decrypt (ks, out, out);
out[0] ^= T[0];
out[1] ^= T[1];
out[2] ^= T[2];
out[3] ^= T[3];
}
// 512 bit
DECLSPEC int verify_header_kuznyechik (__global const tc_t *esalt_bufs, const u32 *ukey1, const u32 *ukey2)
{
u32 ks_kuznyechik[40];
u32 S[4] = { 0 };
u32 T_kuznyechik[4] = { 0 };
u32 data[4];
data[0] = esalt_bufs[0].data_buf[0];
data[1] = esalt_bufs[0].data_buf[1];
data[2] = esalt_bufs[0].data_buf[2];
data[3] = esalt_bufs[0].data_buf[3];
u32 tmp[4];
kuznyechik_decrypt_xts_first (ukey1, ukey2, data, tmp, S, T_kuznyechik, ks_kuznyechik);
const u32 signature = esalt_bufs[0].signature;
if (tmp[0] != signature) return 0;
const u32 crc32_save = swap32_S (~tmp[2]);
// seek to byte 256
for (volatile int i = 4; i < 64 - 16; i += 4)
{
xts_mul2 (T_kuznyechik, T_kuznyechik);
}
// calculate crc32 from here
u32 crc32 = ~0;
for (volatile int i = 64 - 16; i < 128 - 16; i += 4)
{
data[0] = esalt_bufs[0].data_buf[i + 0];
data[1] = esalt_bufs[0].data_buf[i + 1];
data[2] = esalt_bufs[0].data_buf[i + 2];
data[3] = esalt_bufs[0].data_buf[i + 3];
kuznyechik_decrypt_xts_next (data, tmp, T_kuznyechik, ks_kuznyechik);
crc32 = round_crc32_4 (tmp[0], crc32);
crc32 = round_crc32_4 (tmp[1], crc32);
crc32 = round_crc32_4 (tmp[2], crc32);
crc32 = round_crc32_4 (tmp[3], crc32);
}
if (crc32 != crc32_save) return 0;
return 1;
}
// 1024 bit
DECLSPEC int verify_header_kuznyechik_aes (__global const tc_t *esalt_bufs, const u32 *ukey1, const u32 *ukey2, const u32 *ukey3, const u32 *ukey4, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
{
u32 ks_kuznyechik[40];
u32 ks_aes[60];
u32 S[4] = { 0 };
u32 T_kuznyechik[4] = { 0 };
u32 T_aes[4] = { 0 };
u32 data[4];
data[0] = esalt_bufs[0].data_buf[0];
data[1] = esalt_bufs[0].data_buf[1];
data[2] = esalt_bufs[0].data_buf[2];
data[3] = esalt_bufs[0].data_buf[3];
u32 tmp[4];
kuznyechik_decrypt_xts_first (ukey2, ukey4, data, tmp, S, T_kuznyechik, ks_kuznyechik);
aes256_decrypt_xts_first (ukey1, ukey3, tmp, tmp, S, T_aes, ks_aes, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4);
const u32 signature = esalt_bufs[0].signature;
if (tmp[0] != signature) return 0;
const u32 crc32_save = swap32_S (~tmp[2]);
// seek to byte 256
for (volatile int i = 4; i < 64 - 16; i += 4)
{
xts_mul2 (T_kuznyechik, T_kuznyechik);
xts_mul2 (T_aes, T_aes);
}
// calculate crc32 from here
u32 crc32 = ~0;
for (volatile int i = 64 - 16; i < 128 - 16; i += 4)
{
data[0] = esalt_bufs[0].data_buf[i + 0];
data[1] = esalt_bufs[0].data_buf[i + 1];
data[2] = esalt_bufs[0].data_buf[i + 2];
data[3] = esalt_bufs[0].data_buf[i + 3];
kuznyechik_decrypt_xts_next (data, tmp, T_kuznyechik, ks_kuznyechik);
aes256_decrypt_xts_next (tmp, tmp, T_aes, ks_aes, s_td0, s_td1, s_td2, s_td3, s_td4);
crc32 = round_crc32_4 (tmp[0], crc32);
crc32 = round_crc32_4 (tmp[1], crc32);
crc32 = round_crc32_4 (tmp[2], crc32);
crc32 = round_crc32_4 (tmp[3], crc32);
}
if (crc32 != crc32_save) return 0;
return 1;
}
DECLSPEC int verify_header_kuznyechik_twofish (__global const tc_t *esalt_bufs, const u32 *ukey1, const u32 *ukey2, const u32 *ukey3, const u32 *ukey4)
{
u32 ks_kuznyechik[40];
u32 sk_twofish[4];
u32 lk_twofish[40];
u32 S[4] = { 0 };
u32 T_kuznyechik[4] = { 0 };
u32 T_twofish[4] = { 0 };
u32 data[4];
data[0] = esalt_bufs[0].data_buf[0];
data[1] = esalt_bufs[0].data_buf[1];
data[2] = esalt_bufs[0].data_buf[2];
data[3] = esalt_bufs[0].data_buf[3];
u32 tmp[4];
kuznyechik_decrypt_xts_first (ukey2, ukey4, data, tmp, S, T_kuznyechik, ks_kuznyechik);
twofish256_decrypt_xts_first (ukey1, ukey3, tmp, tmp, S, T_twofish, sk_twofish, lk_twofish);
const u32 signature = esalt_bufs[0].signature;
if (tmp[0] != signature) return 0;
const u32 crc32_save = swap32_S (~tmp[2]);
// seek to byte 256
for (volatile int i = 4; i < 64 - 16; i += 4)
{
xts_mul2 (T_kuznyechik, T_kuznyechik);
xts_mul2 (T_twofish, T_twofish);
}
// calculate crc32 from here
u32 crc32 = ~0;
for (volatile int i = 64 - 16; i < 128 - 16; i += 4)
{
data[0] = esalt_bufs[0].data_buf[i + 0];
data[1] = esalt_bufs[0].data_buf[i + 1];
data[2] = esalt_bufs[0].data_buf[i + 2];
data[3] = esalt_bufs[0].data_buf[i + 3];
kuznyechik_decrypt_xts_next (data, tmp, T_kuznyechik, ks_kuznyechik);
twofish256_decrypt_xts_next (tmp, tmp, T_twofish, sk_twofish, lk_twofish);
crc32 = round_crc32_4 (tmp[0], crc32);
crc32 = round_crc32_4 (tmp[1], crc32);
crc32 = round_crc32_4 (tmp[2], crc32);
crc32 = round_crc32_4 (tmp[3], crc32);
}
if (crc32 != crc32_save) return 0;
return 1;
}

View File

@ -16,11 +16,13 @@
#include "inc_cipher_aes.cl"
#include "inc_cipher_twofish.cl"
#include "inc_cipher_serpent.cl"
#include "inc_cipher_kuznyechik.cl"
#include "inc_truecrypt_keyfile.cl"
#include "inc_truecrypt_keyboard.cl"
#include "inc_truecrypt_crc32.cl"
#include "inc_truecrypt_xts.cl"
#include "inc_veracrypt_xts.cl"
DECLSPEC void hmac_sha512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *w4, u32x *w5, u32x *w6, u32x *w7, u64x *ipad, u64x *opad, u64x *digest)
{
@ -521,4 +523,12 @@ __kernel void m06221_comp (KERN_ATTR_TMPS_ESALT (tc64_tmp_t, tc_t))
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
if (verify_header_kuznyechik (esalt_bufs, ukey1, ukey2) == 1)
{
if (atomic_inc (&hashes_shown[0]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
}

View File

@ -16,11 +16,13 @@
#include "inc_cipher_aes.cl"
#include "inc_cipher_twofish.cl"
#include "inc_cipher_serpent.cl"
#include "inc_cipher_kuznyechik.cl"
#include "inc_truecrypt_keyfile.cl"
#include "inc_truecrypt_keyboard.cl"
#include "inc_truecrypt_crc32.cl"
#include "inc_truecrypt_xts.cl"
#include "inc_veracrypt_xts.cl"
DECLSPEC void hmac_sha512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *w4, u32x *w5, u32x *w6, u32x *w7, u64x *ipad, u64x *opad, u64x *digest)
{
@ -522,6 +524,14 @@ __kernel void m06222_comp (KERN_ATTR_TMPS_ESALT (tc64_tmp_t, tc_t))
}
}
if (verify_header_kuznyechik (esalt_bufs, ukey1, ukey2) == 1)
{
if (atomic_inc (&hashes_shown[0]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
u32 ukey3[8];
ukey3[0] = swap32_S (h32_from_64_S (tmps[gid].out[ 8]));
@ -567,4 +577,20 @@ __kernel void m06222_comp (KERN_ATTR_TMPS_ESALT (tc64_tmp_t, tc_t))
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
if (verify_header_kuznyechik_aes (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
{
if (atomic_inc (&hashes_shown[0]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
if (verify_header_kuznyechik_twofish (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
{
if (atomic_inc (&hashes_shown[0]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
}

View File

@ -16,11 +16,13 @@
#include "inc_cipher_aes.cl"
#include "inc_cipher_twofish.cl"
#include "inc_cipher_serpent.cl"
#include "inc_cipher_kuznyechik.cl"
#include "inc_truecrypt_keyfile.cl"
#include "inc_truecrypt_keyboard.cl"
#include "inc_truecrypt_crc32.cl"
#include "inc_truecrypt_xts.cl"
#include "inc_veracrypt_xts.cl"
DECLSPEC void hmac_sha512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *w4, u32x *w5, u32x *w6, u32x *w7, u64x *ipad, u64x *opad, u64x *digest)
{
@ -522,6 +524,14 @@ __kernel void m06223_comp (KERN_ATTR_TMPS_ESALT (tc64_tmp_t, tc_t))
}
}
if (verify_header_kuznyechik (esalt_bufs, ukey1, ukey2) == 1)
{
if (atomic_inc (&hashes_shown[0]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
u32 ukey3[8];
ukey3[0] = swap32_S (h32_from_64_S (tmps[gid].out[ 8]));
@ -568,6 +578,22 @@ __kernel void m06223_comp (KERN_ATTR_TMPS_ESALT (tc64_tmp_t, tc_t))
}
}
if (verify_header_kuznyechik_aes (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
{
if (atomic_inc (&hashes_shown[0]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
if (verify_header_kuznyechik_twofish (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
{
if (atomic_inc (&hashes_shown[0]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
u32 ukey5[8];
ukey5[0] = swap32_S (h32_from_64_S (tmps[gid].out[16]));

View File

@ -16,11 +16,13 @@
#include "inc_cipher_aes.cl"
#include "inc_cipher_twofish.cl"
#include "inc_cipher_serpent.cl"
#include "inc_cipher_kuznyechik.cl"
#include "inc_truecrypt_keyfile.cl"
#include "inc_truecrypt_keyboard.cl"
#include "inc_truecrypt_crc32.cl"
#include "inc_truecrypt_xts.cl"
#include "inc_veracrypt_xts.cl"
DECLSPEC void hmac_whirlpool_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest, __local u32 (*s_Ch)[256], __local u32 (*s_Cl)[256])
{
@ -631,4 +633,12 @@ __kernel void m06231_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
if (verify_header_kuznyechik (esalt_bufs, ukey1, ukey2) == 1)
{
if (atomic_inc (&hashes_shown[0]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
}

View File

@ -16,11 +16,13 @@
#include "inc_cipher_aes.cl"
#include "inc_cipher_twofish.cl"
#include "inc_cipher_serpent.cl"
#include "inc_cipher_kuznyechik.cl"
#include "inc_truecrypt_keyfile.cl"
#include "inc_truecrypt_keyboard.cl"
#include "inc_truecrypt_crc32.cl"
#include "inc_truecrypt_xts.cl"
#include "inc_veracrypt_xts.cl"
DECLSPEC void hmac_whirlpool_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest, __local u32 (*s_Ch)[256], __local u32 (*s_Cl)[256])
{
@ -632,6 +634,14 @@ __kernel void m06232_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
}
}
if (verify_header_kuznyechik (esalt_bufs, ukey1, ukey2) == 1)
{
if (atomic_inc (&hashes_shown[0]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
u32 ukey3[8];
ukey3[0] = swap32_S (tmps[gid].out[16]);
@ -677,4 +687,20 @@ __kernel void m06232_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
if (verify_header_kuznyechik_aes (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
{
if (atomic_inc (&hashes_shown[0]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
if (verify_header_kuznyechik_twofish (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
{
if (atomic_inc (&hashes_shown[0]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
}

View File

@ -16,11 +16,13 @@
#include "inc_cipher_aes.cl"
#include "inc_cipher_twofish.cl"
#include "inc_cipher_serpent.cl"
#include "inc_cipher_kuznyechik.cl"
#include "inc_truecrypt_keyfile.cl"
#include "inc_truecrypt_keyboard.cl"
#include "inc_truecrypt_crc32.cl"
#include "inc_truecrypt_xts.cl"
#include "inc_veracrypt_xts.cl"
DECLSPEC void hmac_whirlpool_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest, __local u32 (*s_Ch)[256], __local u32 (*s_Cl)[256])
{
@ -678,6 +680,14 @@ __kernel void m06233_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
}
}
if (verify_header_kuznyechik (esalt_bufs, ukey1, ukey2) == 1)
{
if (atomic_inc (&hashes_shown[0]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
u32 ukey5[8];
ukey5[0] = swap32_S (tmps[gid].out[32]);
@ -715,4 +725,20 @@ __kernel void m06233_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
if (verify_header_kuznyechik_aes (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
{
if (atomic_inc (&hashes_shown[0]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
if (verify_header_kuznyechik_twofish (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
{
if (atomic_inc (&hashes_shown[0]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
}

View File

@ -16,11 +16,13 @@
#include "inc_cipher_aes.cl"
#include "inc_cipher_twofish.cl"
#include "inc_cipher_serpent.cl"
#include "inc_cipher_kuznyechik.cl"
#include "inc_truecrypt_keyfile.cl"
#include "inc_truecrypt_keyboard.cl"
#include "inc_truecrypt_crc32.cl"
#include "inc_truecrypt_xts.cl"
#include "inc_veracrypt_xts.cl"
DECLSPEC void hmac_sha256_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest)
{
@ -429,4 +431,12 @@ __kernel void m13751_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
if (verify_header_kuznyechik (esalt_bufs, ukey1, ukey2) == 1)
{
if (atomic_inc (&hashes_shown[0]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
}

View File

@ -16,11 +16,13 @@
#include "inc_cipher_aes.cl"
#include "inc_cipher_twofish.cl"
#include "inc_cipher_serpent.cl"
#include "inc_cipher_kuznyechik.cl"
#include "inc_truecrypt_keyfile.cl"
#include "inc_truecrypt_keyboard.cl"
#include "inc_truecrypt_crc32.cl"
#include "inc_truecrypt_xts.cl"
#include "inc_veracrypt_xts.cl"
DECLSPEC void hmac_sha256_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest)
{
@ -430,6 +432,14 @@ __kernel void m13752_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
}
}
if (verify_header_kuznyechik (esalt_bufs, ukey1, ukey2) == 1)
{
if (atomic_inc (&hashes_shown[0]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
u32 ukey3[8];
ukey3[0] = swap32_S (tmps[gid].out[16]);
@ -475,4 +485,20 @@ __kernel void m13752_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
if (verify_header_kuznyechik_aes (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
{
if (atomic_inc (&hashes_shown[0]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
if (verify_header_kuznyechik_twofish (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
{
if (atomic_inc (&hashes_shown[0]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
}

View File

@ -16,11 +16,13 @@
#include "inc_cipher_aes.cl"
#include "inc_cipher_twofish.cl"
#include "inc_cipher_serpent.cl"
#include "inc_cipher_kuznyechik.cl"
#include "inc_truecrypt_keyfile.cl"
#include "inc_truecrypt_keyboard.cl"
#include "inc_truecrypt_crc32.cl"
#include "inc_truecrypt_xts.cl"
#include "inc_veracrypt_xts.cl"
DECLSPEC void hmac_sha256_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest)
{
@ -430,6 +432,14 @@ __kernel void m13753_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
}
}
if (verify_header_kuznyechik (esalt_bufs, ukey1, ukey2) == 1)
{
if (atomic_inc (&hashes_shown[0]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
u32 ukey3[8];
ukey3[0] = swap32_S (tmps[gid].out[16]);
@ -476,6 +486,22 @@ __kernel void m13753_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
}
}
if (verify_header_kuznyechik_aes (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
{
if (atomic_inc (&hashes_shown[0]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
if (verify_header_kuznyechik_twofish (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
{
if (atomic_inc (&hashes_shown[0]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
u32 ukey5[8];
ukey5[0] = swap32_S (tmps[gid].out[32]);

View File

@ -16,11 +16,13 @@
#include "inc_cipher_aes.cl"
#include "inc_cipher_twofish.cl"
#include "inc_cipher_serpent.cl"
#include "inc_cipher_kuznyechik.cl"
#include "inc_truecrypt_keyfile.cl"
#include "inc_truecrypt_keyboard.cl"
#include "inc_truecrypt_crc32.cl"
#include "inc_truecrypt_xts.cl"
#include "inc_veracrypt_xts.cl"
DECLSPEC void hmac_streebog512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u64x *ipad_hash, u64x *opad_hash, u64x *ipad_raw, u64x *opad_raw, u64x *digest, SHM_TYPE u64a (*s_sbob_sl64)[256])
{
@ -558,4 +560,12 @@ __kernel void m13771_comp (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, tc_t))
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
if (verify_header_kuznyechik (esalt_bufs, ukey1, ukey2) == 1)
{
if (atomic_inc (&hashes_shown[0]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
}

View File

@ -16,11 +16,13 @@
#include "inc_cipher_aes.cl"
#include "inc_cipher_twofish.cl"
#include "inc_cipher_serpent.cl"
#include "inc_cipher_kuznyechik.cl"
#include "inc_truecrypt_keyfile.cl"
#include "inc_truecrypt_keyboard.cl"
#include "inc_truecrypt_crc32.cl"
#include "inc_truecrypt_xts.cl"
#include "inc_veracrypt_xts.cl"
DECLSPEC void hmac_streebog512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u64x *ipad_hash, u64x *opad_hash, u64x *ipad_raw, u64x *opad_raw, u64x *digest, SHM_TYPE u64a (*s_sbob_sl64)[256])
{
@ -560,6 +562,14 @@ __kernel void m13772_comp (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, tc_t))
}
}
if (verify_header_kuznyechik (esalt_bufs, ukey1, ukey2) == 1)
{
if (atomic_inc (&hashes_shown[0]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
u32 ukey3[8];
ukey3[0] = swap32_S (h32_from_64_S (tmps[gid].out[15]));
@ -605,4 +615,20 @@ __kernel void m13772_comp (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, tc_t))
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
if (verify_header_kuznyechik_aes (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
{
if (atomic_inc (&hashes_shown[0]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
if (verify_header_kuznyechik_twofish (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
{
if (atomic_inc (&hashes_shown[0]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
}

View File

@ -16,11 +16,13 @@
#include "inc_cipher_aes.cl"
#include "inc_cipher_twofish.cl"
#include "inc_cipher_serpent.cl"
#include "inc_cipher_kuznyechik.cl"
#include "inc_truecrypt_keyfile.cl"
#include "inc_truecrypt_keyboard.cl"
#include "inc_truecrypt_crc32.cl"
#include "inc_truecrypt_xts.cl"
#include "inc_veracrypt_xts.cl"
DECLSPEC void hmac_streebog512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u64x *ipad_hash, u64x *opad_hash, u64x *ipad_raw, u64x *opad_raw, u64x *digest, SHM_TYPE u64a (*s_sbob_sl64)[256])
{
@ -560,6 +562,14 @@ __kernel void m13773_comp (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, tc_t))
}
}
if (verify_header_kuznyechik (esalt_bufs, ukey1, ukey2) == 1)
{
if (atomic_inc (&hashes_shown[0]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
u32 ukey3[8];
ukey3[0] = swap32_S (h32_from_64_S (tmps[gid].out[15]));
@ -643,4 +653,20 @@ __kernel void m13773_comp (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, tc_t))
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
if (verify_header_kuznyechik_aes (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
{
if (atomic_inc (&hashes_shown[0]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
if (verify_header_kuznyechik_twofish (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
{
if (atomic_inc (&hashes_shown[0]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
}

View File

@ -25,6 +25,7 @@
- Added hash-mode 13773 (VeraCrypt PBKDF2-HMAC-Streebog-512 + XTS 1536 bit)
- Added hash-mode 18200 (Kerberos 5 AS-REP etype 23)
- Added hash-mode 18300 (Apple File System (APFS))
- Added Kuznyechik cipher and cascades support for VeraCrypt kernels
##
## Improvements
@ -91,6 +92,7 @@
- Tests: Added hash-mode 13771 (VeraCrypt PBKDF2-HMAC-Streebog-512 + XTS 512 bit)
- Tests: Added hash-mode 13772 (VeraCrypt PBKDF2-HMAC-Streebog-512 + XTS 1024 bit)
- Tests: Added hash-mode 13773 (VeraCrypt PBKDF2-HMAC-Streebog-512 + XTS 1536 bit)
- Tests: Added VeraCrypt containers for Kuznyechik cipher and cascades
* changes v4.2.1 -> v5.0.0

View File

@ -360,10 +360,14 @@ static const char *const USAGE_BIG[] =
" Y | 1 = XTS 512 bit pure AES | Full-Disk Encryption (FDE)",
" Y | 1 = XTS 512 bit pure Serpent | Full-Disk Encryption (FDE)",
" Y | 1 = XTS 512 bit pure Twofish | Full-Disk Encryption (FDE)",
" Y | 1 = XTS 512 bit pure Kuznyechik | Full-Disk Encryption (FDE)",
" Y | 2 = XTS 1024 bit pure AES | Full-Disk Encryption (FDE)",
" Y | 2 = XTS 1024 bit pure Serpent | Full-Disk Encryption (FDE)",
" Y | 2 = XTS 1024 bit pure Twofish | Full-Disk Encryption (FDE)",
" Y | 2 = XTS 1024 bit pure Kuznyechik | Full-Disk Encryption (FDE)",
" Y | 2 = XTS 1024 bit cascaded AES-Twofish | Full-Disk Encryption (FDE)",
" Y | 2 = XTS 1024 bit cascaded Kuznyechik-AES | Full-Disk Encryption (FDE)",
" Y | 2 = XTS 1024 bit cascaded Kuznyechik-Twofish | Full-Disk Encryption (FDE)",
" Y | 2 = XTS 1024 bit cascaded Serpent-AES | Full-Disk Encryption (FDE)",
" Y | 2 = XTS 1024 bit cascaded Twofish-Serpent | Full-Disk Encryption (FDE)",
" Y | 3 = XTS 1536 bit all | Full-Disk Encryption (FDE)",

1
tools/tc_tests/.gitattributes vendored Normal file
View File

@ -0,0 +1 @@
*.tc binary

View File

@ -2253,10 +2253,10 @@ function truecrypt_test()
# Compose and execute hashcat command on a VeraCrypt test container
# Must not be called for hash types other than 137XY
# $1: primary (first layer) cipher id, must be 0-2
# $1: cipher variation, can be 0-4
function veracrypt_test()
{
primary_cipher=$1
cipher_variation=$1
hash_function=""
@ -2274,25 +2274,32 @@ function veracrypt_test()
cipher_digit="${hash_type:4:1}"
case $cipher_digit in
1)
[ $primary_cipher -eq "0" ] && cipher_cascade="aes"
[ $primary_cipher -eq "1" ] && cipher_cascade="serpent"
[ $primary_cipher -eq "2" ] && cipher_cascade="twofish"
[ $cipher_variation -eq "0" ] && cipher_cascade="aes"
[ $cipher_variation -eq "1" ] && cipher_cascade="serpent"
[ $cipher_variation -eq "2" ] && cipher_cascade="twofish"
[ $cipher_variation -eq "3" ] && cipher_cascade="kuznyechik"
;;
2)
[ $primary_cipher -eq "0" ] && cipher_cascade="aes-twofish"
[ $primary_cipher -eq "1" ] && cipher_cascade="serpent-aes"
[ $primary_cipher -eq "2" ] && cipher_cascade="twofish-serpent"
[ $cipher_variation -eq "0" ] && cipher_cascade="aes-twofish"
[ $cipher_variation -eq "1" ] && cipher_cascade="serpent-aes"
[ $cipher_variation -eq "2" ] && cipher_cascade="twofish-serpent"
[ $cipher_variation -eq "3" ] && cipher_cascade="kuznyechik-aes"
[ $cipher_variation -eq "4" ] && cipher_cascade="kuznyechik-twofish"
;;
3)
[ $primary_cipher -eq "0" ] && cipher_cascade="aes-twofish-serpent"
[ $primary_cipher -eq "1" ] && cipher_cascade="serpent-twofish-aes"
[ $primary_cipher -eq "2" ] && cipher_cascade=""
[ $cipher_variation -eq "0" ] && cipher_cascade="aes-twofish-serpent"
[ $cipher_variation -eq "1" ] && cipher_cascade="serpent-twofish-aes"
;;
esac
[ -n "$cipher_cascade" ] || return
CMD="./${BIN} ${OPTS} -a 3 -m ${hash_type} ${TDIR}/vc_tests/hashcat_${hash_function}_${cipher_cascade}.vc hashca?l"
filename="${TDIR}/vc_tests/hashcat_${hash_function}_${cipher_cascade}.vc"
# The hash-cipher combination might be invalid (e.g. RIPEMD-160 + Kuznyechik)
[ -f "${filename}" ] || return
CMD="./${BIN} ${OPTS} -a 3 -m ${hash_type} ${filename} hashca?l"
echo "> Testing hash type ${hash_type} with attack mode 3, markov ${MARKOV}, single hash, Device-Type ${TYPE}, vector-width ${VECTOR}, cipher ${cipher_cascade}" &>> ${OUTD}/logfull.txt
@ -2767,9 +2774,11 @@ if [ "${PACKAGE}" -eq 0 -o -z "${PACKAGE_FOLDER}" ]; then
# Look up if this is one of supported VeraCrypt modes
if is_in_array ${hash_type} ${VC_MODES}; then
veracrypt_test 0
veracrypt_test 1
veracrypt_test 2
veracrypt_test 0 # aes
veracrypt_test 1 # serpent
veracrypt_test 2 # twofish
veracrypt_test 3 # kuznyechik
veracrypt_test 4 # kuznyechik (2nd cascade)
elif [[ ${hash_type} -ge 6211 ]] && [[ ${hash_type} -le 6243 ]]; then
# run truecrypt tests

1
tools/vc_tests/.gitattributes vendored Normal file
View File

@ -0,0 +1 @@
*.vc binary

Binary file not shown.

Binary file not shown.

Binary file not shown.

Binary file not shown.

Binary file not shown.

Binary file not shown.

Binary file not shown.

Binary file not shown.

Binary file not shown.

Binary file not shown.

Binary file not shown.

Binary file not shown.