1
0
mirror of https://github.com/hashcat/hashcat.git synced 2024-11-21 23:58:07 +00:00

Add Camellia support for VeraCrypt kernels

Adds suport for the Japanese cipher Camellia with 256-bit keys as used
by VeraCrypt.

 - Add Camellia header decryption checks to all VeraCrypt kernels
 - Add test containers for remaining cipher combinations
This commit is contained in:
R. Yushaev 2018-11-28 14:21:14 +01:00
parent b457f402c6
commit baf47d409e
36 changed files with 1093 additions and 27 deletions

View File

@ -0,0 +1,365 @@
/* *
* This is an OpenCL implementation of the encryption algorithm: *
* *
* Camellia by Kazumaro Aoki, Masayuki Kanda, Shiho Moriai, *
* Tetsuya Ichikawa, Mitsuru Matsui, *
* Junko Nakajima and Toshio Tokita *
* *
* http://info.isl.ntt.co.jp/crypt/eng/camellia/technology.html *
* *
* Copyright of the ANSI-C implementation: *
* *
* Mitsubishi Electric Corp 2000-2001 *
* *
* Adapted for GPU use with hashcat by Ruslan Yushaev. *
* *
*/
__constant const u32a c_sbox[256] =
{
0x70, 0x82, 0x2c, 0xec, 0xb3, 0x27, 0xc0, 0xe5,
0xe4, 0x85, 0x57, 0x35, 0xea, 0x0c, 0xae, 0x41,
0x23, 0xef, 0x6b, 0x93, 0x45, 0x19, 0xa5, 0x21,
0xed, 0x0e, 0x4f, 0x4e, 0x1d, 0x65, 0x92, 0xbd,
0x86, 0xb8, 0xaf, 0x8f, 0x7c, 0xeb, 0x1f, 0xce,
0x3e, 0x30, 0xdc, 0x5f, 0x5e, 0xc5, 0x0b, 0x1a,
0xa6, 0xe1, 0x39, 0xca, 0xd5, 0x47, 0x5d, 0x3d,
0xd9, 0x01, 0x5a, 0xd6, 0x51, 0x56, 0x6c, 0x4d,
0x8b, 0x0d, 0x9a, 0x66, 0xfb, 0xcc, 0xb0, 0x2d,
0x74, 0x12, 0x2b, 0x20, 0xf0, 0xb1, 0x84, 0x99,
0xdf, 0x4c, 0xcb, 0xc2, 0x34, 0x7e, 0x76, 0x05,
0x6d, 0xb7, 0xa9, 0x31, 0xd1, 0x17, 0x04, 0xd7,
0x14, 0x58, 0x3a, 0x61, 0xde, 0x1b, 0x11, 0x1c,
0x32, 0x0f, 0x9c, 0x16, 0x53, 0x18, 0xf2, 0x22,
0xfe, 0x44, 0xcf, 0xb2, 0xc3, 0xb5, 0x7a, 0x91,
0x24, 0x08, 0xe8, 0xa8, 0x60, 0xfc, 0x69, 0x50,
0xaa, 0xd0, 0xa0, 0x7d, 0xa1, 0x89, 0x62, 0x97,
0x54, 0x5b, 0x1e, 0x95, 0xe0, 0xff, 0x64, 0xd2,
0x10, 0xc4, 0x00, 0x48, 0xa3, 0xf7, 0x75, 0xdb,
0x8a, 0x03, 0xe6, 0xda, 0x09, 0x3f, 0xdd, 0x94,
0x87, 0x5c, 0x83, 0x02, 0xcd, 0x4a, 0x90, 0x33,
0x73, 0x67, 0xf6, 0xf3, 0x9d, 0x7f, 0xbf, 0xe2,
0x52, 0x9b, 0xd8, 0x26, 0xc8, 0x37, 0xc6, 0x3b,
0x81, 0x96, 0x6f, 0x4b, 0x13, 0xbe, 0x63, 0x2e,
0xe9, 0x79, 0xa7, 0x8c, 0x9f, 0x6e, 0xbc, 0x8e,
0x29, 0xf5, 0xf9, 0xb6, 0x2f, 0xfd, 0xb4, 0x59,
0x78, 0x98, 0x06, 0x6a, 0xe7, 0x46, 0x71, 0xba,
0xd4, 0x25, 0xab, 0x42, 0x88, 0xa2, 0x8d, 0xfa,
0x72, 0x07, 0xb9, 0x55, 0xf8, 0xee, 0xac, 0x0a,
0x36, 0x49, 0x2a, 0x68, 0x3c, 0x38, 0xf1, 0xa4,
0x40, 0x28, 0xd3, 0x7b, 0xbb, 0xc9, 0x43, 0xc1,
0x15, 0xe3, 0xad, 0xf4, 0x77, 0xc7, 0x80, 0x9e
};
#define c_sbox1(n) c_sbox[(n)]
#define c_sbox2(n) (((c_sbox[(n)] >> 7) ^ (c_sbox[(n)] << 1)) & 0xff)
#define c_sbox3(n) (((c_sbox[(n)] >> 1) ^ (c_sbox[(n)] << 7)) & 0xff)
#define c_sbox4(n) c_sbox[(((n) << 1) ^ ((n) >> 7)) & 0xff]
#define cam_rotate(a,b,n) swap32_S ((u[(a)] << (n)) ^ (u[(b)] >> (32 - (n))))
#define extract_byte(x,n) (((x) >> (8 * (n))) & 0xff)
DECLSPEC void cam_feistel (const u32 *x, const u32 *k, u32 *y)
{
u32 b[8];
b[0] = c_sbox1 (extract_byte (x[0], 0) ^ extract_byte (k[0], 0));
b[1] = c_sbox2 (extract_byte (x[0], 1) ^ extract_byte (k[0], 1));
b[2] = c_sbox3 (extract_byte (x[0], 2) ^ extract_byte (k[0], 2));
b[3] = c_sbox4 (extract_byte (x[0], 3) ^ extract_byte (k[0], 3));
b[4] = c_sbox2 (extract_byte (x[1], 0) ^ extract_byte (k[1], 0));
b[5] = c_sbox3 (extract_byte (x[1], 1) ^ extract_byte (k[1], 1));
b[6] = c_sbox4 (extract_byte (x[1], 2) ^ extract_byte (k[1], 2));
b[7] = c_sbox1 (extract_byte (x[1], 3) ^ extract_byte (k[1], 3));
u32 tmp[2];
tmp[0] = (b[0] ^ b[2] ^ b[3] ^ b[5] ^ b[6] ^ b[7]) << 0
| (b[0] ^ b[1] ^ b[3] ^ b[4] ^ b[6] ^ b[7]) << 8
| (b[0] ^ b[1] ^ b[2] ^ b[4] ^ b[5] ^ b[7]) << 16
| (b[1] ^ b[2] ^ b[3] ^ b[4] ^ b[5] ^ b[6]) << 24;
tmp[1] = (b[0] ^ b[1] ^ b[5] ^ b[6] ^ b[7]) << 0
| (b[1] ^ b[2] ^ b[4] ^ b[6] ^ b[7]) << 8
| (b[2] ^ b[3] ^ b[4] ^ b[5] ^ b[7]) << 16
| (b[0] ^ b[3] ^ b[4] ^ b[5] ^ b[6]) << 24;
y[0] ^= tmp[0];
y[1] ^= tmp[1];
}
DECLSPEC void cam_fl (u32 *x, const u32 *kl, const u32 *kr)
{
u32 t[4];
u32 u[4];
u32 v[4];
t[0] = swap32_S (x[0]);
t[1] = swap32_S (x[1]);
t[2] = swap32_S (x[2]);
t[3] = swap32_S (x[3]);
u[0] = swap32_S (kl[0]);
u[1] = swap32_S (kl[1]);
u[2] = swap32_S (kl[2]);
u[3] = swap32_S (kl[3]);
v[0] = swap32_S (kr[0]);
v[1] = swap32_S (kr[1]);
v[2] = swap32_S (kr[2]);
v[3] = swap32_S (kr[3]);
t[1] ^= (t[0] & u[0]) << 1;
t[1] ^= (t[0] & u[0]) >> 31;
t[0] ^= t[1] | u[1];
t[2] ^= t[3] | v[1];
t[3] ^= (t[2] & v[0]) << 1;
t[3] ^= (t[2] & v[0]) >> 31;
x[0] = swap32_S (t[0]);
x[1] = swap32_S (t[1]);
x[2] = swap32_S (t[2]);
x[3] = swap32_S (t[3]);
}
DECLSPEC void camellia256_set_key (u32 *ks, const u32 *ukey)
{
const u32 sigma[12] =
{
0x7f669ea0, 0x8b90cc3b, 0x58e87ab6, 0xb273aa4c,
0x2f37efc6, 0xbe824fe9, 0xa553ff54, 0x1c6fd3f1,
0xfa27e510, 0x1d2d68de, 0xc28856b0, 0xfdc1e6b3
};
u32 tmp[8];
tmp[0] = ukey[0] ^ ukey[4];
tmp[1] = ukey[1] ^ ukey[5];
tmp[2] = ukey[2] ^ ukey[6];
tmp[3] = ukey[3] ^ ukey[7];
cam_feistel (&tmp[0], &sigma[0], &tmp[2]);
cam_feistel (&tmp[2], &sigma[2], &tmp[0]);
tmp[0] ^= ukey[0];
tmp[1] ^= ukey[1];
tmp[2] ^= ukey[2];
tmp[3] ^= ukey[3];
cam_feistel (&tmp[0], &sigma[4], &tmp[2]);
cam_feistel (&tmp[2], &sigma[6], &tmp[0]);
tmp[4] = tmp[0] ^ ukey[4];
tmp[5] = tmp[1] ^ ukey[5];
tmp[6] = tmp[2] ^ ukey[6];
tmp[7] = tmp[3] ^ ukey[7];
cam_feistel (&tmp[4], &sigma[8], &tmp[6]);
cam_feistel (&tmp[6], &sigma[10], &tmp[4]);
// used in cam_rotate macro
u32 u[16];
u[0] = swap32_S (ukey[0]);
u[1] = swap32_S (ukey[1]);
u[2] = swap32_S (ukey[2]);
u[3] = swap32_S (ukey[3]);
u[4] = swap32_S (tmp[0]);
u[5] = swap32_S (tmp[1]);
u[6] = swap32_S (tmp[2]);
u[7] = swap32_S (tmp[3]);
u[8] = swap32_S (ukey[4]);
u[9] = swap32_S (ukey[5]);
u[10] = swap32_S (ukey[6]);
u[11] = swap32_S (ukey[7]);
u[12] = swap32_S (tmp[4]);
u[13] = swap32_S (tmp[5]);
u[14] = swap32_S (tmp[6]);
u[15] = swap32_S (tmp[7]);
ks[0] = swap32_S (u[0]);
ks[1] = swap32_S (u[1]);
ks[2] = swap32_S (u[2]);
ks[3] = swap32_S (u[3]);
ks[4] = swap32_S (u[12]);
ks[5] = swap32_S (u[13]);
ks[6] = swap32_S (u[14]);
ks[7] = swap32_S (u[15]);
ks[8] = cam_rotate (8, 9, 15);
ks[9] = cam_rotate (9, 10, 15);
ks[10] = cam_rotate (10, 11, 15);
ks[11] = cam_rotate (11, 8, 15);
ks[12] = cam_rotate (4, 5, 15);
ks[13] = cam_rotate (5, 6, 15);
ks[14] = cam_rotate (6, 7, 15);
ks[15] = cam_rotate (7, 4, 15);
ks[16] = cam_rotate (8, 9, 30);
ks[17] = cam_rotate (9, 10, 30);
ks[18] = cam_rotate (10, 11, 30);
ks[19] = cam_rotate (11, 8, 30);
ks[20] = cam_rotate (12, 13, 30);
ks[21] = cam_rotate (13, 14, 30);
ks[22] = cam_rotate (14, 15, 30);
ks[23] = cam_rotate (15, 12, 30);
ks[24] = cam_rotate (1, 2, 13);
ks[25] = cam_rotate (2, 3, 13);
ks[26] = cam_rotate (3, 0, 13);
ks[27] = cam_rotate (0, 1, 13);
ks[28] = cam_rotate (5, 6, 13);
ks[29] = cam_rotate (6, 7, 13);
ks[30] = cam_rotate (7, 4, 13);
ks[31] = cam_rotate (4, 5, 13);
ks[32] = cam_rotate (1, 2, 28);
ks[33] = cam_rotate (2, 3, 28);
ks[34] = cam_rotate (3, 0, 28);
ks[35] = cam_rotate (0, 1, 28);
ks[36] = cam_rotate (9, 10, 28);
ks[37] = cam_rotate (10, 11, 28);
ks[38] = cam_rotate (11, 8, 28);
ks[39] = cam_rotate (8, 9, 28);
ks[40] = cam_rotate (13, 14, 28);
ks[41] = cam_rotate (14, 15, 28);
ks[42] = cam_rotate (15, 12, 28);
ks[43] = cam_rotate (12, 13, 28);
ks[44] = cam_rotate (2, 3, 13);
ks[45] = cam_rotate (3, 0, 13);
ks[46] = cam_rotate (0, 1, 13);
ks[47] = cam_rotate (1, 2, 13);
ks[48] = cam_rotate (6, 7, 13);
ks[49] = cam_rotate (7, 4, 13);
ks[50] = cam_rotate (4, 5, 13);
ks[51] = cam_rotate (5, 6, 13);
ks[52] = cam_rotate (10, 11, 30);
ks[53] = cam_rotate (11, 8, 30);
ks[54] = cam_rotate (8, 9, 30);
ks[55] = cam_rotate (9, 10, 30);
ks[56] = cam_rotate (6, 7, 30);
ks[57] = cam_rotate (7, 4, 30);
ks[58] = cam_rotate (4, 5, 30);
ks[59] = cam_rotate (5, 6, 30);
ks[60] = cam_rotate (3, 0, 15);
ks[61] = cam_rotate (0, 1, 15);
ks[62] = cam_rotate (1, 2, 15);
ks[63] = cam_rotate (2, 3, 15);
ks[64] = cam_rotate (15, 12, 15);
ks[65] = cam_rotate (12, 13, 15);
ks[66] = cam_rotate (13, 14, 15);
ks[67] = cam_rotate (14, 15, 15);
}
DECLSPEC void camellia256_encrypt (const u32 *ks, const u32 *in, u32 *out)
{
out[0] = in[0] ^ ks[0];
out[1] = in[1] ^ ks[1];
out[2] = in[2] ^ ks[2];
out[3] = in[3] ^ ks[3];
cam_feistel (&out[0], &ks[4], &out[2]);
cam_feistel (&out[2], &ks[6], &out[0]);
cam_feistel (&out[0], &ks[8], &out[2]);
cam_feistel (&out[2], &ks[10], &out[0]);
cam_feistel (&out[0], &ks[12], &out[2]);
cam_feistel (&out[2], &ks[14], &out[0]);
cam_fl (out, &ks[16], &ks[18]);
cam_feistel (&out[0], &ks[20], &out[2]);
cam_feistel (&out[2], &ks[22], &out[0]);
cam_feistel (&out[0], &ks[24], &out[2]);
cam_feistel (&out[2], &ks[26], &out[0]);
cam_feistel (&out[0], &ks[28], &out[2]);
cam_feistel (&out[2], &ks[30], &out[0]);
cam_fl (out, &ks[32], &ks[34]);
cam_feistel (&out[0], &ks[36], &out[2]);
cam_feistel (&out[2], &ks[38], &out[0]);
cam_feistel (&out[0], &ks[40], &out[2]);
cam_feistel (&out[2], &ks[42], &out[0]);
cam_feistel (&out[0], &ks[44], &out[2]);
cam_feistel (&out[2], &ks[46], &out[0]);
cam_fl (out, &ks[48], &ks[50]);
cam_feistel (&out[0], &ks[52], &out[2]);
cam_feistel (&out[2], &ks[54], &out[0]);
cam_feistel (&out[0], &ks[56], &out[2]);
cam_feistel (&out[2], &ks[58], &out[0]);
cam_feistel (&out[0], &ks[60], &out[2]);
cam_feistel (&out[2], &ks[62], &out[0]);
u32 tmp[2];
tmp[0] = out[0];
tmp[1] = out[1];
out[0] = out[2] ^ ks[64];
out[1] = out[3] ^ ks[65];
out[2] = tmp[0] ^ ks[66];
out[3] = tmp[1] ^ ks[67];
}
DECLSPEC void camellia256_decrypt (const u32 *ks, const u32 *in, u32 *out)
{
out[0] = in[0] ^ ks[64];
out[1] = in[1] ^ ks[65];
out[2] = in[2] ^ ks[66];
out[3] = in[3] ^ ks[67];
cam_feistel (&out[0], &ks[62], &out[2]);
cam_feistel (&out[2], &ks[60], &out[0]);
cam_feistel (&out[0], &ks[58], &out[2]);
cam_feistel (&out[2], &ks[56], &out[0]);
cam_feistel (&out[0], &ks[54], &out[2]);
cam_feistel (&out[2], &ks[52], &out[0]);
cam_fl (out, &ks[50], &ks[48]);
cam_feistel (&out[0], &ks[46], &out[2]);
cam_feistel (&out[2], &ks[44], &out[0]);
cam_feistel (&out[0], &ks[42], &out[2]);
cam_feistel (&out[2], &ks[40], &out[0]);
cam_feistel (&out[0], &ks[38], &out[2]);
cam_feistel (&out[2], &ks[36], &out[0]);
cam_fl (out, &ks[34], &ks[32]);
cam_feistel (&out[0], &ks[30], &out[2]);
cam_feistel (&out[2], &ks[28], &out[0]);
cam_feistel (&out[0], &ks[26], &out[2]);
cam_feistel (&out[2], &ks[24], &out[0]);
cam_feistel (&out[0], &ks[22], &out[2]);
cam_feistel (&out[2], &ks[20], &out[0]);
cam_fl (out, &ks[18], &ks[16]);
cam_feistel (&out[0], &ks[14], &out[2]);
cam_feistel (&out[2], &ks[12], &out[0]);
cam_feistel (&out[0], &ks[10], &out[2]);
cam_feistel (&out[2], &ks[8], &out[0]);
cam_feistel (&out[0], &ks[6], &out[2]);
cam_feistel (&out[2], &ks[4], &out[0]);
u32 tmp[2];
tmp[0] = out[0];
tmp[1] = out[1];
out[0] = out[2] ^ ks[0];
out[1] = out[3] ^ ks[1];
out[2] = tmp[0] ^ ks[2];
out[3] = tmp[1] ^ ks[3];
}

