diff --git a/OpenCL/inc_hash_md4.cl b/OpenCL/inc_hash_md4.cl index 3a3f525a2..eff5d3652 100644 --- a/OpenCL/inc_hash_md4.cl +++ b/OpenCL/inc_hash_md4.cl @@ -194,42 +194,42 @@ void md4_update (md4_ctx_t *ctx, const u32 *w, const int len) for (i = 0, j = 0; i < len - 64; i += 64, j += 16) { - w0[0] = w[i + 0]; - w0[1] = w[i + 1]; - w0[2] = w[i + 2]; - w0[3] = w[i + 3]; - w1[0] = w[i + 4]; - w1[1] = w[i + 5]; - w1[2] = w[i + 6]; - w1[3] = w[i + 7]; - w2[0] = w[i + 8]; - w2[1] = w[i + 9]; - w2[2] = w[i + 10]; - w2[3] = w[i + 11]; - w3[0] = w[i + 12]; - w3[1] = w[i + 13]; - w3[2] = w[i + 14]; - w3[3] = w[i + 15]; + w0[0] = w[j + 0]; + w0[1] = w[j + 1]; + w0[2] = w[j + 2]; + w0[3] = w[j + 3]; + w1[0] = w[j + 4]; + w1[1] = w[j + 5]; + w1[2] = w[j + 6]; + w1[3] = w[j + 7]; + w2[0] = w[j + 8]; + w2[1] = w[j + 9]; + w2[2] = w[j + 10]; + w2[3] = w[j + 11]; + w3[0] = w[j + 12]; + w3[1] = w[j + 13]; + w3[2] = w[j + 14]; + w3[3] = w[j + 15]; md4_update_64 (ctx, w0, w1, w2, w3, 64); } - w0[0] = w[i + 0]; - w0[1] = w[i + 1]; - w0[2] = w[i + 2]; - w0[3] = w[i + 3]; - w1[0] = w[i + 4]; - w1[1] = w[i + 5]; - w1[2] = w[i + 6]; - w1[3] = w[i + 7]; - w2[0] = w[i + 8]; - w2[1] = w[i + 9]; - w2[2] = w[i + 10]; - w2[3] = w[i + 11]; - w3[0] = w[i + 12]; - w3[1] = w[i + 13]; - w3[2] = w[i + 14]; - w3[3] = w[i + 15]; + w0[0] = w[j + 0]; + w0[1] = w[j + 1]; + w0[2] = w[j + 2]; + w0[3] = w[j + 3]; + w1[0] = w[j + 4]; + w1[1] = w[j + 5]; + w1[2] = w[j + 6]; + w1[3] = w[j + 7]; + w2[0] = w[j + 8]; + w2[1] = w[j + 9]; + w2[2] = w[j + 10]; + w2[3] = w[j + 11]; + w3[0] = w[j + 12]; + w3[1] = w[j + 13]; + w3[2] = w[j + 14]; + w3[3] = w[j + 15]; md4_update_64 (ctx, w0, w1, w2, w3, len & 63); } @@ -246,42 +246,42 @@ void md4_update_global (md4_ctx_t *ctx, const __global u32 *w, const int len) for (i = 0, j = 0; i < len - 64; i += 64, j += 16) { - w0[0] = w[i + 0]; - w0[1] = w[i + 1]; - w0[2] = w[i + 2]; - w0[3] = w[i + 3]; - w1[0] = w[i + 4]; - w1[1] = w[i + 5]; - w1[2] = w[i + 6]; - w1[3] = w[i + 7]; - w2[0] = w[i + 8]; - w2[1] = w[i + 9]; - w2[2] = w[i + 10]; - w2[3] = w[i + 11]; - w3[0] = w[i + 12]; - w3[1] = w[i + 13]; - w3[2] = w[i + 14]; - w3[3] = w[i + 15]; + w0[0] = w[j + 0]; + w0[1] = w[j + 1]; + w0[2] = w[j + 2]; + w0[3] = w[j + 3]; + w1[0] = w[j + 4]; + w1[1] = w[j + 5]; + w1[2] = w[j + 6]; + w1[3] = w[j + 7]; + w2[0] = w[j + 8]; + w2[1] = w[j + 9]; + w2[2] = w[j + 10]; + w2[3] = w[j + 11]; + w3[0] = w[j + 12]; + w3[1] = w[j + 13]; + w3[2] = w[j + 14]; + w3[3] = w[j + 15]; md4_update_64 (ctx, w0, w1, w2, w3, 64); } - w0[0] = w[i + 0]; - w0[1] = w[i + 1]; - w0[2] = w[i + 2]; - w0[3] = w[i + 3]; - w1[0] = w[i + 4]; - w1[1] = w[i + 5]; - w1[2] = w[i + 6]; - w1[3] = w[i + 7]; - w2[0] = w[i + 8]; - w2[1] = w[i + 9]; - w2[2] = w[i + 10]; - w2[3] = w[i + 11]; - w3[0] = w[i + 12]; - w3[1] = w[i + 13]; - w3[2] = w[i + 14]; - w3[3] = w[i + 15]; + w0[0] = w[j + 0]; + w0[1] = w[j + 1]; + w0[2] = w[j + 2]; + w0[3] = w[j + 3]; + w1[0] = w[j + 4]; + w1[1] = w[j + 5]; + w1[2] = w[j + 6]; + w1[3] = w[j + 7]; + w2[0] = w[j + 8]; + w2[1] = w[j + 9]; + w2[2] = w[j + 10]; + w2[3] = w[j + 11]; + w3[0] = w[j + 12]; + w3[1] = w[j + 13]; + w3[2] = w[j + 14]; + w3[3] = w[j + 15]; md4_update_64 (ctx, w0, w1, w2, w3, len & 63); } @@ -298,14 +298,14 @@ void md4_update_global_utf16le (md4_ctx_t *ctx, const __global u32 *w, const int for (i = 0, j = 0; i < len - 32; i += 32, j += 8) { - w0[0] = w[i + 0]; - w0[1] = w[i + 1]; - w0[2] = w[i + 2]; - w0[3] = w[i + 3]; - w1[0] = w[i + 4]; - w1[1] = w[i + 5]; - w1[2] = w[i + 6]; - w1[3] = w[i + 7]; + w0[0] = w[j + 0]; + w0[1] = w[j + 1]; + w0[2] = w[j + 2]; + w0[3] = w[j + 3]; + w1[0] = w[j + 4]; + w1[1] = w[j + 5]; + w1[2] = w[j + 6]; + w1[3] = w[j + 7]; make_utf16le_S (w1, w2, w3); make_utf16le_S (w0, w0, w1); @@ -313,14 +313,14 @@ void md4_update_global_utf16le (md4_ctx_t *ctx, const __global u32 *w, const int md4_update_64 (ctx, w0, w1, w2, w3, 32 * 2); } - w0[0] = w[i + 0]; - w0[1] = w[i + 1]; - w0[2] = w[i + 2]; - w0[3] = w[i + 3]; - w1[0] = w[i + 4]; - w1[1] = w[i + 5]; - w1[2] = w[i + 6]; - w1[3] = w[i + 7]; + w0[0] = w[j + 0]; + w0[1] = w[j + 1]; + w0[2] = w[j + 2]; + w0[3] = w[j + 3]; + w1[0] = w[j + 4]; + w1[1] = w[j + 5]; + w1[2] = w[j + 6]; + w1[3] = w[j + 7]; make_utf16le_S (w1, w2, w3); make_utf16le_S (w0, w0, w1); @@ -559,42 +559,42 @@ void md4_update_vector (md4_ctx_vector_t *ctx, const u32x *w, const int len) for (i = 0, j = 0; i < len - 64; i += 64, j += 16) { - w0[0] = w[i + 0]; - w0[1] = w[i + 1]; - w0[2] = w[i + 2]; - w0[3] = w[i + 3]; - w1[0] = w[i + 4]; - w1[1] = w[i + 5]; - w1[2] = w[i + 6]; - w1[3] = w[i + 7]; - w2[0] = w[i + 8]; - w2[1] = w[i + 9]; - w2[2] = w[i + 10]; - w2[3] = w[i + 11]; - w3[0] = w[i + 12]; - w3[1] = w[i + 13]; - w3[2] = w[i + 14]; - w3[3] = w[i + 15]; + w0[0] = w[j + 0]; + w0[1] = w[j + 1]; + w0[2] = w[j + 2]; + w0[3] = w[j + 3]; + w1[0] = w[j + 4]; + w1[1] = w[j + 5]; + w1[2] = w[j + 6]; + w1[3] = w[j + 7]; + w2[0] = w[j + 8]; + w2[1] = w[j + 9]; + w2[2] = w[j + 10]; + w2[3] = w[j + 11]; + w3[0] = w[j + 12]; + w3[1] = w[j + 13]; + w3[2] = w[j + 14]; + w3[3] = w[j + 15]; md4_update_vector_64 (ctx, w0, w1, w2, w3, 64); } - w0[0] = w[i + 0]; - w0[1] = w[i + 1]; - w0[2] = w[i + 2]; - w0[3] = w[i + 3]; - w1[0] = w[i + 4]; - w1[1] = w[i + 5]; - w1[2] = w[i + 6]; - w1[3] = w[i + 7]; - w2[0] = w[i + 8]; - w2[1] = w[i + 9]; - w2[2] = w[i + 10]; - w2[3] = w[i + 11]; - w3[0] = w[i + 12]; - w3[1] = w[i + 13]; - w3[2] = w[i + 14]; - w3[3] = w[i + 15]; + w0[0] = w[j + 0]; + w0[1] = w[j + 1]; + w0[2] = w[j + 2]; + w0[3] = w[j + 3]; + w1[0] = w[j + 4]; + w1[1] = w[j + 5]; + w1[2] = w[j + 6]; + w1[3] = w[j + 7]; + w2[0] = w[j + 8]; + w2[1] = w[j + 9]; + w2[2] = w[j + 10]; + w2[3] = w[j + 11]; + w3[0] = w[j + 12]; + w3[1] = w[j + 13]; + w3[2] = w[j + 14]; + w3[3] = w[j + 15]; md4_update_vector_64 (ctx, w0, w1, w2, w3, len & 63); } diff --git a/OpenCL/inc_hash_sha1.cl b/OpenCL/inc_hash_sha1.cl index b1e6f4b59..33a21e643 100644 --- a/OpenCL/inc_hash_sha1.cl +++ b/OpenCL/inc_hash_sha1.cl @@ -259,42 +259,42 @@ void sha1_update (sha1_ctx_t *ctx, const u32 *w, const int len) for (i = 0, j = 0; i < len - 64; i += 64, j += 16) { - w0[0] = w[i + 0]; - w0[1] = w[i + 1]; - w0[2] = w[i + 2]; - w0[3] = w[i + 3]; - w1[0] = w[i + 4]; - w1[1] = w[i + 5]; - w1[2] = w[i + 6]; - w1[3] = w[i + 7]; - w2[0] = w[i + 8]; - w2[1] = w[i + 9]; - w2[2] = w[i + 10]; - w2[3] = w[i + 11]; - w3[0] = w[i + 12]; - w3[1] = w[i + 13]; - w3[2] = w[i + 14]; - w3[3] = w[i + 15]; + w0[0] = w[j + 0]; + w0[1] = w[j + 1]; + w0[2] = w[j + 2]; + w0[3] = w[j + 3]; + w1[0] = w[j + 4]; + w1[1] = w[j + 5]; + w1[2] = w[j + 6]; + w1[3] = w[j + 7]; + w2[0] = w[j + 8]; + w2[1] = w[j + 9]; + w2[2] = w[j + 10]; + w2[3] = w[j + 11]; + w3[0] = w[j + 12]; + w3[1] = w[j + 13]; + w3[2] = w[j + 14]; + w3[3] = w[j + 15]; sha1_update_64 (ctx, w0, w1, w2, w3, 64); } - w0[0] = w[i + 0]; - w0[1] = w[i + 1]; - w0[2] = w[i + 2]; - w0[3] = w[i + 3]; - w1[0] = w[i + 4]; - w1[1] = w[i + 5]; - w1[2] = w[i + 6]; - w1[3] = w[i + 7]; - w2[0] = w[i + 8]; - w2[1] = w[i + 9]; - w2[2] = w[i + 10]; - w2[3] = w[i + 11]; - w3[0] = w[i + 12]; - w3[1] = w[i + 13]; - w3[2] = w[i + 14]; - w3[3] = w[i + 15]; + w0[0] = w[j + 0]; + w0[1] = w[j + 1]; + w0[2] = w[j + 2]; + w0[3] = w[j + 3]; + w1[0] = w[j + 4]; + w1[1] = w[j + 5]; + w1[2] = w[j + 6]; + w1[3] = w[j + 7]; + w2[0] = w[j + 8]; + w2[1] = w[j + 9]; + w2[2] = w[j + 10]; + w2[3] = w[j + 11]; + w3[0] = w[j + 12]; + w3[1] = w[j + 13]; + w3[2] = w[j + 14]; + w3[3] = w[j + 15]; sha1_update_64 (ctx, w0, w1, w2, w3, len & 63); } @@ -311,42 +311,42 @@ void sha1_update_global (sha1_ctx_t *ctx, const __global u32 *w, const int len) for (i = 0, j = 0; i < len - 64; i += 64, j += 16) { - w0[0] = w[i + 0]; - w0[1] = w[i + 1]; - w0[2] = w[i + 2]; - w0[3] = w[i + 3]; - w1[0] = w[i + 4]; - w1[1] = w[i + 5]; - w1[2] = w[i + 6]; - w1[3] = w[i + 7]; - w2[0] = w[i + 8]; - w2[1] = w[i + 9]; - w2[2] = w[i + 10]; - w2[3] = w[i + 11]; - w3[0] = w[i + 12]; - w3[1] = w[i + 13]; - w3[2] = w[i + 14]; - w3[3] = w[i + 15]; + w0[0] = w[j + 0]; + w0[1] = w[j + 1]; + w0[2] = w[j + 2]; + w0[3] = w[j + 3]; + w1[0] = w[j + 4]; + w1[1] = w[j + 5]; + w1[2] = w[j + 6]; + w1[3] = w[j + 7]; + w2[0] = w[j + 8]; + w2[1] = w[j + 9]; + w2[2] = w[j + 10]; + w2[3] = w[j + 11]; + w3[0] = w[j + 12]; + w3[1] = w[j + 13]; + w3[2] = w[j + 14]; + w3[3] = w[j + 15]; sha1_update_64 (ctx, w0, w1, w2, w3, 64); } - w0[0] = w[i + 0]; - w0[1] = w[i + 1]; - w0[2] = w[i + 2]; - w0[3] = w[i + 3]; - w1[0] = w[i + 4]; - w1[1] = w[i + 5]; - w1[2] = w[i + 6]; - w1[3] = w[i + 7]; - w2[0] = w[i + 8]; - w2[1] = w[i + 9]; - w2[2] = w[i + 10]; - w2[3] = w[i + 11]; - w3[0] = w[i + 12]; - w3[1] = w[i + 13]; - w3[2] = w[i + 14]; - w3[3] = w[i + 15]; + w0[0] = w[j + 0]; + w0[1] = w[j + 1]; + w0[2] = w[j + 2]; + w0[3] = w[j + 3]; + w1[0] = w[j + 4]; + w1[1] = w[j + 5]; + w1[2] = w[j + 6]; + w1[3] = w[j + 7]; + w2[0] = w[j + 8]; + w2[1] = w[j + 9]; + w2[2] = w[j + 10]; + w2[3] = w[j + 11]; + w3[0] = w[j + 12]; + w3[1] = w[j + 13]; + w3[2] = w[j + 14]; + w3[3] = w[j + 15]; sha1_update_64 (ctx, w0, w1, w2, w3, len & 63); } @@ -363,42 +363,76 @@ void sha1_update_global_swap (sha1_ctx_t *ctx, const __global u32 *w, const int for (i = 0, j = 0; i < len - 64; i += 64, j += 16) { - w0[0] = swap32_S (w[i + 0]); - w0[1] = swap32_S (w[i + 1]); - w0[2] = swap32_S (w[i + 2]); - w0[3] = swap32_S (w[i + 3]); - w1[0] = swap32_S (w[i + 4]); - w1[1] = swap32_S (w[i + 5]); - w1[2] = swap32_S (w[i + 6]); - w1[3] = swap32_S (w[i + 7]); - w2[0] = swap32_S (w[i + 8]); - w2[1] = swap32_S (w[i + 9]); - w2[2] = swap32_S (w[i + 10]); - w2[3] = swap32_S (w[i + 11]); - w3[0] = swap32_S (w[i + 12]); - w3[1] = swap32_S (w[i + 13]); - w3[2] = swap32_S (w[i + 14]); - w3[3] = swap32_S (w[i + 15]); + w0[0] = w[j + 0]; + w0[1] = w[j + 1]; + w0[2] = w[j + 2]; + w0[3] = w[j + 3]; + w1[0] = w[j + 4]; + w1[1] = w[j + 5]; + w1[2] = w[j + 6]; + w1[3] = w[j + 7]; + w2[0] = w[j + 8]; + w2[1] = w[j + 9]; + w2[2] = w[j + 10]; + w2[3] = w[j + 11]; + w3[0] = w[j + 12]; + w3[1] = w[j + 13]; + w3[2] = w[j + 14]; + w3[3] = w[j + 15]; + + w0[0] = swap32_S (w0[0]); + w0[1] = swap32_S (w0[1]); + w0[2] = swap32_S (w0[2]); + w0[3] = swap32_S (w0[3]); + w1[0] = swap32_S (w1[0]); + w1[1] = swap32_S (w1[1]); + w1[2] = swap32_S (w1[2]); + w1[3] = swap32_S (w1[3]); + w2[0] = swap32_S (w2[0]); + w2[1] = swap32_S (w2[1]); + w2[2] = swap32_S (w2[2]); + w2[3] = swap32_S (w2[3]); + w3[0] = swap32_S (w3[0]); + w3[1] = swap32_S (w3[1]); + w3[2] = swap32_S (w3[2]); + w3[3] = swap32_S (w3[3]); sha1_update_64 (ctx, w0, w1, w2, w3, 64); } - w0[0] = swap32_S (w[i + 0]); - w0[1] = swap32_S (w[i + 1]); - w0[2] = swap32_S (w[i + 2]); - w0[3] = swap32_S (w[i + 3]); - w1[0] = swap32_S (w[i + 4]); - w1[1] = swap32_S (w[i + 5]); - w1[2] = swap32_S (w[i + 6]); - w1[3] = swap32_S (w[i + 7]); - w2[0] = swap32_S (w[i + 8]); - w2[1] = swap32_S (w[i + 9]); - w2[2] = swap32_S (w[i + 10]); - w2[3] = swap32_S (w[i + 11]); - w3[0] = swap32_S (w[i + 12]); - w3[1] = swap32_S (w[i + 13]); - w3[2] = swap32_S (w[i + 14]); - w3[3] = swap32_S (w[i + 15]); + w0[0] = w[j + 0]; + w0[1] = w[j + 1]; + w0[2] = w[j + 2]; + w0[3] = w[j + 3]; + w1[0] = w[j + 4]; + w1[1] = w[j + 5]; + w1[2] = w[j + 6]; + w1[3] = w[j + 7]; + w2[0] = w[j + 8]; + w2[1] = w[j + 9]; + w2[2] = w[j + 10]; + w2[3] = w[j + 11]; + w3[0] = w[j + 12]; + w3[1] = w[j + 13]; + w3[2] = w[j + 14]; + w3[3] = w[j + 15]; + + w0[0] = swap32_S (w0[0]); + w0[1] = swap32_S (w0[1]); + w0[2] = swap32_S (w0[2]); + w0[3] = swap32_S (w0[3]); + w1[0] = swap32_S (w1[0]); + w1[1] = swap32_S (w1[1]); + w1[2] = swap32_S (w1[2]); + w1[3] = swap32_S (w1[3]); + w2[0] = swap32_S (w2[0]); + w2[1] = swap32_S (w2[1]); + w2[2] = swap32_S (w2[2]); + w2[3] = swap32_S (w2[3]); + w3[0] = swap32_S (w3[0]); + w3[1] = swap32_S (w3[1]); + w3[2] = swap32_S (w3[2]); + w3[3] = swap32_S (w3[3]); sha1_update_64 (ctx, w0, w1, w2, w3, len & 63); } @@ -415,14 +449,14 @@ void sha1_update_global_utf16le (sha1_ctx_t *ctx, const __global u32 *w, const i for (i = 0, j = 0; i < len - 32; i += 32, j += 8) { - w0[0] = w[i + 0]; - w0[1] = w[i + 1]; - w0[2] = w[i + 2]; - w0[3] = w[i + 3]; - w1[0] = w[i + 4]; - w1[1] = w[i + 5]; - w1[2] = w[i + 6]; - w1[3] = w[i + 7]; + w0[0] = w[j + 0]; + w0[1] = w[j + 1]; + w0[2] = w[j + 2]; + w0[3] = w[j + 3]; + w1[0] = w[j + 4]; + w1[1] = w[j + 5]; + w1[2] = w[j + 6]; + w1[3] = w[j + 7]; make_utf16le_S (w1, w2, w3); make_utf16le_S (w0, w0, w1); @@ -430,14 +464,14 @@ void sha1_update_global_utf16le (sha1_ctx_t *ctx, const __global u32 *w, const i sha1_update_64 (ctx, w0, w1, w2, w3, 32 * 2); } - w0[0] = w[i + 0]; - w0[1] = w[i + 1]; - w0[2] = w[i + 2]; - w0[3] = w[i + 3]; - w1[0] = w[i + 4]; - w1[1] = w[i + 5]; - w1[2] = w[i + 6]; - w1[3] = w[i + 7]; + w0[0] = w[j + 0]; + w0[1] = w[j + 1]; + w0[2] = w[j + 2]; + w0[3] = w[j + 3]; + w1[0] = w[j + 4]; + w1[1] = w[j + 5]; + w1[2] = w[j + 6]; + w1[3] = w[j + 7]; make_utf16le_S (w1, w2, w3); make_utf16le_S (w0, w0, w1); @@ -445,6 +479,82 @@ void sha1_update_global_utf16le (sha1_ctx_t *ctx, const __global u32 *w, const i sha1_update_64 (ctx, w0, w1, w2, w3, (len & 31) * 2); } +void sha1_update_global_utf16le_swap (sha1_ctx_t *ctx, const __global u32 *w, const int len) +{ + u32 w0[4]; + u32 w1[4]; + u32 w2[4]; + u32 w3[4]; + + int i; + int j; + + for (i = 0, j = 0; i < len - 32; i += 32, j += 8) + { + w0[0] = w[j + 0]; + w0[1] = w[j + 1]; + w0[2] = w[j + 2]; + w0[3] = w[j + 3]; + w1[0] = w[j + 4]; + w1[1] = w[j + 5]; + w1[2] = w[j + 6]; + w1[3] = w[j + 7]; + + make_utf16le_S (w1, w2, w3); + make_utf16le_S (w0, w0, w1); + + w0[0] = swap32_S (w0[0]); + w0[1] = swap32_S (w0[1]); + w0[2] = swap32_S (w0[2]); + w0[3] = swap32_S (w0[3]); + w1[0] = swap32_S (w1[0]); + w1[1] = swap32_S (w1[1]); + w1[2] = swap32_S (w1[2]); + w1[3] = swap32_S (w1[3]); + w2[0] = swap32_S (w2[0]); + w2[1] = swap32_S (w2[1]); + w2[2] = swap32_S (w2[2]); + w2[3] = swap32_S (w2[3]); + w3[0] = swap32_S (w3[0]); + w3[1] = swap32_S (w3[1]); + w3[2] = swap32_S (w3[2]); + w3[3] = swap32_S (w3[3]); + + sha1_update_64 (ctx, w0, w1, w2, w3, 32 * 2); + } + + w0[0] = w[j + 0]; + w0[1] = w[j + 1]; + w0[2] = w[j + 2]; + w0[3] = w[j + 3]; + w1[0] = w[j + 4]; + w1[1] = w[j + 5]; + w1[2] = w[j + 6]; + w1[3] = w[j + 7]; + + make_utf16le_S (w1, w2, w3); + make_utf16le_S (w0, w0, w1); + + w0[0] = swap32_S (w0[0]); + w0[1] = swap32_S (w0[1]); + w0[2] = swap32_S (w0[2]); + w0[3] = swap32_S (w0[3]); + w1[0] = swap32_S (w1[0]); + w1[1] = swap32_S (w1[1]); + w1[2] = swap32_S (w1[2]); + w1[3] = swap32_S (w1[3]); + w2[0] = swap32_S (w2[0]); + w2[1] = swap32_S (w2[1]); + w2[2] = swap32_S (w2[2]); + w2[3] = swap32_S (w2[3]); + w3[0] = swap32_S (w3[0]); + w3[1] = swap32_S (w3[1]); + w3[2] = swap32_S (w3[2]); + w3[3] = swap32_S (w3[3]); + + sha1_update_64 (ctx, w0, w1, w2, w3, (len & 31) * 2); +} + void sha1_final (sha1_ctx_t *ctx) { int pos = ctx->len & 63; @@ -567,6 +677,16 @@ void sha1_hmac_update_global_swap (sha1_hmac_ctx_t *ctx, const __global u32 *w, sha1_update_global_swap (&ctx->ipad, w, len); } +void sha1_hmac_update_global_utf16le (sha1_hmac_ctx_t *ctx, const __global u32 *w, const int len) +{ + sha1_update_global_utf16le (&ctx->ipad, w, len); +} + +void sha1_hmac_update_global_utf16le_swap (sha1_hmac_ctx_t *ctx, const __global u32 *w, const int len) +{ + sha1_update_global_utf16le_swap (&ctx->ipad, w, len); +} + void sha1_hmac_final (sha1_hmac_ctx_t *ctx) { sha1_final (&ctx->ipad); @@ -855,42 +975,42 @@ void sha1_update_vector (sha1_ctx_vector_t *ctx, const u32x *w, const int len) for (i = 0, j = 0; i < len - 64; i += 64, j += 16) { - w0[0] = w[i + 0]; - w0[1] = w[i + 1]; - w0[2] = w[i + 2]; - w0[3] = w[i + 3]; - w1[0] = w[i + 4]; - w1[1] = w[i + 5]; - w1[2] = w[i + 6]; - w1[3] = w[i + 7]; - w2[0] = w[i + 8]; - w2[1] = w[i + 9]; - w2[2] = w[i + 10]; - w2[3] = w[i + 11]; - w3[0] = w[i + 12]; - w3[1] = w[i + 13]; - w3[2] = w[i + 14]; - w3[3] = w[i + 15]; + w0[0] = w[j + 0]; + w0[1] = w[j + 1]; + w0[2] = w[j + 2]; + w0[3] = w[j + 3]; + w1[0] = w[j + 4]; + w1[1] = w[j + 5]; + w1[2] = w[j + 6]; + w1[3] = w[j + 7]; + w2[0] = w[j + 8]; + w2[1] = w[j + 9]; + w2[2] = w[j + 10]; + w2[3] = w[j + 11]; + w3[0] = w[j + 12]; + w3[1] = w[j + 13]; + w3[2] = w[j + 14]; + w3[3] = w[j + 15]; sha1_update_vector_64 (ctx, w0, w1, w2, w3, 64); } - w0[0] = w[i + 0]; - w0[1] = w[i + 1]; - w0[2] = w[i + 2]; - w0[3] = w[i + 3]; - w1[0] = w[i + 4]; - w1[1] = w[i + 5]; - w1[2] = w[i + 6]; - w1[3] = w[i + 7]; - w2[0] = w[i + 8]; - w2[1] = w[i + 9]; - w2[2] = w[i + 10]; - w2[3] = w[i + 11]; - w3[0] = w[i + 12]; - w3[1] = w[i + 13]; - w3[2] = w[i + 14]; - w3[3] = w[i + 15]; + w0[0] = w[j + 0]; + w0[1] = w[j + 1]; + w0[2] = w[j + 2]; + w0[3] = w[j + 3]; + w1[0] = w[j + 4]; + w1[1] = w[j + 5]; + w1[2] = w[j + 6]; + w1[3] = w[j + 7]; + w2[0] = w[j + 8]; + w2[1] = w[j + 9]; + w2[2] = w[j + 10]; + w2[3] = w[j + 11]; + w3[0] = w[j + 12]; + w3[1] = w[j + 13]; + w3[2] = w[j + 14]; + w3[3] = w[j + 15]; sha1_update_vector_64 (ctx, w0, w1, w2, w3, len & 63); } diff --git a/OpenCL/m02100.cl b/OpenCL/m02100.cl index 1b6d11451..eeca0d64b 100644 --- a/OpenCL/m02100.cl +++ b/OpenCL/m02100.cl @@ -84,7 +84,7 @@ __kernel void m02100_init (__global pw_t *pws, __global const kernel_rule_t *rul md4_ctx2.len = 16; - md4_update_global (&md4_ctx2, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len); + md4_update_global_utf16le (&md4_ctx2, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len); md4_final (&md4_ctx2); @@ -133,7 +133,7 @@ __kernel void m02100_init (__global pw_t *pws, __global const kernel_rule_t *rul tmps[gid].opad[3] = sha1_hmac_ctx.opad.h[3]; tmps[gid].opad[4] = sha1_hmac_ctx.opad.h[4]; - sha1_hmac_update_global_swap (&sha1_hmac_ctx, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len); + sha1_hmac_update_global_utf16le_swap (&sha1_hmac_ctx, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len); w0[0] = 1; w0[1] = 0; diff --git a/include/interface.h b/include/interface.h index a21982a78..f5203ef5d 100644 --- a/include/interface.h +++ b/include/interface.h @@ -1047,10 +1047,10 @@ typedef enum display_len DISPLAY_LEN_MAX_1750H = 128 + 1 + 102, DISPLAY_LEN_MIN_1800 = 90 + 0, DISPLAY_LEN_MAX_1800 = 90 + 16, - DISPLAY_LEN_MIN_2100 = 6 + 1 + 1 + 32 + 1 + 0, - DISPLAY_LEN_MAX_2100 = 6 + 5 + 1 + 32 + 1 + 19, - DISPLAY_LEN_MIN_2100H = 6 + 1 + 1 + 32 + 1 + 0, - DISPLAY_LEN_MAX_2100H = 6 + 5 + 1 + 32 + 1 + 38, + DISPLAY_LEN_MIN_2100 = 6 + 1 + 1 + 1 + 1 + 32, + DISPLAY_LEN_MAX_2100 = 6 + 5 + 1 + 256 + 1 + 32, + DISPLAY_LEN_MIN_2100H = 6 + 1 + 1 + 2 + 1 + 32, + DISPLAY_LEN_MAX_2100H = 6 + 5 + 1 + 512 + 1 + 32, DISPLAY_LEN_MIN_2400 = 16, DISPLAY_LEN_MAX_2400 = 16, DISPLAY_LEN_MIN_2410 = 16 + 1 + 0, diff --git a/src/hashes.c b/src/hashes.c index a9130c235..fabe68eb4 100644 --- a/src/hashes.c +++ b/src/hashes.c @@ -64,13 +64,13 @@ int sort_by_salt (const void *v1, const void *v2) if (res2 != 0) return (res2); - for (int n = 0; n < 16; n++) + for (int n = 0; n < 64; n++) { if (s1->salt_buf[n] > s2->salt_buf[n]) return 1; if (s1->salt_buf[n] < s2->salt_buf[n]) return -1; } - for (int n = 0; n < 8; n++) + for (int n = 0; n < 64; n++) { if (s1->salt_buf_pc[n] > s2->salt_buf_pc[n]) return 1; if (s1->salt_buf_pc[n] < s2->salt_buf_pc[n]) return -1; diff --git a/src/interface.c b/src/interface.c index c63b20d90..ab3345047 100644 --- a/src/interface.c +++ b/src/interface.c @@ -12901,7 +12901,7 @@ int krb5tgs_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UN const u8 p1 = data_pos[i + 1]; *checksum_ptr++ = hex_convert (p1) << 0 - | hex_convert (p0) << 4; + | hex_convert (p0) << 4; } u8 *edata_ptr = (u8 *) krb5tgs->edata2; @@ -12913,8 +12913,9 @@ int krb5tgs_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UN { const u8 p0 = data_pos[i + 0]; const u8 p1 = data_pos[i + 1]; + *edata_ptr++ = hex_convert (p1) << 0 - | hex_convert (p0) << 4; + | hex_convert (p0) << 4; } /* this is needed for hmac_md5 */ @@ -21045,8 +21046,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx) hashconfig->salt_type = SALT_TYPE_EMBEDDED; hashconfig->attack_exec = ATTACK_EXEC_OUTSIDE_KERNEL; hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_LE - | OPTS_TYPE_ST_LOWER - | OPTS_TYPE_ST_UTF16LE; + | OPTS_TYPE_ST_LOWER; hashconfig->kern_type = KERN_TYPE_DCC2; hashconfig->dgst_size = DGST_SIZE_4_4; hashconfig->parse_func = dcc2_parse_hash; diff --git a/src/potfile.c b/src/potfile.c index b84cd4a9e..00be82cad 100644 --- a/src/potfile.c +++ b/src/potfile.c @@ -439,7 +439,7 @@ int potfile_remove_parse (hashcat_ctx_t *hashcat_ctx) if (hashconfig->hash_mode == 6800) { - if (line_hash_len < 64) // 64 = 16 * u32 in salt_buf[] + if (line_hash_len < 256) // 64 = 64 * u32 in salt_buf[] { // manipulate salt_buf memcpy (hash_buf.salt->salt_buf, line_hash_buf, line_hash_len); diff --git a/tools/test.pl b/tools/test.pl index 222d83f19..599865856 100755 --- a/tools/test.pl +++ b/tools/test.pl @@ -50,7 +50,8 @@ my $MAX_LEN = 55; my @modes = (0, 10, 11, 12, 20, 21, 22, 23, 30, 40, 50, 60, 100, 101, 110, 111, 112, 120, 121, 122, 125, 130, 131, 132, 133, 140, 141, 150, 160, 200, 300, 400, 500, 600, 900, 1000, 1100, 1300, 1400, 1410, 1411, 1420, 1430, 1440, 1441, 1450, 1460, 1500, 1600, 1700, 1710, 1711, 1720, 1730, 1740, 1722, 1731, 1750, 1760, 1800, 2100, 2400, 2410, 2500, 2600, 2611, 2612, 2711, 2811, 3000, 3100, 3200, 3710, 3711, 3300, 3500, 3610, 3720, 3800, 3910, 4010, 4110, 4210, 4300, 4400, 4500, 4520, 4521, 4522, 4600, 4700, 4800, 4900, 5000, 5100, 5300, 5400, 5500, 5600, 5700, 5800, 6000, 6100, 6300, 6400, 6500, 6600, 6700, 6800, 6900, 7000, 7100, 7200, 7300, 7400, 7500, 7700, 7800, 7900, 8000, 8100, 8200, 8300, 8400, 8500, 8600, 8700, 8900, 9100, 9200, 9300, 9400, 9500, 9600, 9700, 9800, 9900, 10000, 10100, 10200, 10300, 10400, 10500, 10600, 10700, 10800, 10900, 11000, 11100, 11200, 11300, 11400, 11500, 11600, 11900, 12000, 12001, 12100, 12200, 12300, 12400, 12600, 12700, 12800, 12900, 13000, 13100, 13200, 13300, 13400, 13500, 13600, 13800, 13900, 14000, 14100, 14400, 14700, 14800, 14900, 15000, 15100, 15200, 15300, 15400, 15500, 15600, 15700, 99999); -my %is_utf16le = map { $_ => 1 } qw (30 40 130 131 132 133 140 141 1000 1100 1430 1440 1441 1730 1740 1731 5500 5600 8000 9400 9500 9600 9700 9800 11600 13500 13800); +#my %is_utf16le = map { $_ => 1 } qw (30 40 130 131 132 133 140 141 1000 1100 1430 1440 1441 1730 1740 1731 5500 5600 8000 9400 9500 9600 9700 9800 11600 13500 13800); +my %is_utf16le = map { $_ => 1 } qw (30 40 130 131 132 133 140 141 1000 1100 1430 1440 1441 1730 1740 1731 2100 5500 5600 8000 9400 9500 9600 9700 9800 11600 13500 13800); my %less_fifteen = map { $_ => 1 } qw (500 1600 1800 2400 2410 3200 6300 7400 10500 10700); my %allow_long_salt = map { $_ => 1 } qw (2500 4520 4521 5500 5600 7100 7200 7300 9400 9500 9600 9700 9800 10400 10500 10600 10700 1100 11000 11200 11300 11400 11600 12600 13500 13800 15000); @@ -3339,9 +3340,9 @@ sub passthrough } elsif ($mode == 2100) { - next if length ($word_buf) > 13; + next if length ($word_buf) >= 256; - my $salt_len = get_random_num (1, 19); + my $salt_len = get_random_num (1, 256); $tmp_hash = gen_hash ($mode, $word_buf, substr ($salt_buf, 0, $salt_len)); } @@ -3878,9 +3879,9 @@ sub single } elsif ($mode == 2100) { - my $salt_len = get_random_num (1, 19); + my $salt_len = get_random_num (1, 256); - for (my $i = 1; $i < 13; $i++) + for (my $i = 1; $i < 256; $i++) { if ($len != 0) { @@ -8839,7 +8840,6 @@ sub dpapi_pbkdf2 return substr ($t, 0, $keylen); } - sub rnd { my $mode = shift;