From 30681e5151066af89943be1ee95d88c464931439 Mon Sep 17 00:00:00 2001 From: Jens Steube Date: Fri, 8 Mar 2019 11:14:33 +0100 Subject: [PATCH] Move remaining module specific typedefs and structures from inc_types.cl to kernels --- OpenCL/inc_types.cl | 9 +++++++++ OpenCL/m06400-pure.cl | 10 ++++++++++ OpenCL/m06500-pure.cl | 10 ++++++++++ OpenCL/m06600-pure.cl | 10 ++++++++++ OpenCL/m06700-pure.cl | 10 ++++++++++ OpenCL/m06800-pure.cl | 10 ++++++++++ OpenCL/m07100-pure.cl | 16 ++++++++++++++++ OpenCL/m07300_a0-optimized.cl | 7 +++++++ OpenCL/m07300_a0-pure.cl | 7 +++++++ OpenCL/m07300_a1-optimized.cl | 7 +++++++ OpenCL/m07300_a1-pure.cl | 7 +++++++ OpenCL/m07300_a3-optimized.cl | 7 +++++++ OpenCL/m07300_a3-pure.cl | 7 +++++++ OpenCL/m07400-optimized.cl | 10 ++++++++++ OpenCL/m07400-pure.cl | 10 ++++++++++ OpenCL/m07500_a0-optimized.cl | 10 ++++++++++ OpenCL/m07500_a0-pure.cl | 10 ++++++++++ OpenCL/m07500_a1-optimized.cl | 10 ++++++++++ OpenCL/m07500_a1-pure.cl | 10 ++++++++++ OpenCL/m07500_a3-optimized.cl | 10 ++++++++++ OpenCL/m07500_a3-pure.cl | 10 ++++++++++ OpenCL/m07900-pure.cl | 6 ++++++ OpenCL/m08200-pure.cl | 17 +++++++++++++++++ OpenCL/m08800-pure.cl | 16 ++++++++++++++++ OpenCL/m09000-pure.cl | 13 +++++++++++++ OpenCL/m09100-pure.cl | 10 ++++++++++ OpenCL/m09400-pure.cl | 15 +++++++++++++++ OpenCL/m09500-pure.cl | 13 +++++++++++++ OpenCL/m09600-pure.cl | 13 +++++++++++++ OpenCL/m09700_a0-optimized.cl | 9 +++++++++ OpenCL/m09700_a1-optimized.cl | 9 +++++++++ OpenCL/m09700_a3-optimized.cl | 9 +++++++++ OpenCL/m09710_a0-optimized.cl | 9 +++++++++ OpenCL/m09710_a1-optimized.cl | 9 +++++++++ OpenCL/m09710_a3-optimized.cl | 9 +++++++++ OpenCL/m09720_a0-optimized.cl | 9 +++++++++ OpenCL/m09720_a1-optimized.cl | 9 +++++++++ OpenCL/m09720_a3-optimized.cl | 9 +++++++++ OpenCL/m09800_a0-optimized.cl | 9 +++++++++ OpenCL/m09800_a1-optimized.cl | 9 +++++++++ OpenCL/m09800_a3-optimized.cl | 9 +++++++++ OpenCL/m09810_a0-optimized.cl | 9 +++++++++ OpenCL/m09810_a1-optimized.cl | 9 +++++++++ OpenCL/m09810_a3-optimized.cl | 9 +++++++++ OpenCL/m09820_a0-optimized.cl | 9 +++++++++ OpenCL/m09820_a1-optimized.cl | 9 +++++++++ OpenCL/m09820_a3-optimized.cl | 9 +++++++++ OpenCL/m10300-pure.cl | 6 ++++++ OpenCL/m10400_a0-optimized.cl | 21 +++++++++++++++++++++ OpenCL/m10400_a1-optimized.cl | 21 +++++++++++++++++++++ OpenCL/m10400_a3-optimized.cl | 21 +++++++++++++++++++++ OpenCL/m10410_a0-optimized.cl | 21 +++++++++++++++++++++ OpenCL/m10410_a1-optimized.cl | 21 +++++++++++++++++++++ OpenCL/m10410_a3-optimized.cl | 21 +++++++++++++++++++++ OpenCL/m10420_a0-optimized.cl | 21 +++++++++++++++++++++ OpenCL/m10420_a1-optimized.cl | 21 +++++++++++++++++++++ OpenCL/m10420_a3-optimized.cl | 21 +++++++++++++++++++++ OpenCL/m10500-pure.cl | 28 ++++++++++++++++++++++++++++ OpenCL/m10900-pure.cl | 16 ++++++++++++++++ OpenCL/m11300-pure.cl | 13 +++++++++++++ OpenCL/m11400_a0-pure.cl | 10 ++++++++++ OpenCL/m11400_a1-pure.cl | 10 ++++++++++ OpenCL/m11400_a3-pure.cl | 10 ++++++++++ OpenCL/m11600-pure.cl | 21 +++++++++++++++++++++ OpenCL/m11900-pure.cl | 16 ++++++++++++++++ OpenCL/m12000-pure.cl | 16 ++++++++++++++++ OpenCL/m12200-pure.cl | 6 ++++++ OpenCL/m12300-pure.cl | 10 ++++++++++ OpenCL/m12500-pure.cl | 12 ++++++++++++ OpenCL/m12700-pure.cl | 10 ++++++++++ OpenCL/m12800-pure.cl | 16 ++++++++++++++++ OpenCL/m12900-pure.cl | 16 ++++++++++++++++ OpenCL/m13000-pure.cl | 16 +++++++++++++--- OpenCL/m13100_a0-optimized.cl | 9 +++++++++ OpenCL/m13100_a0-pure.cl | 9 +++++++++ OpenCL/m13100_a1-optimized.cl | 9 +++++++++ OpenCL/m13100_a1-pure.cl | 9 +++++++++ OpenCL/m13100_a3-optimized.cl | 9 +++++++++ OpenCL/m13100_a3-pure.cl | 9 +++++++++ OpenCL/m13200-pure.cl | 8 ++++++++ OpenCL/m13400-pure.cl | 29 +++++++++++++++++++++++++++++ OpenCL/m13500_a0-optimized.cl | 10 ++++++++++ OpenCL/m13500_a0-pure.cl | 10 ++++++++++ OpenCL/m13500_a1-optimized.cl | 10 ++++++++++ OpenCL/m13500_a1-pure.cl | 10 ++++++++++ OpenCL/m13500_a3-optimized.cl | 10 ++++++++++ OpenCL/m13500_a3-pure.cl | 10 ++++++++++ OpenCL/m13600-pure.cl | 26 ++++++++++++++++++++++++++ OpenCL/m13751-pure.cl | 22 ++++++++++++++++++++++ OpenCL/m13752-pure.cl | 22 ++++++++++++++++++++++ OpenCL/m13753-pure.cl | 22 ++++++++++++++++++++++ OpenCL/m13771-pure.cl | 25 +++++++++++++++++++++++++ OpenCL/m13772-pure.cl | 25 +++++++++++++++++++++++++ OpenCL/m13773-pure.cl | 25 +++++++++++++++++++++++++ OpenCL/m13800_a0-optimized.cl | 6 ++++++ OpenCL/m13800_a0-pure.cl | 6 ++++++ OpenCL/m13800_a1-optimized.cl | 6 ++++++ OpenCL/m13800_a1-pure.cl | 6 ++++++ OpenCL/m13800_a3-optimized.cl | 6 ++++++ OpenCL/m13800_a3-pure.cl | 6 ++++++ 100 files changed, 1234 insertions(+), 3 deletions(-) diff --git a/OpenCL/inc_types.cl b/OpenCL/inc_types.cl index 0d883871b..99e1c8321 100644 --- a/OpenCL/inc_types.cl +++ b/OpenCL/inc_types.cl @@ -1329,3 +1329,12 @@ typedef struct plain u32 hash_pos; } plain_t; + +typedef struct keyboard_layout_mapping +{ + u32 src_char; + int src_len; + u32 dst_char; + int dst_len; + +} keyboard_layout_mapping_t; diff --git a/OpenCL/m06400-pure.cl b/OpenCL/m06400-pure.cl index a312bf990..fdcacad78 100644 --- a/OpenCL/m06400-pure.cl +++ b/OpenCL/m06400-pure.cl @@ -16,6 +16,16 @@ #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" +typedef struct sha256aix_tmp +{ + u32 ipad[8]; + u32 opad[8]; + + u32 dgst[8]; + u32 out[8]; + +} sha256aix_tmp_t; + DECLSPEC void hmac_sha256_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest) { digest[0] = ipad[0]; diff --git a/OpenCL/m06500-pure.cl b/OpenCL/m06500-pure.cl index acb70f006..3983ba181 100644 --- a/OpenCL/m06500-pure.cl +++ b/OpenCL/m06500-pure.cl @@ -16,6 +16,16 @@ #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" +typedef struct sha512aix_tmp +{ + u64 ipad[8]; + u64 opad[8]; + + u64 dgst[8]; + u64 out[8]; + +} sha512aix_tmp_t; + 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) { digest[0] = ipad[0]; diff --git a/OpenCL/m06600-pure.cl b/OpenCL/m06600-pure.cl index cee322339..41f6e3f8d 100644 --- a/OpenCL/m06600-pure.cl +++ b/OpenCL/m06600-pure.cl @@ -17,6 +17,16 @@ #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" +typedef struct agilekey_tmp +{ + u32 ipad[5]; + u32 opad[5]; + + u32 dgst[5]; + u32 out[5]; + +} agilekey_tmp_t; + DECLSPEC void hmac_sha1_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest) { digest[0] = ipad[0]; diff --git a/OpenCL/m06700-pure.cl b/OpenCL/m06700-pure.cl index 7a6fb736f..3cea8b579 100644 --- a/OpenCL/m06700-pure.cl +++ b/OpenCL/m06700-pure.cl @@ -16,6 +16,16 @@ #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" +typedef struct sha1aix_tmp +{ + u32 ipad[5]; + u32 opad[5]; + + u32 dgst[5]; + u32 out[5]; + +} sha1aix_tmp_t; + DECLSPEC void hmac_sha1_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest) { digest[0] = ipad[0]; diff --git a/OpenCL/m06800-pure.cl b/OpenCL/m06800-pure.cl index bb4fbbab2..eb23a5d00 100644 --- a/OpenCL/m06800-pure.cl +++ b/OpenCL/m06800-pure.cl @@ -17,6 +17,16 @@ #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" +typedef struct lastpass_tmp +{ + u32 ipad[8]; + u32 opad[8]; + + u32 dgst[8]; + u32 out[8]; + +} lastpass_tmp_t; + DECLSPEC void hmac_sha256_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest) { digest[0] = ipad[0]; diff --git a/OpenCL/m07100-pure.cl b/OpenCL/m07100-pure.cl index 3651b1dd0..acbd0e82a 100644 --- a/OpenCL/m07100-pure.cl +++ b/OpenCL/m07100-pure.cl @@ -16,6 +16,22 @@ #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" +typedef struct pbkdf2_sha512_tmp +{ + u64 ipad[8]; + u64 opad[8]; + + u64 dgst[16]; + u64 out[16]; + +} pbkdf2_sha512_tmp_t; + +typedef struct pbkdf2_sha512 +{ + u32 salt_buf[64]; + +} pbkdf2_sha512_t; + 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) { digest[0] = ipad[0]; diff --git a/OpenCL/m07300_a0-optimized.cl b/OpenCL/m07300_a0-optimized.cl index 82335c3a5..911889c09 100644 --- a/OpenCL/m07300_a0-optimized.cl +++ b/OpenCL/m07300_a0-optimized.cl @@ -15,6 +15,13 @@ #include "inc_simd.cl" #include "inc_hash_sha1.cl" +typedef struct rakp +{ + u32 salt_buf[128]; + u32 salt_len; + +} rakp_t; + DECLSPEC void hmac_sha1_pad (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad) { w0[0] = w0[0] ^ 0x36363636; diff --git a/OpenCL/m07300_a0-pure.cl b/OpenCL/m07300_a0-pure.cl index 7121f970e..346d8c199 100644 --- a/OpenCL/m07300_a0-pure.cl +++ b/OpenCL/m07300_a0-pure.cl @@ -15,6 +15,13 @@ #include "inc_scalar.cl" #include "inc_hash_sha1.cl" +typedef struct rakp +{ + u32 salt_buf[128]; + u32 salt_len; + +} rakp_t; + __kernel void m07300_mxx (KERN_ATTR_RULES_ESALT (rakp_t)) { /** diff --git a/OpenCL/m07300_a1-optimized.cl b/OpenCL/m07300_a1-optimized.cl index 6f096a277..28e18eab7 100644 --- a/OpenCL/m07300_a1-optimized.cl +++ b/OpenCL/m07300_a1-optimized.cl @@ -13,6 +13,13 @@ #include "inc_simd.cl" #include "inc_hash_sha1.cl" +typedef struct rakp +{ + u32 salt_buf[128]; + u32 salt_len; + +} rakp_t; + DECLSPEC void hmac_sha1_pad (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad) { w0[0] = w0[0] ^ 0x36363636; diff --git a/OpenCL/m07300_a1-pure.cl b/OpenCL/m07300_a1-pure.cl index 34e087878..1353e5290 100644 --- a/OpenCL/m07300_a1-pure.cl +++ b/OpenCL/m07300_a1-pure.cl @@ -13,6 +13,13 @@ #include "inc_scalar.cl" #include "inc_hash_sha1.cl" +typedef struct rakp +{ + u32 salt_buf[128]; + u32 salt_len; + +} rakp_t; + __kernel void m07300_mxx (KERN_ATTR_ESALT (rakp_t)) { /** diff --git a/OpenCL/m07300_a3-optimized.cl b/OpenCL/m07300_a3-optimized.cl index 0929992db..fad09a0a4 100644 --- a/OpenCL/m07300_a3-optimized.cl +++ b/OpenCL/m07300_a3-optimized.cl @@ -13,6 +13,13 @@ #include "inc_simd.cl" #include "inc_hash_sha1.cl" +typedef struct rakp +{ + u32 salt_buf[128]; + u32 salt_len; + +} rakp_t; + DECLSPEC void hmac_sha1_pad (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad) { w0[0] = w0[0] ^ 0x36363636; diff --git a/OpenCL/m07300_a3-pure.cl b/OpenCL/m07300_a3-pure.cl index 05525566e..ca16a4dcc 100644 --- a/OpenCL/m07300_a3-pure.cl +++ b/OpenCL/m07300_a3-pure.cl @@ -13,6 +13,13 @@ #include "inc_scalar.cl" #include "inc_hash_sha1.cl" +typedef struct rakp +{ + u32 salt_buf[128]; + u32 salt_len; + +} rakp_t; + __kernel void m07300_mxx (KERN_ATTR_VECTOR_ESALT (rakp_t)) { /** diff --git a/OpenCL/m07400-optimized.cl b/OpenCL/m07400-optimized.cl index 0ee8ad118..b5db1fc7e 100644 --- a/OpenCL/m07400-optimized.cl +++ b/OpenCL/m07400-optimized.cl @@ -13,6 +13,16 @@ #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" +typedef struct sha256crypt_tmp +{ + // pure version + + u32 alt_result[8]; + u32 p_bytes[64]; + u32 s_bytes[64]; + +} sha256crypt_tmp_t; + DECLSPEC void sha256_transform_transport (const u32 *w, u32 *digest) { u32 w0[4]; diff --git a/OpenCL/m07400-pure.cl b/OpenCL/m07400-pure.cl index 2935d79cb..87aa97322 100644 --- a/OpenCL/m07400-pure.cl +++ b/OpenCL/m07400-pure.cl @@ -13,6 +13,16 @@ #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" +typedef struct sha256crypt_tmp +{ + // pure version + + u32 alt_result[8]; + u32 p_bytes[64]; + u32 s_bytes[64]; + +} sha256crypt_tmp_t; + __kernel void m07400_init (KERN_ATTR_TMPS (sha256crypt_tmp_t)) { /** diff --git a/OpenCL/m07500_a0-optimized.cl b/OpenCL/m07500_a0-optimized.cl index f1a76fde8..d564fcc1f 100644 --- a/OpenCL/m07500_a0-optimized.cl +++ b/OpenCL/m07500_a0-optimized.cl @@ -17,6 +17,16 @@ #include "inc_hash_md4.cl" #include "inc_hash_md5.cl" +typedef struct krb5pa +{ + u32 user[16]; + u32 realm[16]; + u32 salt[32]; + u32 timestamp[16]; + u32 checksum[4]; + +} krb5pa_t; + typedef struct { u8 S[256]; diff --git a/OpenCL/m07500_a0-pure.cl b/OpenCL/m07500_a0-pure.cl index 73384904b..e48d78871 100644 --- a/OpenCL/m07500_a0-pure.cl +++ b/OpenCL/m07500_a0-pure.cl @@ -16,6 +16,16 @@ #include "inc_hash_md4.cl" #include "inc_hash_md5.cl" +typedef struct krb5pa +{ + u32 user[16]; + u32 realm[16]; + u32 salt[32]; + u32 timestamp[16]; + u32 checksum[4]; + +} krb5pa_t; + typedef struct { u8 S[256]; diff --git a/OpenCL/m07500_a1-optimized.cl b/OpenCL/m07500_a1-optimized.cl index e732afb77..55d753454 100644 --- a/OpenCL/m07500_a1-optimized.cl +++ b/OpenCL/m07500_a1-optimized.cl @@ -15,6 +15,16 @@ #include "inc_hash_md4.cl" #include "inc_hash_md5.cl" +typedef struct krb5pa +{ + u32 user[16]; + u32 realm[16]; + u32 salt[32]; + u32 timestamp[16]; + u32 checksum[4]; + +} krb5pa_t; + typedef struct { u8 S[256]; diff --git a/OpenCL/m07500_a1-pure.cl b/OpenCL/m07500_a1-pure.cl index 7109a9a59..df2b8c286 100644 --- a/OpenCL/m07500_a1-pure.cl +++ b/OpenCL/m07500_a1-pure.cl @@ -14,6 +14,16 @@ #include "inc_hash_md4.cl" #include "inc_hash_md5.cl" +typedef struct krb5pa +{ + u32 user[16]; + u32 realm[16]; + u32 salt[32]; + u32 timestamp[16]; + u32 checksum[4]; + +} krb5pa_t; + typedef struct { u8 S[256]; diff --git a/OpenCL/m07500_a3-optimized.cl b/OpenCL/m07500_a3-optimized.cl index 14cd3680b..99ea812be 100644 --- a/OpenCL/m07500_a3-optimized.cl +++ b/OpenCL/m07500_a3-optimized.cl @@ -15,6 +15,16 @@ #include "inc_hash_md4.cl" #include "inc_hash_md5.cl" +typedef struct krb5pa +{ + u32 user[16]; + u32 realm[16]; + u32 salt[32]; + u32 timestamp[16]; + u32 checksum[4]; + +} krb5pa_t; + typedef struct { u8 S[256]; diff --git a/OpenCL/m07500_a3-pure.cl b/OpenCL/m07500_a3-pure.cl index 4fd5b124f..98420a3eb 100644 --- a/OpenCL/m07500_a3-pure.cl +++ b/OpenCL/m07500_a3-pure.cl @@ -14,6 +14,16 @@ #include "inc_hash_md4.cl" #include "inc_hash_md5.cl" +typedef struct krb5pa +{ + u32 user[16]; + u32 realm[16]; + u32 salt[32]; + u32 timestamp[16]; + u32 checksum[4]; + +} krb5pa_t; + typedef struct { u8 S[256]; diff --git a/OpenCL/m07900-pure.cl b/OpenCL/m07900-pure.cl index 61b393c2c..fa25a1354 100644 --- a/OpenCL/m07900-pure.cl +++ b/OpenCL/m07900-pure.cl @@ -13,6 +13,12 @@ #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" +typedef struct drupal7_tmp +{ + u64 digest_buf[8]; + +} drupal7_tmp_t; + __kernel void m07900_init (KERN_ATTR_TMPS (drupal7_tmp_t)) { /** diff --git a/OpenCL/m08200-pure.cl b/OpenCL/m08200-pure.cl index 5be1abfc0..e405b885c 100644 --- a/OpenCL/m08200-pure.cl +++ b/OpenCL/m08200-pure.cl @@ -17,6 +17,23 @@ #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" +typedef struct pbkdf2_sha512_tmp +{ + u64 ipad[8]; + u64 opad[8]; + + u64 dgst[16]; + u64 out[16]; + +} pbkdf2_sha512_tmp_t; + +typedef struct cloudkey +{ + u32 data_len; + u32 data_buf[512]; + +} cloudkey_t; + 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) { digest[0] = ipad[0]; diff --git a/OpenCL/m08800-pure.cl b/OpenCL/m08800-pure.cl index dd5c35366..8a042a542 100644 --- a/OpenCL/m08800-pure.cl +++ b/OpenCL/m08800-pure.cl @@ -15,6 +15,22 @@ #include "inc_hash_sha256.cl" #include "inc_cipher_aes.cl" +typedef struct androidfde_tmp +{ + u32 ipad[5]; + u32 opad[5]; + + u32 dgst[10]; + u32 out[10]; + +} androidfde_tmp_t; + +typedef struct androidfde +{ + u32 data[384]; + +} androidfde_t; + DECLSPEC void hmac_sha1_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest) { digest[0] = ipad[0]; diff --git a/OpenCL/m09000-pure.cl b/OpenCL/m09000-pure.cl index 6f4496e48..19c518a45 100644 --- a/OpenCL/m09000-pure.cl +++ b/OpenCL/m09000-pure.cl @@ -12,6 +12,19 @@ #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" +typedef struct pwsafe2_tmp +{ + u32 digest[2]; + + u32 P[18]; + + u32 S0[256]; + u32 S1[256]; + u32 S2[256]; + u32 S3[256]; + +} pwsafe2_tmp_t; + // http://www.schneier.com/code/constants.txt __constant u32a c_sbox0[256] = diff --git a/OpenCL/m09100-pure.cl b/OpenCL/m09100-pure.cl index 58a05f2c5..7d286d448 100644 --- a/OpenCL/m09100-pure.cl +++ b/OpenCL/m09100-pure.cl @@ -16,6 +16,16 @@ #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" +typedef struct lotus8_tmp +{ + u32 ipad[5]; + u32 opad[5]; + + u32 dgst[5]; + u32 out[5]; + +} lotus8_tmp_t; + __constant u32 lotus64_table[64] = { '0', '1', '2', '3', '4', '5', '6', '7', diff --git a/OpenCL/m09400-pure.cl b/OpenCL/m09400-pure.cl index 6950da73b..196a087a3 100644 --- a/OpenCL/m09400-pure.cl +++ b/OpenCL/m09400-pure.cl @@ -17,6 +17,21 @@ #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" +typedef struct office2007_tmp +{ + u32 out[5]; + +} office2007_tmp_t; + +typedef struct office2007 +{ + u32 encryptedVerifier[4]; + u32 encryptedVerifierHash[5]; + + u32 keySize; + +} office2007_t; + __kernel void m09400_init (KERN_ATTR_TMPS_ESALT (office2007_tmp_t, office2007_t)) { /** diff --git a/OpenCL/m09500-pure.cl b/OpenCL/m09500-pure.cl index a90e973c2..f56e347de 100644 --- a/OpenCL/m09500-pure.cl +++ b/OpenCL/m09500-pure.cl @@ -17,6 +17,19 @@ #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" +typedef struct office2010 +{ + u32 encryptedVerifier[4]; + u32 encryptedVerifierHash[8]; + +} office2010_t; + +typedef struct office2010_tmp +{ + u32 out[5]; + +} office2010_tmp_t; + __kernel void m09500_init (KERN_ATTR_TMPS_ESALT (office2010_tmp_t, office2010_t)) { /** diff --git a/OpenCL/m09600-pure.cl b/OpenCL/m09600-pure.cl index 56840d548..9467ef360 100644 --- a/OpenCL/m09600-pure.cl +++ b/OpenCL/m09600-pure.cl @@ -17,6 +17,19 @@ #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" +typedef struct office2013_tmp +{ + u64 out[8]; + +} office2013_tmp_t; + +typedef struct office2013 +{ + u32 encryptedVerifier[4]; + u32 encryptedVerifierHash[8]; + +} office2013_t; + __kernel void m09600_init (KERN_ATTR_TMPS_ESALT (office2013_tmp_t, office2013_t)) { /** diff --git a/OpenCL/m09700_a0-optimized.cl b/OpenCL/m09700_a0-optimized.cl index 4e62f287f..6a6a3d81a 100644 --- a/OpenCL/m09700_a0-optimized.cl +++ b/OpenCL/m09700_a0-optimized.cl @@ -16,6 +16,15 @@ #include "inc_simd.cl" #include "inc_hash_md5.cl" +typedef struct oldoffice01 +{ + u32 version; + u32 encryptedVerifier[4]; + u32 encryptedVerifierHash[4]; + u32 rc4key[2]; + +} oldoffice01_t; + typedef struct { u8 S[256]; diff --git a/OpenCL/m09700_a1-optimized.cl b/OpenCL/m09700_a1-optimized.cl index 2fedbafdc..e8f7da1a8 100644 --- a/OpenCL/m09700_a1-optimized.cl +++ b/OpenCL/m09700_a1-optimized.cl @@ -14,6 +14,15 @@ #include "inc_simd.cl" #include "inc_hash_md5.cl" +typedef struct oldoffice01 +{ + u32 version; + u32 encryptedVerifier[4]; + u32 encryptedVerifierHash[4]; + u32 rc4key[2]; + +} oldoffice01_t; + typedef struct { u8 S[256]; diff --git a/OpenCL/m09700_a3-optimized.cl b/OpenCL/m09700_a3-optimized.cl index e136fe2d1..dd20e0d35 100644 --- a/OpenCL/m09700_a3-optimized.cl +++ b/OpenCL/m09700_a3-optimized.cl @@ -11,6 +11,15 @@ #include "inc_simd.cl" #include "inc_hash_md5.cl" +typedef struct oldoffice01 +{ + u32 version; + u32 encryptedVerifier[4]; + u32 encryptedVerifierHash[4]; + u32 rc4key[2]; + +} oldoffice01_t; + typedef struct { u8 S[256]; diff --git a/OpenCL/m09710_a0-optimized.cl b/OpenCL/m09710_a0-optimized.cl index 09fd289e1..8bd0a85fb 100644 --- a/OpenCL/m09710_a0-optimized.cl +++ b/OpenCL/m09710_a0-optimized.cl @@ -16,6 +16,15 @@ #include "inc_simd.cl" #include "inc_hash_md5.cl" +typedef struct oldoffice01 +{ + u32 version; + u32 encryptedVerifier[4]; + u32 encryptedVerifierHash[4]; + u32 rc4key[2]; + +} oldoffice01_t; + typedef struct { u8 S[256]; diff --git a/OpenCL/m09710_a1-optimized.cl b/OpenCL/m09710_a1-optimized.cl index 5cd5a923d..acfe98652 100644 --- a/OpenCL/m09710_a1-optimized.cl +++ b/OpenCL/m09710_a1-optimized.cl @@ -14,6 +14,15 @@ #include "inc_simd.cl" #include "inc_hash_md5.cl" +typedef struct oldoffice01 +{ + u32 version; + u32 encryptedVerifier[4]; + u32 encryptedVerifierHash[4]; + u32 rc4key[2]; + +} oldoffice01_t; + typedef struct { u8 S[256]; diff --git a/OpenCL/m09710_a3-optimized.cl b/OpenCL/m09710_a3-optimized.cl index 0dcc6ffd3..41cf52024 100644 --- a/OpenCL/m09710_a3-optimized.cl +++ b/OpenCL/m09710_a3-optimized.cl @@ -14,6 +14,15 @@ #include "inc_simd.cl" #include "inc_hash_md5.cl" +typedef struct oldoffice01 +{ + u32 version; + u32 encryptedVerifier[4]; + u32 encryptedVerifierHash[4]; + u32 rc4key[2]; + +} oldoffice01_t; + typedef struct { u8 S[256]; diff --git a/OpenCL/m09720_a0-optimized.cl b/OpenCL/m09720_a0-optimized.cl index b1dd35f30..698a9cf01 100644 --- a/OpenCL/m09720_a0-optimized.cl +++ b/OpenCL/m09720_a0-optimized.cl @@ -15,6 +15,15 @@ #include "inc_simd.cl" #include "inc_hash_md5.cl" +typedef struct oldoffice01 +{ + u32 version; + u32 encryptedVerifier[4]; + u32 encryptedVerifierHash[4]; + u32 rc4key[2]; + +} oldoffice01_t; + DECLSPEC void gen336 (u32x *digest_pre, u32 *salt_buf, u32x *digest) { u32x digest_t0[2]; diff --git a/OpenCL/m09720_a1-optimized.cl b/OpenCL/m09720_a1-optimized.cl index 956d9b5a6..b590a9486 100644 --- a/OpenCL/m09720_a1-optimized.cl +++ b/OpenCL/m09720_a1-optimized.cl @@ -13,6 +13,15 @@ #include "inc_simd.cl" #include "inc_hash_md5.cl" +typedef struct oldoffice01 +{ + u32 version; + u32 encryptedVerifier[4]; + u32 encryptedVerifierHash[4]; + u32 rc4key[2]; + +} oldoffice01_t; + DECLSPEC void gen336 (u32x *digest_pre, u32 *salt_buf, u32x *digest) { u32x digest_t0[2]; diff --git a/OpenCL/m09720_a3-optimized.cl b/OpenCL/m09720_a3-optimized.cl index b0cff0fc1..c6fad8678 100644 --- a/OpenCL/m09720_a3-optimized.cl +++ b/OpenCL/m09720_a3-optimized.cl @@ -13,6 +13,15 @@ #include "inc_simd.cl" #include "inc_hash_md5.cl" +typedef struct oldoffice01 +{ + u32 version; + u32 encryptedVerifier[4]; + u32 encryptedVerifierHash[4]; + u32 rc4key[2]; + +} oldoffice01_t; + DECLSPEC void gen336 (u32x *digest_pre, u32 *salt_buf, u32x *digest) { u32x digest_t0[2]; diff --git a/OpenCL/m09800_a0-optimized.cl b/OpenCL/m09800_a0-optimized.cl index 44c65b558..fd536bcdc 100644 --- a/OpenCL/m09800_a0-optimized.cl +++ b/OpenCL/m09800_a0-optimized.cl @@ -16,6 +16,15 @@ #include "inc_simd.cl" #include "inc_hash_sha1.cl" +typedef struct oldoffice34 +{ + u32 version; + u32 encryptedVerifier[4]; + u32 encryptedVerifierHash[5]; + u32 rc4key[2]; + +} oldoffice34_t; + typedef struct { u8 S[256]; diff --git a/OpenCL/m09800_a1-optimized.cl b/OpenCL/m09800_a1-optimized.cl index ba64ac196..5f553215e 100644 --- a/OpenCL/m09800_a1-optimized.cl +++ b/OpenCL/m09800_a1-optimized.cl @@ -14,6 +14,15 @@ #include "inc_simd.cl" #include "inc_hash_sha1.cl" +typedef struct oldoffice34 +{ + u32 version; + u32 encryptedVerifier[4]; + u32 encryptedVerifierHash[5]; + u32 rc4key[2]; + +} oldoffice34_t; + typedef struct { u8 S[256]; diff --git a/OpenCL/m09800_a3-optimized.cl b/OpenCL/m09800_a3-optimized.cl index 38d07fc7d..25d38db55 100644 --- a/OpenCL/m09800_a3-optimized.cl +++ b/OpenCL/m09800_a3-optimized.cl @@ -11,6 +11,15 @@ #include "inc_simd.cl" #include "inc_hash_sha1.cl" +typedef struct oldoffice34 +{ + u32 version; + u32 encryptedVerifier[4]; + u32 encryptedVerifierHash[5]; + u32 rc4key[2]; + +} oldoffice34_t; + typedef struct { u8 S[256]; diff --git a/OpenCL/m09810_a0-optimized.cl b/OpenCL/m09810_a0-optimized.cl index a90363f7a..2909a58e1 100644 --- a/OpenCL/m09810_a0-optimized.cl +++ b/OpenCL/m09810_a0-optimized.cl @@ -16,6 +16,15 @@ #include "inc_simd.cl" #include "inc_hash_sha1.cl" +typedef struct oldoffice34 +{ + u32 version; + u32 encryptedVerifier[4]; + u32 encryptedVerifierHash[5]; + u32 rc4key[2]; + +} oldoffice34_t; + typedef struct { u8 S[256]; diff --git a/OpenCL/m09810_a1-optimized.cl b/OpenCL/m09810_a1-optimized.cl index e0cce2e77..6a0e03626 100644 --- a/OpenCL/m09810_a1-optimized.cl +++ b/OpenCL/m09810_a1-optimized.cl @@ -14,6 +14,15 @@ #include "inc_simd.cl" #include "inc_hash_sha1.cl" +typedef struct oldoffice34 +{ + u32 version; + u32 encryptedVerifier[4]; + u32 encryptedVerifierHash[5]; + u32 rc4key[2]; + +} oldoffice34_t; + typedef struct { u8 S[256]; diff --git a/OpenCL/m09810_a3-optimized.cl b/OpenCL/m09810_a3-optimized.cl index 5607f7258..1a5c9ee73 100644 --- a/OpenCL/m09810_a3-optimized.cl +++ b/OpenCL/m09810_a3-optimized.cl @@ -14,6 +14,15 @@ #include "inc_simd.cl" #include "inc_hash_sha1.cl" +typedef struct oldoffice34 +{ + u32 version; + u32 encryptedVerifier[4]; + u32 encryptedVerifierHash[5]; + u32 rc4key[2]; + +} oldoffice34_t; + typedef struct { u8 S[256]; diff --git a/OpenCL/m09820_a0-optimized.cl b/OpenCL/m09820_a0-optimized.cl index 9c3f06cb3..d312a212c 100644 --- a/OpenCL/m09820_a0-optimized.cl +++ b/OpenCL/m09820_a0-optimized.cl @@ -15,6 +15,15 @@ #include "inc_simd.cl" #include "inc_hash_sha1.cl" +typedef struct oldoffice34 +{ + u32 version; + u32 encryptedVerifier[4]; + u32 encryptedVerifierHash[5]; + u32 rc4key[2]; + +} oldoffice34_t; + __kernel void m09820_m04 (KERN_ATTR_RULES_ESALT (oldoffice34_t)) { /** diff --git a/OpenCL/m09820_a1-optimized.cl b/OpenCL/m09820_a1-optimized.cl index cc07760a4..89bb2249e 100644 --- a/OpenCL/m09820_a1-optimized.cl +++ b/OpenCL/m09820_a1-optimized.cl @@ -13,6 +13,15 @@ #include "inc_simd.cl" #include "inc_hash_sha1.cl" +typedef struct oldoffice34 +{ + u32 version; + u32 encryptedVerifier[4]; + u32 encryptedVerifierHash[5]; + u32 rc4key[2]; + +} oldoffice34_t; + __kernel void m09820_m04 (KERN_ATTR_ESALT (oldoffice34_t)) { /** diff --git a/OpenCL/m09820_a3-optimized.cl b/OpenCL/m09820_a3-optimized.cl index efa6ad967..5c23d43b4 100644 --- a/OpenCL/m09820_a3-optimized.cl +++ b/OpenCL/m09820_a3-optimized.cl @@ -13,6 +13,15 @@ #include "inc_simd.cl" #include "inc_hash_sha1.cl" +typedef struct oldoffice34 +{ + u32 version; + u32 encryptedVerifier[4]; + u32 encryptedVerifierHash[5]; + u32 rc4key[2]; + +} oldoffice34_t; + DECLSPEC void m09820m (u32 *w0, u32 *w1, u32 *w2, u32 *w3, const u32 pw_len, KERN_ATTR_ESALT (oldoffice34_t)) { /** diff --git a/OpenCL/m10300-pure.cl b/OpenCL/m10300-pure.cl index 6541ce2b2..cf956e741 100644 --- a/OpenCL/m10300-pure.cl +++ b/OpenCL/m10300-pure.cl @@ -13,6 +13,12 @@ #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" +typedef struct saph_sha1_tmp +{ + u32 digest_buf[5]; + +} saph_sha1_tmp_t; + __kernel void m10300_init (KERN_ATTR_TMPS (saph_sha1_tmp_t)) { /** diff --git a/OpenCL/m10400_a0-optimized.cl b/OpenCL/m10400_a0-optimized.cl index f74f233e8..95e999e80 100644 --- a/OpenCL/m10400_a0-optimized.cl +++ b/OpenCL/m10400_a0-optimized.cl @@ -28,6 +28,27 @@ __constant u32a padding[8] = 0x7a695364 }; +typedef struct pdf +{ + int V; + int R; + int P; + + int enc_md; + + u32 id_buf[8]; + u32 u_buf[32]; + u32 o_buf[32]; + + int id_len; + int o_len; + int u_len; + + u32 rc4key[2]; + u32 rc4data[2]; + +} pdf_t; + typedef struct { u8 S[256]; diff --git a/OpenCL/m10400_a1-optimized.cl b/OpenCL/m10400_a1-optimized.cl index 2b843c695..84be38b4c 100644 --- a/OpenCL/m10400_a1-optimized.cl +++ b/OpenCL/m10400_a1-optimized.cl @@ -26,6 +26,27 @@ __constant u32a padding[8] = 0x7a695364 }; +typedef struct pdf +{ + int V; + int R; + int P; + + int enc_md; + + u32 id_buf[8]; + u32 u_buf[32]; + u32 o_buf[32]; + + int id_len; + int o_len; + int u_len; + + u32 rc4key[2]; + u32 rc4data[2]; + +} pdf_t; + typedef struct { u8 S[256]; diff --git a/OpenCL/m10400_a3-optimized.cl b/OpenCL/m10400_a3-optimized.cl index a625d7534..0035c0cf9 100644 --- a/OpenCL/m10400_a3-optimized.cl +++ b/OpenCL/m10400_a3-optimized.cl @@ -26,6 +26,27 @@ __constant u32a padding[8] = 0x7a695364 }; +typedef struct pdf +{ + int V; + int R; + int P; + + int enc_md; + + u32 id_buf[8]; + u32 u_buf[32]; + u32 o_buf[32]; + + int id_len; + int o_len; + int u_len; + + u32 rc4key[2]; + u32 rc4data[2]; + +} pdf_t; + typedef struct { u8 S[256]; diff --git a/OpenCL/m10410_a0-optimized.cl b/OpenCL/m10410_a0-optimized.cl index d50247790..afd5f61e4 100644 --- a/OpenCL/m10410_a0-optimized.cl +++ b/OpenCL/m10410_a0-optimized.cl @@ -28,6 +28,27 @@ __constant u32a padding[8] = 0x7a695364 }; +typedef struct pdf +{ + int V; + int R; + int P; + + int enc_md; + + u32 id_buf[8]; + u32 u_buf[32]; + u32 o_buf[32]; + + int id_len; + int o_len; + int u_len; + + u32 rc4key[2]; + u32 rc4data[2]; + +} pdf_t; + typedef struct { u8 S[256]; diff --git a/OpenCL/m10410_a1-optimized.cl b/OpenCL/m10410_a1-optimized.cl index a5d5a6ed7..09abe207f 100644 --- a/OpenCL/m10410_a1-optimized.cl +++ b/OpenCL/m10410_a1-optimized.cl @@ -26,6 +26,27 @@ __constant u32a padding[8] = 0x7a695364 }; +typedef struct pdf +{ + int V; + int R; + int P; + + int enc_md; + + u32 id_buf[8]; + u32 u_buf[32]; + u32 o_buf[32]; + + int id_len; + int o_len; + int u_len; + + u32 rc4key[2]; + u32 rc4data[2]; + +} pdf_t; + typedef struct { u8 S[256]; diff --git a/OpenCL/m10410_a3-optimized.cl b/OpenCL/m10410_a3-optimized.cl index 13e7de4a8..cad024dfa 100644 --- a/OpenCL/m10410_a3-optimized.cl +++ b/OpenCL/m10410_a3-optimized.cl @@ -26,6 +26,27 @@ __constant u32a padding[8] = 0x7a695364 }; +typedef struct pdf +{ + int V; + int R; + int P; + + int enc_md; + + u32 id_buf[8]; + u32 u_buf[32]; + u32 o_buf[32]; + + int id_len; + int o_len; + int u_len; + + u32 rc4key[2]; + u32 rc4data[2]; + +} pdf_t; + typedef struct { u8 S[256]; diff --git a/OpenCL/m10420_a0-optimized.cl b/OpenCL/m10420_a0-optimized.cl index 38ef96a0a..0174fb91a 100644 --- a/OpenCL/m10420_a0-optimized.cl +++ b/OpenCL/m10420_a0-optimized.cl @@ -27,6 +27,27 @@ __constant u32a padding[8] = 0x7a695364 }; +typedef struct pdf +{ + int V; + int R; + int P; + + int enc_md; + + u32 id_buf[8]; + u32 u_buf[32]; + u32 o_buf[32]; + + int id_len; + int o_len; + int u_len; + + u32 rc4key[2]; + u32 rc4data[2]; + +} pdf_t; + __kernel void m10420_m04 (KERN_ATTR_RULES_ESALT (pdf_t)) { /** diff --git a/OpenCL/m10420_a1-optimized.cl b/OpenCL/m10420_a1-optimized.cl index 853135537..70fffdb6b 100644 --- a/OpenCL/m10420_a1-optimized.cl +++ b/OpenCL/m10420_a1-optimized.cl @@ -25,6 +25,27 @@ __constant u32a padding[8] = 0x7a695364 }; +typedef struct pdf +{ + int V; + int R; + int P; + + int enc_md; + + u32 id_buf[8]; + u32 u_buf[32]; + u32 o_buf[32]; + + int id_len; + int o_len; + int u_len; + + u32 rc4key[2]; + u32 rc4data[2]; + +} pdf_t; + __kernel void m10420_m04 (KERN_ATTR_ESALT (pdf_t)) { /** diff --git a/OpenCL/m10420_a3-optimized.cl b/OpenCL/m10420_a3-optimized.cl index df61ce41d..847197d20 100644 --- a/OpenCL/m10420_a3-optimized.cl +++ b/OpenCL/m10420_a3-optimized.cl @@ -25,6 +25,27 @@ __constant u32a padding[8] = 0x7a695364 }; +typedef struct pdf +{ + int V; + int R; + int P; + + int enc_md; + + u32 id_buf[8]; + u32 u_buf[32]; + u32 o_buf[32]; + + int id_len; + int o_len; + int u_len; + + u32 rc4key[2]; + u32 rc4data[2]; + +} pdf_t; + DECLSPEC void m10420m (u32 *w0, u32 *w1, u32 *w2, u32 *w3, const u32 pw_len, KERN_ATTR_ESALT (pdf_t)) { /** diff --git a/OpenCL/m10500-pure.cl b/OpenCL/m10500-pure.cl index 4669f65ea..18a59f65b 100644 --- a/OpenCL/m10500-pure.cl +++ b/OpenCL/m10500-pure.cl @@ -25,6 +25,34 @@ __constant u32a padding[8] = 0x7a695364 }; +typedef struct pdf +{ + int V; + int R; + int P; + + int enc_md; + + u32 id_buf[8]; + u32 u_buf[32]; + u32 o_buf[32]; + + int id_len; + int o_len; + int u_len; + + u32 rc4key[2]; + u32 rc4data[2]; + +} pdf_t; + +typedef struct pdf14_tmp +{ + u32 digest[4]; + u32 out[4]; + +} pdf14_tmp_t; + typedef struct { u8 S[256]; diff --git a/OpenCL/m10900-pure.cl b/OpenCL/m10900-pure.cl index c228d3403..15b1763c3 100644 --- a/OpenCL/m10900-pure.cl +++ b/OpenCL/m10900-pure.cl @@ -16,6 +16,22 @@ #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" +typedef struct pbkdf2_sha256_tmp +{ + u32 ipad[8]; + u32 opad[8]; + + u32 dgst[32]; + u32 out[32]; + +} pbkdf2_sha256_tmp_t; + +typedef struct pbkdf2_sha256 +{ + u32 salt_buf[64]; + +} pbkdf2_sha256_t; + DECLSPEC void hmac_sha256_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest) { digest[0] = ipad[0]; diff --git a/OpenCL/m11300-pure.cl b/OpenCL/m11300-pure.cl index f0e59d4a8..309ec370a 100644 --- a/OpenCL/m11300-pure.cl +++ b/OpenCL/m11300-pure.cl @@ -14,6 +14,19 @@ #include "inc_hash_sha512.cl" #include "inc_cipher_aes.cl" +typedef struct bitcoin_wallet_tmp +{ + u64 dgst[8]; + +} bitcoin_wallet_tmp_t; + +typedef struct bitcoin_wallet +{ + u32 cry_master_buf[64]; + u32 cry_master_len; + +} bitcoin_wallet_t; + 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) { digest[0] = ipad[0]; diff --git a/OpenCL/m11400_a0-pure.cl b/OpenCL/m11400_a0-pure.cl index 4e72a4936..d267e5b0f 100644 --- a/OpenCL/m11400_a0-pure.cl +++ b/OpenCL/m11400_a0-pure.cl @@ -15,6 +15,16 @@ #include "inc_scalar.cl" #include "inc_hash_md5.cl" +typedef struct sip +{ + u32 salt_buf[32]; + u32 salt_len; + + u32 esalt_buf[256]; + u32 esalt_len; + +} sip_t; + #if VECT_SIZE == 1 #define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i)]) #elif VECT_SIZE == 2 diff --git a/OpenCL/m11400_a1-pure.cl b/OpenCL/m11400_a1-pure.cl index b73f363fc..b408d2ec7 100644 --- a/OpenCL/m11400_a1-pure.cl +++ b/OpenCL/m11400_a1-pure.cl @@ -13,6 +13,16 @@ #include "inc_scalar.cl" #include "inc_hash_md5.cl" +typedef struct sip +{ + u32 salt_buf[32]; + u32 salt_len; + + u32 esalt_buf[256]; + u32 esalt_len; + +} sip_t; + #if VECT_SIZE == 1 #define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i)]) #elif VECT_SIZE == 2 diff --git a/OpenCL/m11400_a3-pure.cl b/OpenCL/m11400_a3-pure.cl index 75332bcc7..6c2e48797 100644 --- a/OpenCL/m11400_a3-pure.cl +++ b/OpenCL/m11400_a3-pure.cl @@ -13,6 +13,16 @@ #include "inc_simd.cl" #include "inc_hash_md5.cl" +typedef struct sip +{ + u32 salt_buf[32]; + u32 salt_len; + + u32 esalt_buf[256]; + u32 esalt_len; + +} sip_t; + #if VECT_SIZE == 1 #define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i)]) #elif VECT_SIZE == 2 diff --git a/OpenCL/m11600-pure.cl b/OpenCL/m11600-pure.cl index 7a762ab09..71e8a41c2 100644 --- a/OpenCL/m11600-pure.cl +++ b/OpenCL/m11600-pure.cl @@ -10,6 +10,27 @@ #include "inc_common.cl" #include "inc_hash_sha256.cl" +typedef struct seven_zip_tmp +{ + u32 h[8]; + + u32 w0[4]; + u32 w1[4]; + u32 w2[4]; + u32 w3[4]; + + int len; + +} seven_zip_tmp_t; + +typedef struct +{ + u32 ukey[8]; + + u32 hook_success; + +} seven_zip_hook_t; + DECLSPEC void memcat8c_be (u32 *w0, u32 *w1, u32 *w2, u32 *w3, const u32 len, const u32 append, u32 *digest) { const u32 func_len = len & 63; diff --git a/OpenCL/m11900-pure.cl b/OpenCL/m11900-pure.cl index 12e00ed2f..7fbe43814 100644 --- a/OpenCL/m11900-pure.cl +++ b/OpenCL/m11900-pure.cl @@ -16,6 +16,22 @@ #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" +typedef struct pbkdf2_md5 +{ + u32 salt_buf[64]; + +} pbkdf2_md5_t; + +typedef struct pbkdf2_md5_tmp +{ + u32 ipad[4]; + u32 opad[4]; + + u32 dgst[32]; + u32 out[32]; + +} pbkdf2_md5_tmp_t; + DECLSPEC void hmac_md5_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest) { digest[0] = ipad[0]; diff --git a/OpenCL/m12000-pure.cl b/OpenCL/m12000-pure.cl index e9ac20e59..e89a9d10e 100644 --- a/OpenCL/m12000-pure.cl +++ b/OpenCL/m12000-pure.cl @@ -16,6 +16,22 @@ #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" +typedef struct pbkdf2_sha1_tmp +{ + u32 ipad[5]; + u32 opad[5]; + + u32 dgst[32]; + u32 out[32]; + +} pbkdf2_sha1_tmp_t; + +typedef struct pbkdf2_sha1 +{ + u32 salt_buf[64]; + +} pbkdf2_sha1_t; + DECLSPEC void hmac_sha1_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest) { digest[0] = ipad[0]; diff --git a/OpenCL/m12200-pure.cl b/OpenCL/m12200-pure.cl index 2f8ddc03f..70578ea08 100644 --- a/OpenCL/m12200-pure.cl +++ b/OpenCL/m12200-pure.cl @@ -16,6 +16,12 @@ #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" +typedef struct ecryptfs_tmp +{ + u64 out[8]; + +} ecryptfs_tmp_t; + __kernel void m12200_init (KERN_ATTR_TMPS (ecryptfs_tmp_t)) { /** diff --git a/OpenCL/m12300-pure.cl b/OpenCL/m12300-pure.cl index 0dddcacf5..e68e9970e 100644 --- a/OpenCL/m12300-pure.cl +++ b/OpenCL/m12300-pure.cl @@ -16,6 +16,16 @@ #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" +typedef struct oraclet_tmp +{ + u64 ipad[8]; + u64 opad[8]; + + u64 dgst[16]; + u64 out[16]; + +} oraclet_tmp_t; + 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) { digest[0] = ipad[0]; diff --git a/OpenCL/m12500-pure.cl b/OpenCL/m12500-pure.cl index ae1f9a7f8..1ba2ab6d5 100644 --- a/OpenCL/m12500-pure.cl +++ b/OpenCL/m12500-pure.cl @@ -23,6 +23,18 @@ #define MIN(a,b) (((a) < (b)) ? (a) : (b)) +typedef struct pbkdf2_sha1 +{ + u32 salt_buf[64]; + +} pbkdf2_sha1_t; + +typedef struct rar3_tmp +{ + u32 dgst[17][5]; + +} rar3_tmp_t; + DECLSPEC void sha1_transform (const u32 *w, u32 *digest) { u32 A = digest[0]; diff --git a/OpenCL/m12700-pure.cl b/OpenCL/m12700-pure.cl index 1fde0a42a..6c0ebbc11 100644 --- a/OpenCL/m12700-pure.cl +++ b/OpenCL/m12700-pure.cl @@ -17,6 +17,16 @@ #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" +typedef struct mywallet_tmp +{ + u32 ipad[5]; + u32 opad[5]; + + u32 dgst[10]; + u32 out[10]; + +} mywallet_tmp_t; + DECLSPEC void hmac_sha1_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest) { digest[0] = ipad[0]; diff --git a/OpenCL/m12800-pure.cl b/OpenCL/m12800-pure.cl index 8d472aabf..737da4093 100644 --- a/OpenCL/m12800-pure.cl +++ b/OpenCL/m12800-pure.cl @@ -17,6 +17,22 @@ #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" +typedef struct pbkdf2_sha256_tmp +{ + u32 ipad[8]; + u32 opad[8]; + + u32 dgst[32]; + u32 out[32]; + +} pbkdf2_sha256_tmp_t; + +typedef struct pbkdf2_sha256 +{ + u32 salt_buf[64]; + +} pbkdf2_sha256_t; + DECLSPEC void hmac_sha256_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest) { digest[0] = ipad[0]; diff --git a/OpenCL/m12900-pure.cl b/OpenCL/m12900-pure.cl index a6999ceab..19eaac2fd 100644 --- a/OpenCL/m12900-pure.cl +++ b/OpenCL/m12900-pure.cl @@ -16,6 +16,22 @@ #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" +typedef struct pbkdf2_sha256_tmp +{ + u32 ipad[8]; + u32 opad[8]; + + u32 dgst[32]; + u32 out[32]; + +} pbkdf2_sha256_tmp_t; + +typedef struct pbkdf2_sha256 +{ + u32 salt_buf[64]; + +} pbkdf2_sha256_t; + DECLSPEC void hmac_sha256_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest) { digest[0] = ipad[0]; diff --git a/OpenCL/m13000-pure.cl b/OpenCL/m13000-pure.cl index b371def0c..b9c6fd4bf 100644 --- a/OpenCL/m13000-pure.cl +++ b/OpenCL/m13000-pure.cl @@ -16,6 +16,16 @@ #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" +typedef struct pbkdf2_sha256_tmp +{ + u32 ipad[8]; + u32 opad[8]; + + u32 dgst[32]; + u32 out[32]; + +} pbkdf2_sha256_tmp_t; + DECLSPEC void hmac_sha256_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest) { digest[0] = ipad[0]; @@ -58,7 +68,7 @@ DECLSPEC void hmac_sha256_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *i sha256_transform_vector (w0, w1, w2, w3, digest); } -__kernel void m13000_init (KERN_ATTR_TMPS_ESALT (pbkdf2_sha256_tmp_t, pbkdf2_sha256_t)) +__kernel void m13000_init (KERN_ATTR_TMPS (pbkdf2_sha256_tmp_t)) { /** * base @@ -142,7 +152,7 @@ __kernel void m13000_init (KERN_ATTR_TMPS_ESALT (pbkdf2_sha256_tmp_t, pbkdf2_sha } } -__kernel void m13000_loop (KERN_ATTR_TMPS_ESALT (pbkdf2_sha256_tmp_t, pbkdf2_sha256_t)) +__kernel void m13000_loop (KERN_ATTR_TMPS (pbkdf2_sha256_tmp_t)) { const u64 gid = get_global_id (0); @@ -248,7 +258,7 @@ __kernel void m13000_loop (KERN_ATTR_TMPS_ESALT (pbkdf2_sha256_tmp_t, pbkdf2_sha } } -__kernel void m13000_comp (KERN_ATTR_TMPS_ESALT (pbkdf2_sha256_tmp_t, pbkdf2_sha256_t)) +__kernel void m13000_comp (KERN_ATTR_TMPS (pbkdf2_sha256_tmp_t)) { /** * base diff --git a/OpenCL/m13100_a0-optimized.cl b/OpenCL/m13100_a0-optimized.cl index e61092d08..c43041314 100644 --- a/OpenCL/m13100_a0-optimized.cl +++ b/OpenCL/m13100_a0-optimized.cl @@ -17,6 +17,15 @@ #include "inc_hash_md4.cl" #include "inc_hash_md5.cl" +typedef struct krb5tgs +{ + u32 account_info[512]; + u32 checksum[4]; + u32 edata2[5120]; + u32 edata2_len; + +} krb5tgs_t; + typedef struct { u8 S[256]; diff --git a/OpenCL/m13100_a0-pure.cl b/OpenCL/m13100_a0-pure.cl index 470d580dd..5a9a71aa2 100644 --- a/OpenCL/m13100_a0-pure.cl +++ b/OpenCL/m13100_a0-pure.cl @@ -16,6 +16,15 @@ #include "inc_hash_md4.cl" #include "inc_hash_md5.cl" +typedef struct krb5tgs +{ + u32 account_info[512]; + u32 checksum[4]; + u32 edata2[5120]; + u32 edata2_len; + +} krb5tgs_t; + typedef struct { u8 S[256]; diff --git a/OpenCL/m13100_a1-optimized.cl b/OpenCL/m13100_a1-optimized.cl index 0e257946b..3e60501e5 100644 --- a/OpenCL/m13100_a1-optimized.cl +++ b/OpenCL/m13100_a1-optimized.cl @@ -15,6 +15,15 @@ #include "inc_hash_md4.cl" #include "inc_hash_md5.cl" +typedef struct krb5tgs +{ + u32 account_info[512]; + u32 checksum[4]; + u32 edata2[5120]; + u32 edata2_len; + +} krb5tgs_t; + typedef struct { u8 S[256]; diff --git a/OpenCL/m13100_a1-pure.cl b/OpenCL/m13100_a1-pure.cl index 5e0988168..36748313f 100644 --- a/OpenCL/m13100_a1-pure.cl +++ b/OpenCL/m13100_a1-pure.cl @@ -14,6 +14,15 @@ #include "inc_hash_md4.cl" #include "inc_hash_md5.cl" +typedef struct krb5tgs +{ + u32 account_info[512]; + u32 checksum[4]; + u32 edata2[5120]; + u32 edata2_len; + +} krb5tgs_t; + typedef struct { u8 S[256]; diff --git a/OpenCL/m13100_a3-optimized.cl b/OpenCL/m13100_a3-optimized.cl index 7c09680a5..ffc10e99e 100644 --- a/OpenCL/m13100_a3-optimized.cl +++ b/OpenCL/m13100_a3-optimized.cl @@ -15,6 +15,15 @@ #include "inc_hash_md4.cl" #include "inc_hash_md5.cl" +typedef struct krb5tgs +{ + u32 account_info[512]; + u32 checksum[4]; + u32 edata2[5120]; + u32 edata2_len; + +} krb5tgs_t; + typedef struct { u8 S[256]; diff --git a/OpenCL/m13100_a3-pure.cl b/OpenCL/m13100_a3-pure.cl index 114f11017..f08dd2b89 100644 --- a/OpenCL/m13100_a3-pure.cl +++ b/OpenCL/m13100_a3-pure.cl @@ -14,6 +14,15 @@ #include "inc_hash_md4.cl" #include "inc_hash_md5.cl" +typedef struct krb5tgs +{ + u32 account_info[512]; + u32 checksum[4]; + u32 edata2[5120]; + u32 edata2_len; + +} krb5tgs_t; + typedef struct { u8 S[256]; diff --git a/OpenCL/m13200-pure.cl b/OpenCL/m13200-pure.cl index d85dd5e05..066f5aeaa 100644 --- a/OpenCL/m13200-pure.cl +++ b/OpenCL/m13200-pure.cl @@ -11,6 +11,14 @@ #include "inc_hash_sha1.cl" #include "inc_cipher_aes.cl" +typedef struct axcrypt_tmp +{ + u32 KEK[4]; + u32 lsb[4]; + u32 cipher[4]; + +} axcrypt_tmp_t; + __kernel void m13200_init (KERN_ATTR_TMPS (axcrypt_tmp_t)) { /** diff --git a/OpenCL/m13400-pure.cl b/OpenCL/m13400-pure.cl index 9dc28e6f5..8fa9da925 100644 --- a/OpenCL/m13400-pure.cl +++ b/OpenCL/m13400-pure.cl @@ -15,6 +15,35 @@ #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" +typedef struct keepass_tmp +{ + u32 tmp_digest[8]; + +} keepass_tmp_t; + +typedef struct keepass +{ + u32 version; + u32 algorithm; + + /* key-file handling */ + u32 keyfile_len; + u32 keyfile[8]; + + u32 final_random_seed[8]; + u32 transf_random_seed[8]; + u32 enc_iv[4]; + u32 contents_hash[8]; + + /* specific to version 1 */ + u32 contents_len; + u32 contents[75000]; + + /* specific to version 2 */ + u32 expected_bytes[8]; + +} keepass_t; + __kernel void m13400_init (KERN_ATTR_TMPS_ESALT (keepass_tmp_t, keepass_t)) { /** diff --git a/OpenCL/m13500_a0-optimized.cl b/OpenCL/m13500_a0-optimized.cl index 0c553ec56..20feceb7c 100644 --- a/OpenCL/m13500_a0-optimized.cl +++ b/OpenCL/m13500_a0-optimized.cl @@ -15,6 +15,16 @@ #include "inc_rp_optimized.cl" #include "inc_simd.cl" +typedef struct pstoken +{ + u32 salt_buf[128]; + u32 salt_len; + + u32 pc_digest[5]; + u32 pc_offset; + +} pstoken_t; + __kernel void m13500_m04 (KERN_ATTR_RULES_ESALT (pstoken_t)) { /** diff --git a/OpenCL/m13500_a0-pure.cl b/OpenCL/m13500_a0-pure.cl index 09311b92d..70d11b517 100644 --- a/OpenCL/m13500_a0-pure.cl +++ b/OpenCL/m13500_a0-pure.cl @@ -15,6 +15,16 @@ #include "inc_scalar.cl" #include "inc_hash_sha1.cl" +typedef struct pstoken +{ + u32 salt_buf[128]; + u32 salt_len; + + u32 pc_digest[5]; + u32 pc_offset; + +} pstoken_t; + __kernel void m13500_mxx (KERN_ATTR_RULES_ESALT (pstoken_t)) { /** diff --git a/OpenCL/m13500_a1-optimized.cl b/OpenCL/m13500_a1-optimized.cl index 453392b4e..f9defa844 100644 --- a/OpenCL/m13500_a1-optimized.cl +++ b/OpenCL/m13500_a1-optimized.cl @@ -13,6 +13,16 @@ #include "inc_common.cl" #include "inc_simd.cl" +typedef struct pstoken +{ + u32 salt_buf[128]; + u32 salt_len; + + u32 pc_digest[5]; + u32 pc_offset; + +} pstoken_t; + __kernel void m13500_m04 (KERN_ATTR_ESALT (pstoken_t)) { /** diff --git a/OpenCL/m13500_a1-pure.cl b/OpenCL/m13500_a1-pure.cl index 526609e0e..25096e242 100644 --- a/OpenCL/m13500_a1-pure.cl +++ b/OpenCL/m13500_a1-pure.cl @@ -13,6 +13,16 @@ #include "inc_scalar.cl" #include "inc_hash_sha1.cl" +typedef struct pstoken +{ + u32 salt_buf[128]; + u32 salt_len; + + u32 pc_digest[5]; + u32 pc_offset; + +} pstoken_t; + __kernel void m13500_mxx (KERN_ATTR_ESALT (pstoken_t)) { /** diff --git a/OpenCL/m13500_a3-optimized.cl b/OpenCL/m13500_a3-optimized.cl index db2035b51..74f41133c 100644 --- a/OpenCL/m13500_a3-optimized.cl +++ b/OpenCL/m13500_a3-optimized.cl @@ -13,6 +13,16 @@ #include "inc_common.cl" #include "inc_simd.cl" +typedef struct pstoken +{ + u32 salt_buf[128]; + u32 salt_len; + + u32 pc_digest[5]; + u32 pc_offset; + +} pstoken_t; + DECLSPEC void m13500m (u32 *w0, u32 *w1, u32 *w2, u32 *w3, const u32 pw_len, KERN_ATTR_ESALT (pstoken_t)) { /** diff --git a/OpenCL/m13500_a3-pure.cl b/OpenCL/m13500_a3-pure.cl index 26af0447a..9e594244c 100644 --- a/OpenCL/m13500_a3-pure.cl +++ b/OpenCL/m13500_a3-pure.cl @@ -13,6 +13,16 @@ #include "inc_simd.cl" #include "inc_hash_sha1.cl" +typedef struct pstoken +{ + u32 salt_buf[128]; + u32 salt_len; + + u32 pc_digest[5]; + u32 pc_offset; + +} pstoken_t; + __kernel void m13500_mxx (KERN_ATTR_VECTOR_ESALT (pstoken_t)) { /** diff --git a/OpenCL/m13600-pure.cl b/OpenCL/m13600-pure.cl index cb4f967d0..eb07e44a1 100644 --- a/OpenCL/m13600-pure.cl +++ b/OpenCL/m13600-pure.cl @@ -16,6 +16,32 @@ #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" +typedef struct pbkdf2_sha1_tmp +{ + u32 ipad[5]; + u32 opad[5]; + + u32 dgst[32]; + u32 out[32]; + +} pbkdf2_sha1_tmp_t; + +typedef struct zip2 +{ + u32 type; + u32 mode; + u32 magic; + u32 salt_len; + u32 salt_buf[4]; + u32 verify_bytes; + u32 compress_length; + u32 data_len; + u32 data_buf[2048]; + u32 auth_len; + u32 auth_buf[4]; + +} zip2_t; + DECLSPEC void hmac_sha1_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest) { digest[0] = ipad[0]; diff --git a/OpenCL/m13751-pure.cl b/OpenCL/m13751-pure.cl index 96c4ee0f9..761e79a72 100644 --- a/OpenCL/m13751-pure.cl +++ b/OpenCL/m13751-pure.cl @@ -24,6 +24,28 @@ #include "inc_truecrypt_xts.cl" #include "inc_veracrypt_xts.cl" +typedef struct tc +{ + u32 salt_buf[32]; + u32 data_buf[112]; + u32 keyfile_buf[16]; + u32 signature; + + keyboard_layout_mapping_t keyboard_layout_mapping_buf[256]; + int keyboard_layout_mapping_cnt; + +} tc_t; + +typedef struct tc_tmp +{ + u32 ipad[16]; + u32 opad[16]; + + u32 dgst[64]; + u32 out[64]; + +} tc_tmp_t; + DECLSPEC void hmac_sha256_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest) { digest[0] = ipad[0]; diff --git a/OpenCL/m13752-pure.cl b/OpenCL/m13752-pure.cl index 9c251821b..4909b76fa 100644 --- a/OpenCL/m13752-pure.cl +++ b/OpenCL/m13752-pure.cl @@ -24,6 +24,28 @@ #include "inc_truecrypt_xts.cl" #include "inc_veracrypt_xts.cl" +typedef struct tc +{ + u32 salt_buf[32]; + u32 data_buf[112]; + u32 keyfile_buf[16]; + u32 signature; + + keyboard_layout_mapping_t keyboard_layout_mapping_buf[256]; + int keyboard_layout_mapping_cnt; + +} tc_t; + +typedef struct tc_tmp +{ + u32 ipad[16]; + u32 opad[16]; + + u32 dgst[64]; + u32 out[64]; + +} tc_tmp_t; + DECLSPEC void hmac_sha256_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest) { digest[0] = ipad[0]; diff --git a/OpenCL/m13753-pure.cl b/OpenCL/m13753-pure.cl index 152556f19..77d97d115 100644 --- a/OpenCL/m13753-pure.cl +++ b/OpenCL/m13753-pure.cl @@ -24,6 +24,28 @@ #include "inc_truecrypt_xts.cl" #include "inc_veracrypt_xts.cl" +typedef struct tc +{ + u32 salt_buf[32]; + u32 data_buf[112]; + u32 keyfile_buf[16]; + u32 signature; + + keyboard_layout_mapping_t keyboard_layout_mapping_buf[256]; + int keyboard_layout_mapping_cnt; + +} tc_t; + +typedef struct tc_tmp +{ + u32 ipad[16]; + u32 opad[16]; + + u32 dgst[64]; + u32 out[64]; + +} tc_tmp_t; + DECLSPEC void hmac_sha256_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest) { digest[0] = ipad[0]; diff --git a/OpenCL/m13771-pure.cl b/OpenCL/m13771-pure.cl index ff7a57623..d75822f6c 100644 --- a/OpenCL/m13771-pure.cl +++ b/OpenCL/m13771-pure.cl @@ -24,6 +24,31 @@ #include "inc_truecrypt_xts.cl" #include "inc_veracrypt_xts.cl" +typedef struct tc +{ + u32 salt_buf[32]; + u32 data_buf[112]; + u32 keyfile_buf[16]; + u32 signature; + + keyboard_layout_mapping_t keyboard_layout_mapping_buf[256]; + int keyboard_layout_mapping_cnt; + +} tc_t; + +typedef struct vc64_sbog_tmp +{ + u64 ipad_raw[8]; + u64 opad_raw[8]; + + u64 ipad_hash[8]; + u64 opad_hash[8]; + + u64 dgst[32]; + u64 out[32]; + +} vc64_sbog_tmp_t; + 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]) { const u64x nullbuf[8] = { 0 }; diff --git a/OpenCL/m13772-pure.cl b/OpenCL/m13772-pure.cl index c4107c4b4..7eeaa1ed8 100644 --- a/OpenCL/m13772-pure.cl +++ b/OpenCL/m13772-pure.cl @@ -24,6 +24,31 @@ #include "inc_truecrypt_xts.cl" #include "inc_veracrypt_xts.cl" +typedef struct tc +{ + u32 salt_buf[32]; + u32 data_buf[112]; + u32 keyfile_buf[16]; + u32 signature; + + keyboard_layout_mapping_t keyboard_layout_mapping_buf[256]; + int keyboard_layout_mapping_cnt; + +} tc_t; + +typedef struct vc64_sbog_tmp +{ + u64 ipad_raw[8]; + u64 opad_raw[8]; + + u64 ipad_hash[8]; + u64 opad_hash[8]; + + u64 dgst[32]; + u64 out[32]; + +} vc64_sbog_tmp_t; + 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]) { const u64x nullbuf[8] = { 0 }; diff --git a/OpenCL/m13773-pure.cl b/OpenCL/m13773-pure.cl index a63ad3a0f..2097bddc3 100644 --- a/OpenCL/m13773-pure.cl +++ b/OpenCL/m13773-pure.cl @@ -24,6 +24,31 @@ #include "inc_truecrypt_xts.cl" #include "inc_veracrypt_xts.cl" +typedef struct tc +{ + u32 salt_buf[32]; + u32 data_buf[112]; + u32 keyfile_buf[16]; + u32 signature; + + keyboard_layout_mapping_t keyboard_layout_mapping_buf[256]; + int keyboard_layout_mapping_cnt; + +} tc_t; + +typedef struct vc64_sbog_tmp +{ + u64 ipad_raw[8]; + u64 opad_raw[8]; + + u64 ipad_hash[8]; + u64 opad_hash[8]; + + u64 dgst[32]; + u64 out[32]; + +} vc64_sbog_tmp_t; + 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]) { const u64x nullbuf[8] = { 0 }; diff --git a/OpenCL/m13800_a0-optimized.cl b/OpenCL/m13800_a0-optimized.cl index 1f7bad962..d973fcc78 100644 --- a/OpenCL/m13800_a0-optimized.cl +++ b/OpenCL/m13800_a0-optimized.cl @@ -16,6 +16,12 @@ #include "inc_simd.cl" #include "inc_hash_sha256.cl" +typedef struct win8phone +{ + u32 salt_buf[32]; + +} win8phone_t; + DECLSPEC void sha256_transform_transport_vector (const u32x *w, u32x *digest) { sha256_transform_vector (w + 0, w + 4, w + 8, w + 12, digest); diff --git a/OpenCL/m13800_a0-pure.cl b/OpenCL/m13800_a0-pure.cl index 57fd5b261..3b704f036 100644 --- a/OpenCL/m13800_a0-pure.cl +++ b/OpenCL/m13800_a0-pure.cl @@ -15,6 +15,12 @@ #include "inc_scalar.cl" #include "inc_hash_sha256.cl" +typedef struct win8phone +{ + u32 salt_buf[32]; + +} win8phone_t; + __kernel void m13800_mxx (KERN_ATTR_RULES_ESALT (win8phone_t)) { /** diff --git a/OpenCL/m13800_a1-optimized.cl b/OpenCL/m13800_a1-optimized.cl index b0905ab50..b4118da9f 100644 --- a/OpenCL/m13800_a1-optimized.cl +++ b/OpenCL/m13800_a1-optimized.cl @@ -14,6 +14,12 @@ #include "inc_simd.cl" #include "inc_hash_sha256.cl" +typedef struct win8phone +{ + u32 salt_buf[32]; + +} win8phone_t; + DECLSPEC void sha256_transform_transport_vector (const u32x *w, u32x *digest) { sha256_transform_vector (w + 0, w + 4, w + 8, w + 12, digest); diff --git a/OpenCL/m13800_a1-pure.cl b/OpenCL/m13800_a1-pure.cl index 9a3d73538..24699c4d6 100644 --- a/OpenCL/m13800_a1-pure.cl +++ b/OpenCL/m13800_a1-pure.cl @@ -13,6 +13,12 @@ #include "inc_scalar.cl" #include "inc_hash_sha256.cl" +typedef struct win8phone +{ + u32 salt_buf[32]; + +} win8phone_t; + __kernel void m13800_mxx (KERN_ATTR_ESALT (win8phone_t)) { /** diff --git a/OpenCL/m13800_a3-optimized.cl b/OpenCL/m13800_a3-optimized.cl index 4ea3aa390..73f0c7223 100644 --- a/OpenCL/m13800_a3-optimized.cl +++ b/OpenCL/m13800_a3-optimized.cl @@ -13,6 +13,12 @@ #include "inc_simd.cl" #include "inc_hash_sha256.cl" +typedef struct win8phone +{ + u32 salt_buf[32]; + +} win8phone_t; + DECLSPEC void sha256_transform_transport_vector (const u32x *w, u32x *digest) { sha256_transform_vector (w + 0, w + 4, w + 8, w + 12, digest); diff --git a/OpenCL/m13800_a3-pure.cl b/OpenCL/m13800_a3-pure.cl index 2a0de8d04..a3b2ccff4 100644 --- a/OpenCL/m13800_a3-pure.cl +++ b/OpenCL/m13800_a3-pure.cl @@ -13,6 +13,12 @@ #include "inc_simd.cl" #include "inc_hash_sha256.cl" +typedef struct win8phone +{ + u32 salt_buf[32]; + +} win8phone_t; + __kernel void m13800_mxx (KERN_ATTR_VECTOR_ESALT (win8phone_t)) { /**