View File

@ -1,3 +1,49 @@
DECLSPEC void camellia256_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];
camellia256_set_key (ks, ukey2);
camellia256_encrypt (ks, S, T);
out[0] ^= T[0];
out[1] ^= T[1];
out[2] ^= T[2];
out[3] ^= T[3];
camellia256_set_key (ks, ukey1);
camellia256_decrypt (ks, out, out);
out[0] ^= T[0];
out[1] ^= T[1];
out[2] ^= T[2];
out[3] ^= T[3];
}
DECLSPEC void camellia256_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];
camellia256_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_first (const u32 *ukey1, const u32 *ukey2, const u32 *in, u32 *out, u32 *S, u32 *T, u32 *ks)
{
out[0] = in[0];
@ -46,6 +92,62 @@ DECLSPEC void kuznyechik_decrypt_xts_next (const u32 *in, u32 *out, u32 *T, u32
// 512 bit
DECLSPEC int verify_header_camellia (__global const tc_t *esalt_bufs, const u32 *ukey1, const u32 *ukey2)
{
u32 ks_camellia[68];
u32 S[4] = { 0 };
u32 T_camellia[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];
camellia256_decrypt_xts_first (ukey1, ukey2, data, tmp, S, T_camellia, ks_camellia);
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_camellia, T_camellia);
}
// 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];
camellia256_decrypt_xts_next (data, tmp, T_camellia, ks_camellia);
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 (__global const tc_t *esalt_bufs, const u32 *ukey1, const u32 *ukey2)
{
u32 ks_kuznyechik[40];
@ -104,6 +206,128 @@ DECLSPEC int verify_header_kuznyechik (__global const tc_t *esalt_bufs, const u3
// 1024 bit
DECLSPEC int verify_header_camellia_kuznyechik (__global const tc_t *esalt_bufs, const u32 *ukey1, const u32 *ukey2, const u32 *ukey3, const u32 *ukey4)
{
u32 ks_camellia[68];
u32 ks_kuznyechik[40];
u32 S[4] = { 0 };
u32 T_camellia[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];
camellia256_decrypt_xts_first (ukey2, ukey4, data, tmp, S, T_camellia, ks_camellia);
kuznyechik_decrypt_xts_first (ukey1, ukey3, tmp, 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_camellia, T_camellia);
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];
camellia256_decrypt_xts_next (data, tmp, T_camellia, ks_camellia);
kuznyechik_decrypt_xts_next (tmp, 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;
}
DECLSPEC int verify_header_camellia_serpent (__global const tc_t *esalt_bufs, const u32 *ukey1, const u32 *ukey2, const u32 *ukey3, const u32 *ukey4)
{
u32 ks_camellia[68];
u32 ks_serpent[140];
u32 S[4] = { 0 };
u32 T_camellia[4] = { 0 };
u32 T_serpent[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];
camellia256_decrypt_xts_first (ukey2, ukey4, data, tmp, S, T_camellia, ks_camellia);
serpent256_decrypt_xts_first (ukey1, ukey3, tmp, tmp, S, T_serpent, ks_serpent);
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_camellia, T_camellia);
xts_mul2 (T_serpent, T_serpent);
}
// 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];
camellia256_decrypt_xts_next (data, tmp, T_camellia, ks_camellia);
serpent256_decrypt_xts_next (tmp, tmp, T_serpent, ks_serpent);
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_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];
@ -227,3 +451,71 @@ DECLSPEC int verify_header_kuznyechik_twofish (__global const tc_t *esalt_bufs,
return 1;
}
// 1536 bit
DECLSPEC int verify_header_kuznyechik_serpent_camellia (__global const tc_t *esalt_bufs, const u32 *ukey1, const u32 *ukey2, const u32 *ukey3, const u32 *ukey4, const u32 *ukey5, const u32 *ukey6)
{
u32 ks_kuznyechik[40];
u32 ks_serpent[140];
u32 ks_camellia[68];
u32 S[4] = { 0 };
u32 T_kuznyechik[4] = { 0 };
u32 T_serpent[4] = { 0 };
u32 T_camellia[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 (ukey3, ukey6, data, tmp, S, T_kuznyechik, ks_kuznyechik);
serpent256_decrypt_xts_first (ukey2, ukey5, tmp, tmp, S, T_serpent, ks_serpent);
camellia256_decrypt_xts_first (ukey1, ukey4, tmp, tmp, S, T_camellia, ks_camellia);
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_serpent, T_serpent);
xts_mul2 (T_camellia, T_camellia);
}
// 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);
serpent256_decrypt_xts_next (tmp, tmp, T_serpent, ks_serpent);
camellia256_decrypt_xts_next (tmp, tmp, T_camellia, ks_camellia);
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,10 +16,13 @@
#include "inc_cipher_aes.cl"
#include "inc_cipher_twofish.cl"
#include "inc_cipher_serpent.cl"
#include "inc_cipher_camellia.cl"
#include "inc_cipher_kuznyechik.cl"
#include "inc_truecrypt_keyfile.cl"
#include "inc_truecrypt_crc32.cl"
#include "inc_truecrypt_xts.cl"
#include "inc_veracrypt_xts.cl"
DECLSPEC void hmac_ripemd160_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest)
{
@ -373,4 +376,20 @@ __kernel void m06211_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_camellia (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);
}
}
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,10 +16,13 @@
#include "inc_cipher_aes.cl"
#include "inc_cipher_twofish.cl"
#include "inc_cipher_serpent.cl"
#include "inc_cipher_camellia.cl"
#include "inc_cipher_kuznyechik.cl"
#include "inc_truecrypt_keyfile.cl"
#include "inc_truecrypt_crc32.cl"
#include "inc_truecrypt_xts.cl"
#include "inc_veracrypt_xts.cl"
DECLSPEC void hmac_ripemd160_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest)
{
@ -374,6 +377,22 @@ __kernel void m06212_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
}
}
if (verify_header_camellia (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);
}
}
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] = tmps[gid].out[16];
@ -419,4 +438,36 @@ __kernel void m06212_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_camellia_kuznyechik (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
{
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
if (verify_header_camellia_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
{
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{
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,10 +16,13 @@
#include "inc_cipher_aes.cl"
#include "inc_cipher_twofish.cl"
#include "inc_cipher_serpent.cl"
#include "inc_cipher_camellia.cl"
#include "inc_cipher_kuznyechik.cl"
#include "inc_truecrypt_keyfile.cl"
#include "inc_truecrypt_crc32.cl"
#include "inc_truecrypt_xts.cl"
#include "inc_veracrypt_xts.cl"
DECLSPEC void hmac_ripemd160_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest)
{
@ -374,6 +377,22 @@ __kernel void m06213_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
}
}
if (verify_header_camellia (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);
}
}
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] = tmps[gid].out[16];
@ -420,6 +439,38 @@ __kernel void m06213_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
}
}
if (verify_header_camellia_kuznyechik (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
{
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
if (verify_header_camellia_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
{
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{
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);
}
}
u32 ukey5[8];
ukey5[0] = tmps[gid].out[32];
@ -457,4 +508,12 @@ __kernel void m06213_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_serpent_camellia (esalt_bufs, ukey1, ukey2, ukey3, ukey4, ukey5, ukey6) == 1)
{
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
}

View File

@ -16,6 +16,7 @@
#include "inc_cipher_aes.cl"
#include "inc_cipher_twofish.cl"
#include "inc_cipher_serpent.cl"
#include "inc_cipher_camellia.cl"
#include "inc_cipher_kuznyechik.cl"
#include "inc_truecrypt_keyfile.cl"
@ -520,6 +521,14 @@ __kernel void m06221_comp (KERN_ATTR_TMPS_ESALT (tc64_tmp_t, tc_t))
}
}
if (verify_header_camellia (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);
}
}
if (verify_header_kuznyechik (esalt_bufs, ukey1, ukey2) == 1)
{
if (atomic_inc (&hashes_shown[0]) == 0)

View File

@ -16,6 +16,7 @@
#include "inc_cipher_aes.cl"
#include "inc_cipher_twofish.cl"
#include "inc_cipher_serpent.cl"
#include "inc_cipher_camellia.cl"
#include "inc_cipher_kuznyechik.cl"
#include "inc_truecrypt_keyfile.cl"
@ -520,6 +521,14 @@ __kernel void m06222_comp (KERN_ATTR_TMPS_ESALT (tc64_tmp_t, tc_t))
}
}
if (verify_header_camellia (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);
}
}
if (verify_header_kuznyechik (esalt_bufs, ukey1, ukey2) == 1)
{
if (atomic_inc (&hashes_shown[0]) == 0)
@ -574,6 +583,22 @@ __kernel void m06222_comp (KERN_ATTR_TMPS_ESALT (tc64_tmp_t, tc_t))
}
}
if (verify_header_camellia_kuznyechik (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
{
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
if (verify_header_camellia_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
{
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{
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)

View File

@ -16,6 +16,7 @@
#include "inc_cipher_aes.cl"
#include "inc_cipher_twofish.cl"
#include "inc_cipher_serpent.cl"
#include "inc_cipher_camellia.cl"
#include "inc_cipher_kuznyechik.cl"
#include "inc_truecrypt_keyfile.cl"
@ -520,6 +521,14 @@ __kernel void m06223_comp (KERN_ATTR_TMPS_ESALT (tc64_tmp_t, tc_t))
}
}
if (verify_header_camellia (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);
}
}
if (verify_header_kuznyechik (esalt_bufs, ukey1, ukey2) == 1)
{
if (atomic_inc (&hashes_shown[0]) == 0)
@ -574,6 +583,22 @@ __kernel void m06223_comp (KERN_ATTR_TMPS_ESALT (tc64_tmp_t, tc_t))
}
}
if (verify_header_camellia_kuznyechik (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
{
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
if (verify_header_camellia_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
{
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{
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)
@ -627,4 +652,12 @@ __kernel void m06223_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_serpent_camellia (esalt_bufs, ukey1, ukey2, ukey3, ukey4, ukey5, ukey6) == 1)
{
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
}

View File

@ -16,6 +16,7 @@
#include "inc_cipher_aes.cl"
#include "inc_cipher_twofish.cl"
#include "inc_cipher_serpent.cl"
#include "inc_cipher_camellia.cl"
#include "inc_cipher_kuznyechik.cl"
#include "inc_truecrypt_keyfile.cl"
@ -634,6 +635,14 @@ __kernel void m06231_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
}
}
if (verify_header_camellia (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);
}
}
if (verify_header_kuznyechik (esalt_bufs, ukey1, ukey2) == 1)
{
if (atomic_inc (&hashes_shown[0]) == 0)

View File

@ -16,6 +16,7 @@
#include "inc_cipher_aes.cl"
#include "inc_cipher_twofish.cl"
#include "inc_cipher_serpent.cl"
#include "inc_cipher_camellia.cl"
#include "inc_cipher_kuznyechik.cl"
#include "inc_truecrypt_keyfile.cl"
@ -634,6 +635,14 @@ __kernel void m06232_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
}
}
if (verify_header_camellia (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);
}
}
if (verify_header_kuznyechik (esalt_bufs, ukey1, ukey2) == 1)
{
if (atomic_inc (&hashes_shown[0]) == 0)
@ -688,6 +697,22 @@ __kernel void m06232_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
}
}
if (verify_header_camellia_kuznyechik (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
{
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
if (verify_header_camellia_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
{
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{
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)

View File

@ -16,6 +16,7 @@
#include "inc_cipher_aes.cl"
#include "inc_cipher_twofish.cl"
#include "inc_cipher_serpent.cl"
#include "inc_cipher_camellia.cl"
#include "inc_cipher_kuznyechik.cl"
#include "inc_truecrypt_keyfile.cl"
@ -634,6 +635,22 @@ __kernel void m06233_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
}
}
if (verify_header_camellia (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);
}
}
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]);
@ -680,7 +697,31 @@ __kernel void m06233_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
}
}
if (verify_header_kuznyechik (esalt_bufs, ukey1, ukey2) == 1)
if (verify_header_camellia_kuznyechik (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
{
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
if (verify_header_camellia_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
{
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{
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)
{
@ -726,17 +767,9 @@ __kernel void m06233_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 (verify_header_kuznyechik_serpent_camellia (esalt_bufs, ukey1, ukey2, ukey3, ukey4, ukey5, ukey6) == 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)
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}

View File

@ -16,6 +16,7 @@
#include "inc_cipher_aes.cl"
#include "inc_cipher_twofish.cl"
#include "inc_cipher_serpent.cl"
#include "inc_cipher_camellia.cl"
#include "inc_cipher_kuznyechik.cl"
#include "inc_truecrypt_keyfile.cl"
@ -432,6 +433,14 @@ __kernel void m13751_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
}
}
if (verify_header_camellia (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);
}
}
if (verify_header_kuznyechik (esalt_bufs, ukey1, ukey2) == 1)
{
if (atomic_inc (&hashes_shown[0]) == 0)

View File

@ -16,6 +16,7 @@
#include "inc_cipher_aes.cl"
#include "inc_cipher_twofish.cl"
#include "inc_cipher_serpent.cl"
#include "inc_cipher_camellia.cl"
#include "inc_cipher_kuznyechik.cl"
#include "inc_truecrypt_keyfile.cl"
@ -432,6 +433,14 @@ __kernel void m13752_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
}
}
if (verify_header_camellia (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);
}
}
if (verify_header_kuznyechik (esalt_bufs, ukey1, ukey2) == 1)
{
if (atomic_inc (&hashes_shown[0]) == 0)
@ -486,6 +495,22 @@ __kernel void m13752_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
}
}
if (verify_header_camellia_kuznyechik (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
{
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
if (verify_header_camellia_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
{
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{
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)

View File

@ -16,6 +16,7 @@
#include "inc_cipher_aes.cl"
#include "inc_cipher_twofish.cl"
#include "inc_cipher_serpent.cl"
#include "inc_cipher_camellia.cl"
#include "inc_cipher_kuznyechik.cl"
#include "inc_truecrypt_keyfile.cl"
@ -432,6 +433,14 @@ __kernel void m13753_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
}
}
if (verify_header_camellia (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);
}
}
if (verify_header_kuznyechik (esalt_bufs, ukey1, ukey2) == 1)
{
if (atomic_inc (&hashes_shown[0]) == 0)
@ -486,6 +495,22 @@ __kernel void m13753_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
}
}
if (verify_header_camellia_kuznyechik (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
{
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
if (verify_header_camellia_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
{
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{
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)
@ -539,4 +564,12 @@ __kernel void m13753_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_serpent_camellia (esalt_bufs, ukey1, ukey2, ukey3, ukey4, ukey5, ukey6) == 1)
{
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
}

View File

@ -16,6 +16,7 @@
#include "inc_cipher_aes.cl"
#include "inc_cipher_twofish.cl"
#include "inc_cipher_serpent.cl"
#include "inc_cipher_camellia.cl"
#include "inc_cipher_kuznyechik.cl"
#include "inc_truecrypt_keyfile.cl"
@ -561,6 +562,14 @@ __kernel void m13771_comp (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, tc_t))
}
}
if (verify_header_camellia (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);
}
}
if (verify_header_kuznyechik (esalt_bufs, ukey1, ukey2) == 1)
{
if (atomic_inc (&hashes_shown[0]) == 0)

View File

@ -16,6 +16,7 @@
#include "inc_cipher_aes.cl"
#include "inc_cipher_twofish.cl"
#include "inc_cipher_serpent.cl"
#include "inc_cipher_camellia.cl"
#include "inc_cipher_kuznyechik.cl"
#include "inc_truecrypt_keyfile.cl"
@ -562,6 +563,14 @@ __kernel void m13772_comp (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, tc_t))
}
}
if (verify_header_camellia (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);
}
}
if (verify_header_kuznyechik (esalt_bufs, ukey1, ukey2) == 1)
{
if (atomic_inc (&hashes_shown[0]) == 0)
@ -616,6 +625,22 @@ __kernel void m13772_comp (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, tc_t))
}
}
if (verify_header_camellia_kuznyechik (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
{
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
if (verify_header_camellia_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
{
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{
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)

View File

@ -16,6 +16,7 @@
#include "inc_cipher_aes.cl"
#include "inc_cipher_twofish.cl"
#include "inc_cipher_serpent.cl"
#include "inc_cipher_camellia.cl"
#include "inc_cipher_kuznyechik.cl"
#include "inc_truecrypt_keyfile.cl"
@ -562,6 +563,14 @@ __kernel void m13773_comp (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, tc_t))
}
}
if (verify_header_camellia (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);
}
}
if (verify_header_kuznyechik (esalt_bufs, ukey1, ukey2) == 1)
{
if (atomic_inc (&hashes_shown[0]) == 0)
@ -616,6 +625,38 @@ __kernel void m13773_comp (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, tc_t))
}
}
if (verify_header_camellia_kuznyechik (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
{
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
}
}
if (verify_header_camellia_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
{
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{
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);
}
}
u32 ukey5[8];
ukey5[0] = swap32_S (h32_from_64_S (tmps[gid].out[23]));
@ -654,17 +695,9 @@ __kernel void m13773_comp (KERN_ATTR_TMPS_ESALT (vc64_sbog_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 (verify_header_kuznyechik_serpent_camellia (esalt_bufs, ukey1, ukey2, ukey3, ukey4, ukey5, ukey6) == 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)
if (atomic_inc (&hashes_shown[digests_offset]) == 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 18200 (Kerberos 5 AS-REP etype 23)
- Added hash-mode 18300 (Apple File System (APFS))
- Added Kuznyechik cipher and cascades support for VeraCrypt kernels
- Added Camellia cipher and cascades support for VeraCrypt kernels
##
## Improvements
@ -92,6 +93,7 @@
- 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
- Tests: Added VeraCrypt containers for Camellia cipher and cascades
* changes v4.2.1 -> v5.0.0

View File

@ -359,12 +359,16 @@ 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 Camellia | 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 Camellia | 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 Camellia-Kuznyechik | Full-Disk Encryption (FDE)",
" Y | 2 = XTS 1024 bit cascaded Camellia-Serpent | 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)",

View File

@ -2253,7 +2253,7 @@ function truecrypt_test()
# Compose and execute hashcat command on a VeraCrypt test container
# Must not be called for hash types other than 137XY
# $1: cipher variation, can be 0-4
# $1: cipher variation, can be 0-6
function veracrypt_test()
{
cipher_variation=$1
@ -2277,18 +2277,22 @@ function veracrypt_test()
[ $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"
[ $cipher_variation -eq "3" ] && cipher_cascade="camellia"
[ $cipher_variation -eq "5" ] && cipher_cascade="kuznyechik"
;;
2)
[ $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"
[ $cipher_variation -eq "3" ] && cipher_cascade="camellia-kuznyechik"
[ $cipher_variation -eq "4" ] && cipher_cascade="camellia-serpent"
[ $cipher_variation -eq "5" ] && cipher_cascade="kuznyechik-aes"
[ $cipher_variation -eq "6" ] && cipher_cascade="kuznyechik-twofish"
;;
3)
[ $cipher_variation -eq "0" ] && cipher_cascade="aes-twofish-serpent"
[ $cipher_variation -eq "1" ] && cipher_cascade="serpent-twofish-aes"
[ $cipher_variation -eq "5" ] && cipher_cascade="kuznyechik-serpent-camellia"
;;
esac
@ -2777,8 +2781,10 @@ if [ "${PACKAGE}" -eq 0 -o -z "${PACKAGE_FOLDER}" ]; then
veracrypt_test 0 # aes
veracrypt_test 1 # serpent
veracrypt_test 2 # twofish
veracrypt_test 3 # kuznyechik
veracrypt_test 4 # kuznyechik (2nd cascade)
veracrypt_test 3 # camellia
veracrypt_test 4 # camellia (alternative cascade)
veracrypt_test 5 # kuznyechik
veracrypt_test 6 # kuznyechik (alternative cascade)
elif [[ ${hash_type} -ge 6211 ]] && [[ ${hash_type} -le 6243 ]]; then
# run truecrypt tests

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.