From 74abff44d9b660ebdc4b16a5ac00d74846044307 Mon Sep 17 00:00:00 2001 From: Jens Steube Date: Fri, 8 Mar 2019 12:50:31 +0100 Subject: [PATCH] Move remaining module specific typedefs and structures from inc_types.cl to kernels --- OpenCL/m08900-pure.cl | 10 ++++++++++ OpenCL/m14611-pure.cl | 31 +++++++++++++++++++++++++++++++ OpenCL/m14612-pure.cl | 31 +++++++++++++++++++++++++++++++ OpenCL/m14613-pure.cl | 31 +++++++++++++++++++++++++++++++ OpenCL/m14621-pure.cl | 31 +++++++++++++++++++++++++++++++ OpenCL/m14622-pure.cl | 31 +++++++++++++++++++++++++++++++ OpenCL/m14623-pure.cl | 31 +++++++++++++++++++++++++++++++ OpenCL/m14631-pure.cl | 31 +++++++++++++++++++++++++++++++ OpenCL/m14632-pure.cl | 31 +++++++++++++++++++++++++++++++ OpenCL/m14633-pure.cl | 31 +++++++++++++++++++++++++++++++ OpenCL/m14641-pure.cl | 31 +++++++++++++++++++++++++++++++ OpenCL/m14642-pure.cl | 31 +++++++++++++++++++++++++++++++ OpenCL/m14643-pure.cl | 31 +++++++++++++++++++++++++++++++ OpenCL/m14700-pure.cl | 17 +++++++++++++++++ OpenCL/m14800-pure.cl | 17 +++++++++++++++++ OpenCL/m15100-pure.cl | 11 +++++++++++ OpenCL/m15300-pure.cl | 31 +++++++++++++++++++++++++++++++ OpenCL/m15400_a0-optimized.cl | 9 +++++++++ OpenCL/m15400_a1-optimized.cl | 9 +++++++++ OpenCL/m15400_a3-optimized.cl | 9 +++++++++ OpenCL/m15600-pure.cl | 17 +++++++++++++++++ OpenCL/m15700-pure.cl | 17 +++++++++++++++++ OpenCL/m15900-pure.cl | 31 +++++++++++++++++++++++++++++++ OpenCL/m16100_a0-optimized.cl | 11 +++++++++++ OpenCL/m16100_a0-pure.cl | 11 +++++++++++ OpenCL/m16100_a1-optimized.cl | 11 +++++++++++ OpenCL/m16100_a1-pure.cl | 11 +++++++++++ OpenCL/m16100_a3-optimized.cl | 11 +++++++++++ OpenCL/m16100_a3-pure.cl | 11 +++++++++++ OpenCL/m16200-pure.cl | 19 +++++++++++++++++++ OpenCL/m16300-pure.cl | 18 ++++++++++++++++++ OpenCL/m16511_a0-pure.cl | 9 +++++++++ OpenCL/m16511_a1-pure.cl | 9 +++++++++ OpenCL/m16511_a3-pure.cl | 9 +++++++++ OpenCL/m16512_a0-pure.cl | 9 +++++++++ OpenCL/m16512_a1-pure.cl | 9 +++++++++ OpenCL/m16512_a3-pure.cl | 9 +++++++++ OpenCL/m16513_a0-pure.cl | 9 +++++++++ OpenCL/m16513_a1-pure.cl | 9 +++++++++ OpenCL/m16513_a3-pure.cl | 9 +++++++++ OpenCL/m16600_a0-optimized.cl | 8 ++++++++ OpenCL/m16600_a0-pure.cl | 8 ++++++++ OpenCL/m16600_a1-optimized.cl | 8 ++++++++ OpenCL/m16600_a1-pure.cl | 8 ++++++++ OpenCL/m16600_a3-optimized.cl | 8 ++++++++ OpenCL/m16600_a3-pure.cl | 8 ++++++++ OpenCL/m16800-pure.cl | 21 +++++++++++++++++++++ OpenCL/m16801-pure.cl | 21 +++++++++++++++++++++ OpenCL/m16900-pure.cl | 19 +++++++++++++++++++ OpenCL/m18200_a0-optimized.cl | 9 +++++++++ OpenCL/m18200_a0-pure.cl | 9 +++++++++ OpenCL/m18200_a1-optimized.cl | 9 +++++++++ OpenCL/m18200_a1-pure.cl | 9 +++++++++ OpenCL/m18200_a3-optimized.cl | 9 +++++++++ OpenCL/m18200_a3-pure.cl | 9 +++++++++ OpenCL/m18300-pure.cl | 19 +++++++++++++++++++ OpenCL/m18400-pure.cl | 19 +++++++++++++++++++ OpenCL/m18600-pure.cl | 19 +++++++++++++++++++ 58 files changed, 954 insertions(+) diff --git a/OpenCL/m08900-pure.cl b/OpenCL/m08900-pure.cl index 86421f48a..32b3f7fb6 100644 --- a/OpenCL/m08900-pure.cl +++ b/OpenCL/m08900-pure.cl @@ -13,6 +13,16 @@ #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" +typedef struct +{ + #ifndef SCRYPT_TMP_ELEM + #define SCRYPT_TMP_ELEM 1 + #endif + + uint4 P[SCRYPT_TMP_ELEM]; + +} scrypt_tmp_t; + DECLSPEC uint4 swap32_4 (uint4 v) { return (rotate ((v & 0x00FF00FF), 24u) | rotate ((v & 0xFF00FF00), 8u)); diff --git a/OpenCL/m14611-pure.cl b/OpenCL/m14611-pure.cl index 080ca0f53..d4639a6ff 100644 --- a/OpenCL/m14611-pure.cl +++ b/OpenCL/m14611-pure.cl @@ -25,6 +25,37 @@ #define MAX_ENTROPY 7.0 +#define LUKS_STRIPES 4000 + +typedef struct luks +{ + int hash_type; // hc_luks_hash_type_t + int key_size; // hc_luks_key_size_t + int cipher_type; // hc_luks_cipher_type_t + int cipher_mode; // hc_luks_cipher_mode_t + + u32 ct_buf[128]; + + u32 af_src_buf[((HC_LUKS_KEY_SIZE_512 / 8) * LUKS_STRIPES) / 4]; + +} luks_t; + +typedef struct luks_tmp +{ + u32 ipad32[8]; + u64 ipad64[8]; + + u32 opad32[8]; + u64 opad64[8]; + + u32 dgst32[32]; + u64 dgst64[16]; + + u32 out32[32]; + u64 out64[16]; + +} luks_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/m14612-pure.cl b/OpenCL/m14612-pure.cl index d8db96cdb..dad2b0d8b 100644 --- a/OpenCL/m14612-pure.cl +++ b/OpenCL/m14612-pure.cl @@ -25,6 +25,37 @@ #define MAX_ENTROPY 7.0 +#define LUKS_STRIPES 4000 + +typedef struct luks +{ + int hash_type; // hc_luks_hash_type_t + int key_size; // hc_luks_key_size_t + int cipher_type; // hc_luks_cipher_type_t + int cipher_mode; // hc_luks_cipher_mode_t + + u32 ct_buf[128]; + + u32 af_src_buf[((HC_LUKS_KEY_SIZE_512 / 8) * LUKS_STRIPES) / 4]; + +} luks_t; + +typedef struct luks_tmp +{ + u32 ipad32[8]; + u64 ipad64[8]; + + u32 opad32[8]; + u64 opad64[8]; + + u32 dgst32[32]; + u64 dgst64[16]; + + u32 out32[32]; + u64 out64[16]; + +} luks_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/m14613-pure.cl b/OpenCL/m14613-pure.cl index c7315c6fe..34e823c19 100644 --- a/OpenCL/m14613-pure.cl +++ b/OpenCL/m14613-pure.cl @@ -25,6 +25,37 @@ #define MAX_ENTROPY 7.0 +#define LUKS_STRIPES 4000 + +typedef struct luks +{ + int hash_type; // hc_luks_hash_type_t + int key_size; // hc_luks_key_size_t + int cipher_type; // hc_luks_cipher_type_t + int cipher_mode; // hc_luks_cipher_mode_t + + u32 ct_buf[128]; + + u32 af_src_buf[((HC_LUKS_KEY_SIZE_512 / 8) * LUKS_STRIPES) / 4]; + +} luks_t; + +typedef struct luks_tmp +{ + u32 ipad32[8]; + u64 ipad64[8]; + + u32 opad32[8]; + u64 opad64[8]; + + u32 dgst32[32]; + u64 dgst64[16]; + + u32 out32[32]; + u64 out64[16]; + +} luks_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/m14621-pure.cl b/OpenCL/m14621-pure.cl index 779fd2bbb..8e137babf 100644 --- a/OpenCL/m14621-pure.cl +++ b/OpenCL/m14621-pure.cl @@ -25,6 +25,37 @@ #define MAX_ENTROPY 7.0 +#define LUKS_STRIPES 4000 + +typedef struct luks +{ + int hash_type; // hc_luks_hash_type_t + int key_size; // hc_luks_key_size_t + int cipher_type; // hc_luks_cipher_type_t + int cipher_mode; // hc_luks_cipher_mode_t + + u32 ct_buf[128]; + + u32 af_src_buf[((HC_LUKS_KEY_SIZE_512 / 8) * LUKS_STRIPES) / 4]; + +} luks_t; + +typedef struct luks_tmp +{ + u32 ipad32[8]; + u64 ipad64[8]; + + u32 opad32[8]; + u64 opad64[8]; + + u32 dgst32[32]; + u64 dgst64[16]; + + u32 out32[32]; + u64 out64[16]; + +} luks_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/m14622-pure.cl b/OpenCL/m14622-pure.cl index 7594ca53c..1f2eae2a5 100644 --- a/OpenCL/m14622-pure.cl +++ b/OpenCL/m14622-pure.cl @@ -25,6 +25,37 @@ #define MAX_ENTROPY 7.0 +#define LUKS_STRIPES 4000 + +typedef struct luks +{ + int hash_type; // hc_luks_hash_type_t + int key_size; // hc_luks_key_size_t + int cipher_type; // hc_luks_cipher_type_t + int cipher_mode; // hc_luks_cipher_mode_t + + u32 ct_buf[128]; + + u32 af_src_buf[((HC_LUKS_KEY_SIZE_512 / 8) * LUKS_STRIPES) / 4]; + +} luks_t; + +typedef struct luks_tmp +{ + u32 ipad32[8]; + u64 ipad64[8]; + + u32 opad32[8]; + u64 opad64[8]; + + u32 dgst32[32]; + u64 dgst64[16]; + + u32 out32[32]; + u64 out64[16]; + +} luks_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/m14623-pure.cl b/OpenCL/m14623-pure.cl index 1f70d96d3..1983be349 100644 --- a/OpenCL/m14623-pure.cl +++ b/OpenCL/m14623-pure.cl @@ -25,6 +25,37 @@ #define MAX_ENTROPY 7.0 +#define LUKS_STRIPES 4000 + +typedef struct luks +{ + int hash_type; // hc_luks_hash_type_t + int key_size; // hc_luks_key_size_t + int cipher_type; // hc_luks_cipher_type_t + int cipher_mode; // hc_luks_cipher_mode_t + + u32 ct_buf[128]; + + u32 af_src_buf[((HC_LUKS_KEY_SIZE_512 / 8) * LUKS_STRIPES) / 4]; + +} luks_t; + +typedef struct luks_tmp +{ + u32 ipad32[8]; + u64 ipad64[8]; + + u32 opad32[8]; + u64 opad64[8]; + + u32 dgst32[32]; + u64 dgst64[16]; + + u32 out32[32]; + u64 out64[16]; + +} luks_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/m14631-pure.cl b/OpenCL/m14631-pure.cl index c31ee140f..a8a8f6117 100644 --- a/OpenCL/m14631-pure.cl +++ b/OpenCL/m14631-pure.cl @@ -25,6 +25,37 @@ #define MAX_ENTROPY 7.0 +#define LUKS_STRIPES 4000 + +typedef struct luks +{ + int hash_type; // hc_luks_hash_type_t + int key_size; // hc_luks_key_size_t + int cipher_type; // hc_luks_cipher_type_t + int cipher_mode; // hc_luks_cipher_mode_t + + u32 ct_buf[128]; + + u32 af_src_buf[((HC_LUKS_KEY_SIZE_512 / 8) * LUKS_STRIPES) / 4]; + +} luks_t; + +typedef struct luks_tmp +{ + u32 ipad32[8]; + u64 ipad64[8]; + + u32 opad32[8]; + u64 opad64[8]; + + u32 dgst32[32]; + u64 dgst64[16]; + + u32 out32[32]; + u64 out64[16]; + +} luks_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/m14632-pure.cl b/OpenCL/m14632-pure.cl index 24b647db0..cc6b6d61d 100644 --- a/OpenCL/m14632-pure.cl +++ b/OpenCL/m14632-pure.cl @@ -25,6 +25,37 @@ #define MAX_ENTROPY 7.0 +#define LUKS_STRIPES 4000 + +typedef struct luks +{ + int hash_type; // hc_luks_hash_type_t + int key_size; // hc_luks_key_size_t + int cipher_type; // hc_luks_cipher_type_t + int cipher_mode; // hc_luks_cipher_mode_t + + u32 ct_buf[128]; + + u32 af_src_buf[((HC_LUKS_KEY_SIZE_512 / 8) * LUKS_STRIPES) / 4]; + +} luks_t; + +typedef struct luks_tmp +{ + u32 ipad32[8]; + u64 ipad64[8]; + + u32 opad32[8]; + u64 opad64[8]; + + u32 dgst32[32]; + u64 dgst64[16]; + + u32 out32[32]; + u64 out64[16]; + +} luks_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/m14633-pure.cl b/OpenCL/m14633-pure.cl index c12632dfc..2f5733419 100644 --- a/OpenCL/m14633-pure.cl +++ b/OpenCL/m14633-pure.cl @@ -25,6 +25,37 @@ #define MAX_ENTROPY 7.0 +#define LUKS_STRIPES 4000 + +typedef struct luks +{ + int hash_type; // hc_luks_hash_type_t + int key_size; // hc_luks_key_size_t + int cipher_type; // hc_luks_cipher_type_t + int cipher_mode; // hc_luks_cipher_mode_t + + u32 ct_buf[128]; + + u32 af_src_buf[((HC_LUKS_KEY_SIZE_512 / 8) * LUKS_STRIPES) / 4]; + +} luks_t; + +typedef struct luks_tmp +{ + u32 ipad32[8]; + u64 ipad64[8]; + + u32 opad32[8]; + u64 opad64[8]; + + u32 dgst32[32]; + u64 dgst64[16]; + + u32 out32[32]; + u64 out64[16]; + +} luks_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/m14641-pure.cl b/OpenCL/m14641-pure.cl index 4c97d27d6..20a8bcbf1 100644 --- a/OpenCL/m14641-pure.cl +++ b/OpenCL/m14641-pure.cl @@ -25,6 +25,37 @@ #define MAX_ENTROPY 7.0 +#define LUKS_STRIPES 4000 + +typedef struct luks +{ + int hash_type; // hc_luks_hash_type_t + int key_size; // hc_luks_key_size_t + int cipher_type; // hc_luks_cipher_type_t + int cipher_mode; // hc_luks_cipher_mode_t + + u32 ct_buf[128]; + + u32 af_src_buf[((HC_LUKS_KEY_SIZE_512 / 8) * LUKS_STRIPES) / 4]; + +} luks_t; + +typedef struct luks_tmp +{ + u32 ipad32[8]; + u64 ipad64[8]; + + u32 opad32[8]; + u64 opad64[8]; + + u32 dgst32[32]; + u64 dgst64[16]; + + u32 out32[32]; + u64 out64[16]; + +} luks_tmp_t; + DECLSPEC void hmac_ripemd160_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest) { digest[0] = ipad[0]; diff --git a/OpenCL/m14642-pure.cl b/OpenCL/m14642-pure.cl index 2fbbadb8b..9236f7677 100644 --- a/OpenCL/m14642-pure.cl +++ b/OpenCL/m14642-pure.cl @@ -25,6 +25,37 @@ #define MAX_ENTROPY 7.0 +#define LUKS_STRIPES 4000 + +typedef struct luks +{ + int hash_type; // hc_luks_hash_type_t + int key_size; // hc_luks_key_size_t + int cipher_type; // hc_luks_cipher_type_t + int cipher_mode; // hc_luks_cipher_mode_t + + u32 ct_buf[128]; + + u32 af_src_buf[((HC_LUKS_KEY_SIZE_512 / 8) * LUKS_STRIPES) / 4]; + +} luks_t; + +typedef struct luks_tmp +{ + u32 ipad32[8]; + u64 ipad64[8]; + + u32 opad32[8]; + u64 opad64[8]; + + u32 dgst32[32]; + u64 dgst64[16]; + + u32 out32[32]; + u64 out64[16]; + +} luks_tmp_t; + DECLSPEC void hmac_ripemd160_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest) { digest[0] = ipad[0]; diff --git a/OpenCL/m14643-pure.cl b/OpenCL/m14643-pure.cl index 23895de0e..2cd539ba1 100644 --- a/OpenCL/m14643-pure.cl +++ b/OpenCL/m14643-pure.cl @@ -25,6 +25,37 @@ #define MAX_ENTROPY 7.0 +#define LUKS_STRIPES 4000 + +typedef struct luks +{ + int hash_type; // hc_luks_hash_type_t + int key_size; // hc_luks_key_size_t + int cipher_type; // hc_luks_cipher_type_t + int cipher_mode; // hc_luks_cipher_mode_t + + u32 ct_buf[128]; + + u32 af_src_buf[((HC_LUKS_KEY_SIZE_512 / 8) * LUKS_STRIPES) / 4]; + +} luks_t; + +typedef struct luks_tmp +{ + u32 ipad32[8]; + u64 ipad64[8]; + + u32 opad32[8]; + u64 opad64[8]; + + u32 dgst32[32]; + u64 dgst64[16]; + + u32 out32[32]; + u64 out64[16]; + +} luks_tmp_t; + DECLSPEC void hmac_ripemd160_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest) { digest[0] = ipad[0]; diff --git a/OpenCL/m14700-pure.cl b/OpenCL/m14700-pure.cl index e8f84071b..a2f71b569 100644 --- a/OpenCL/m14700-pure.cl +++ b/OpenCL/m14700-pure.cl @@ -17,6 +17,23 @@ #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" +typedef struct itunes_backup +{ + u32 wpky[10]; + u32 dpsl[5]; + +} itunes_backup_t; + +typedef struct pbkdf2_sha1_tmp +{ + u32 ipad[5]; + u32 opad[5]; + + u32 dgst[32]; + u32 out[32]; + +} pbkdf2_sha1_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/m14800-pure.cl b/OpenCL/m14800-pure.cl index d01de2098..d85a11100 100644 --- a/OpenCL/m14800-pure.cl +++ b/OpenCL/m14800-pure.cl @@ -15,6 +15,23 @@ #include "inc_hash_sha256.cl" #include "inc_cipher_aes.cl" +typedef struct itunes_backup +{ + u32 wpky[10]; + u32 dpsl[5]; + +} itunes_backup_t; + +typedef struct pbkdf2_sha256_tmp +{ + u32 ipad[8]; + u32 opad[8]; + + u32 dgst[32]; + u32 out[32]; + +} pbkdf2_sha256_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/m15100-pure.cl b/OpenCL/m15100-pure.cl index e79ffb5e2..f3a439a25 100644 --- a/OpenCL/m15100-pure.cl +++ b/OpenCL/m15100-pure.cl @@ -16,6 +16,17 @@ #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" +typedef struct pbkdf1_sha1_tmp +{ + // pbkdf1-sha1 is limited to 160 bits + + u32 ipad[5]; + u32 opad[5]; + + u32 out[5]; + +} pbkdf1_sha1_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/m15300-pure.cl b/OpenCL/m15300-pure.cl index d3f68d6dc..a43a61718 100644 --- a/OpenCL/m15300-pure.cl +++ b/OpenCL/m15300-pure.cl @@ -17,6 +17,37 @@ #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" +typedef struct dpapimk +{ + u32 context; + + u32 SID[32]; + u32 SID_len; + u32 SID_offset; + + /* here only for possible + forward compatibiliy + */ + // u8 cipher_algo[16]; + // u8 hash_algo[16]; + + u32 iv[4]; + u32 contents_len; + u32 contents[128]; + +} dpapimk_t; + +typedef struct dpapimk_tmp_v1 +{ + u32 ipad[5]; + u32 opad[5]; + u32 dgst[10]; + u32 out[10]; + + u32 userKey[5]; + +} dpapimk_tmp_v1_t; + #define PERM_OP(a,b,tt,n,m) \ { \ tt = a >> n; \ diff --git a/OpenCL/m15400_a0-optimized.cl b/OpenCL/m15400_a0-optimized.cl index 2e5e7536c..cb018105c 100644 --- a/OpenCL/m15400_a0-optimized.cl +++ b/OpenCL/m15400_a0-optimized.cl @@ -14,6 +14,15 @@ #include "inc_rp_optimized.cl" #include "inc_simd.cl" +typedef struct chacha20 +{ + u32 iv[2]; + u32 plain[2]; + u32 position[2]; + u32 offset; + +} chacha20_t; + #define CHACHA_CONST_00 0x61707865 #define CHACHA_CONST_01 0x3320646e #define CHACHA_CONST_02 0x79622d32 diff --git a/OpenCL/m15400_a1-optimized.cl b/OpenCL/m15400_a1-optimized.cl index 13378b279..5ffedf3dd 100644 --- a/OpenCL/m15400_a1-optimized.cl +++ b/OpenCL/m15400_a1-optimized.cl @@ -14,6 +14,15 @@ #include "inc_rp_optimized.cl" #include "inc_simd.cl" +typedef struct chacha20 +{ + u32 iv[2]; + u32 plain[2]; + u32 position[2]; + u32 offset; + +} chacha20_t; + #define CHACHA_CONST_00 0x61707865 #define CHACHA_CONST_01 0x3320646e #define CHACHA_CONST_02 0x79622d32 diff --git a/OpenCL/m15400_a3-optimized.cl b/OpenCL/m15400_a3-optimized.cl index 4db30c3bc..57e24d579 100644 --- a/OpenCL/m15400_a3-optimized.cl +++ b/OpenCL/m15400_a3-optimized.cl @@ -12,6 +12,15 @@ #include "inc_common.cl" #include "inc_simd.cl" +typedef struct chacha20 +{ + u32 iv[2]; + u32 plain[2]; + u32 position[2]; + u32 offset; + +} chacha20_t; + #define CHACHA_CONST_00 0x61707865 #define CHACHA_CONST_01 0x3320646e #define CHACHA_CONST_02 0x79622d32 diff --git a/OpenCL/m15600-pure.cl b/OpenCL/m15600-pure.cl index 711903a13..2cb0523b8 100644 --- a/OpenCL/m15600-pure.cl +++ b/OpenCL/m15600-pure.cl @@ -13,6 +13,23 @@ #include "inc_simd.cl" #include "inc_hash_sha256.cl" +typedef struct pbkdf2_sha256_tmp +{ + u32 ipad[8]; + u32 opad[8]; + + u32 dgst[32]; + u32 out[32]; + +} pbkdf2_sha256_tmp_t; + +typedef struct ethereum_pbkdf2 +{ + u32 salt_buf[16]; + u32 ciphertext[8]; + +} ethereum_pbkdf2_t; + #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" diff --git a/OpenCL/m15700-pure.cl b/OpenCL/m15700-pure.cl index 20dd83009..8829b1dc8 100644 --- a/OpenCL/m15700-pure.cl +++ b/OpenCL/m15700-pure.cl @@ -13,6 +13,23 @@ #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" +typedef struct +{ + #ifndef SCRYPT_TMP_ELEM + #define SCRYPT_TMP_ELEM 1 + #endif + + uint4 P[SCRYPT_TMP_ELEM]; + +} scrypt_tmp_t; + +typedef struct ethereum_scrypt +{ + u32 salt_buf[16]; + u32 ciphertext[8]; + +} ethereum_scrypt_t; + DECLSPEC uint4 swap32_4 (uint4 v) { return (rotate ((v & 0x00FF00FF), 24u) | rotate ((v & 0xFF00FF00), 8u)); diff --git a/OpenCL/m15900-pure.cl b/OpenCL/m15900-pure.cl index 807b30b64..9d9d62807 100644 --- a/OpenCL/m15900-pure.cl +++ b/OpenCL/m15900-pure.cl @@ -19,6 +19,37 @@ #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" +typedef struct dpapimk_tmp_v2 +{ + u64 ipad64[8]; + u64 opad64[8]; + u64 dgst64[16]; + u64 out64[16]; + + u32 userKey[8]; + +} dpapimk_tmp_v2_t; + +typedef struct dpapimk +{ + u32 context; + + u32 SID[32]; + u32 SID_len; + u32 SID_offset; + + /* here only for possible + forward compatibiliy + */ + // u8 cipher_algo[16]; + // u8 hash_algo[16]; + + u32 iv[4]; + u32 contents_len; + u32 contents[128]; + +} dpapimk_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/m16100_a0-optimized.cl b/OpenCL/m16100_a0-optimized.cl index be6caa422..44b90ad63 100644 --- a/OpenCL/m16100_a0-optimized.cl +++ b/OpenCL/m16100_a0-optimized.cl @@ -14,6 +14,17 @@ #include "inc_rp_optimized.cl" #include "inc_simd.cl" +typedef struct tacacs_plus +{ + u32 session_buf[16]; + + u32 ct_data_buf[64]; + u32 ct_data_len; + + u32 sequence_buf[16]; + +} tacacs_plus_t; + __kernel void m16100_m04 (KERN_ATTR_RULES_ESALT (tacacs_plus_t)) { /** diff --git a/OpenCL/m16100_a0-pure.cl b/OpenCL/m16100_a0-pure.cl index 061b26a96..ac2e81c66 100644 --- a/OpenCL/m16100_a0-pure.cl +++ b/OpenCL/m16100_a0-pure.cl @@ -15,6 +15,17 @@ #include "inc_scalar.cl" #include "inc_hash_md5.cl" +typedef struct tacacs_plus +{ + u32 session_buf[16]; + + u32 ct_data_buf[64]; + u32 ct_data_len; + + u32 sequence_buf[16]; + +} tacacs_plus_t; + __kernel void m16100_mxx (KERN_ATTR_RULES_ESALT (tacacs_plus_t)) { /** diff --git a/OpenCL/m16100_a1-optimized.cl b/OpenCL/m16100_a1-optimized.cl index 8ae70c8ff..a24d654d1 100644 --- a/OpenCL/m16100_a1-optimized.cl +++ b/OpenCL/m16100_a1-optimized.cl @@ -12,6 +12,17 @@ #include "inc_common.cl" #include "inc_simd.cl" +typedef struct tacacs_plus +{ + u32 session_buf[16]; + + u32 ct_data_buf[64]; + u32 ct_data_len; + + u32 sequence_buf[16]; + +} tacacs_plus_t; + __kernel void m16100_m04 (KERN_ATTR_ESALT (tacacs_plus_t)) { /** diff --git a/OpenCL/m16100_a1-pure.cl b/OpenCL/m16100_a1-pure.cl index 749db59aa..7eb1ee0bd 100644 --- a/OpenCL/m16100_a1-pure.cl +++ b/OpenCL/m16100_a1-pure.cl @@ -12,6 +12,17 @@ #include "inc_common.cl" #include "inc_hash_md5.cl" +typedef struct tacacs_plus +{ + u32 session_buf[16]; + + u32 ct_data_buf[64]; + u32 ct_data_len; + + u32 sequence_buf[16]; + +} tacacs_plus_t; + __kernel void m16100_mxx (KERN_ATTR_ESALT (tacacs_plus_t)) { /** diff --git a/OpenCL/m16100_a3-optimized.cl b/OpenCL/m16100_a3-optimized.cl index 2d4fae198..023ab906c 100644 --- a/OpenCL/m16100_a3-optimized.cl +++ b/OpenCL/m16100_a3-optimized.cl @@ -12,6 +12,17 @@ #include "inc_common.cl" #include "inc_simd.cl" +typedef struct tacacs_plus +{ + u32 session_buf[16]; + + u32 ct_data_buf[64]; + u32 ct_data_len; + + u32 sequence_buf[16]; + +} tacacs_plus_t; + DECLSPEC void m16100m (u32 *w0, u32 *w1, u32 *w2, u32 *w3, const u32 pw_len, KERN_ATTR_ESALT (tacacs_plus_t)) { /** diff --git a/OpenCL/m16100_a3-pure.cl b/OpenCL/m16100_a3-pure.cl index 3133b5c8b..6201122b2 100644 --- a/OpenCL/m16100_a3-pure.cl +++ b/OpenCL/m16100_a3-pure.cl @@ -12,6 +12,17 @@ #include "inc_common.cl" #include "inc_hash_md5.cl" +typedef struct tacacs_plus +{ + u32 session_buf[16]; + + u32 ct_data_buf[64]; + u32 ct_data_len; + + u32 sequence_buf[16]; + +} tacacs_plus_t; + __kernel void m16100_mxx (KERN_ATTR_VECTOR_ESALT (tacacs_plus_t)) { /** diff --git a/OpenCL/m16200-pure.cl b/OpenCL/m16200-pure.cl index 34d13117d..066e919e7 100644 --- a/OpenCL/m16200-pure.cl +++ b/OpenCL/m16200-pure.cl @@ -17,6 +17,25 @@ #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" +typedef struct apple_secure_notes_tmp +{ + u32 ipad[8]; + u32 opad[8]; + + u32 dgst[8]; + u32 out[8]; + +} apple_secure_notes_tmp_t; + +typedef struct apple_secure_notes +{ + u32 Z_PK; + u32 ZCRYPTOITERATIONCOUNT; + u32 ZCRYPTOSALT[16]; + u32 ZCRYPTOWRAPPEDKEY[16]; + +} apple_secure_notes_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/m16300-pure.cl b/OpenCL/m16300-pure.cl index ae6c0261d..c23dc5c64 100644 --- a/OpenCL/m16300-pure.cl +++ b/OpenCL/m16300-pure.cl @@ -17,6 +17,24 @@ #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" +typedef struct ethereum_presale +{ + u32 iv[4]; + u32 enc_seed[152]; + u32 enc_seed_len; + +} ethereum_presale_t; + +typedef struct pbkdf2_sha256_tmp +{ + u32 ipad[8]; + u32 opad[8]; + + u32 dgst[32]; + u32 out[32]; + +} pbkdf2_sha256_tmp_t; + __constant u64a keccakf_rndc[24] = { 0x0000000000000001, 0x0000000000008082, 0x800000000000808a, diff --git a/OpenCL/m16511_a0-pure.cl b/OpenCL/m16511_a0-pure.cl index 895e4fb77..bad4b7030 100644 --- a/OpenCL/m16511_a0-pure.cl +++ b/OpenCL/m16511_a0-pure.cl @@ -15,6 +15,15 @@ #include "inc_scalar.cl" #include "inc_hash_sha256.cl" +typedef struct jwt +{ + u32 salt_buf[1024]; + u32 salt_len; + + u32 signature_len; + +} jwt_t; + __kernel void m16511_mxx (KERN_ATTR_RULES_ESALT (jwt_t)) { /** diff --git a/OpenCL/m16511_a1-pure.cl b/OpenCL/m16511_a1-pure.cl index 742b6dc39..e6df24d78 100644 --- a/OpenCL/m16511_a1-pure.cl +++ b/OpenCL/m16511_a1-pure.cl @@ -13,6 +13,15 @@ #include "inc_scalar.cl" #include "inc_hash_sha256.cl" +typedef struct jwt +{ + u32 salt_buf[1024]; + u32 salt_len; + + u32 signature_len; + +} jwt_t; + __kernel void m16511_mxx (KERN_ATTR_ESALT (jwt_t)) { /** diff --git a/OpenCL/m16511_a3-pure.cl b/OpenCL/m16511_a3-pure.cl index 8cef28d2c..2812273da 100644 --- a/OpenCL/m16511_a3-pure.cl +++ b/OpenCL/m16511_a3-pure.cl @@ -13,6 +13,15 @@ #include "inc_simd.cl" #include "inc_hash_sha256.cl" +typedef struct jwt +{ + u32 salt_buf[1024]; + u32 salt_len; + + u32 signature_len; + +} jwt_t; + __kernel void m16511_mxx (KERN_ATTR_VECTOR_ESALT (jwt_t)) { /** diff --git a/OpenCL/m16512_a0-pure.cl b/OpenCL/m16512_a0-pure.cl index 356f2f9bf..4a048297b 100644 --- a/OpenCL/m16512_a0-pure.cl +++ b/OpenCL/m16512_a0-pure.cl @@ -15,6 +15,15 @@ #include "inc_scalar.cl" #include "inc_hash_sha384.cl" +typedef struct jwt +{ + u32 salt_buf[1024]; + u32 salt_len; + + u32 signature_len; + +} jwt_t; + __kernel void m16512_mxx (KERN_ATTR_RULES_ESALT (jwt_t)) { /** diff --git a/OpenCL/m16512_a1-pure.cl b/OpenCL/m16512_a1-pure.cl index 2ce9234ae..9a3d9baeb 100644 --- a/OpenCL/m16512_a1-pure.cl +++ b/OpenCL/m16512_a1-pure.cl @@ -13,6 +13,15 @@ #include "inc_scalar.cl" #include "inc_hash_sha384.cl" +typedef struct jwt +{ + u32 salt_buf[1024]; + u32 salt_len; + + u32 signature_len; + +} jwt_t; + __kernel void m16512_mxx (KERN_ATTR_ESALT (jwt_t)) { /** diff --git a/OpenCL/m16512_a3-pure.cl b/OpenCL/m16512_a3-pure.cl index 0df0be7f5..52b453031 100644 --- a/OpenCL/m16512_a3-pure.cl +++ b/OpenCL/m16512_a3-pure.cl @@ -13,6 +13,15 @@ #include "inc_simd.cl" #include "inc_hash_sha384.cl" +typedef struct jwt +{ + u32 salt_buf[1024]; + u32 salt_len; + + u32 signature_len; + +} jwt_t; + __kernel void m16512_mxx (KERN_ATTR_VECTOR_ESALT (jwt_t)) { /** diff --git a/OpenCL/m16513_a0-pure.cl b/OpenCL/m16513_a0-pure.cl index ac756e765..29490ae05 100644 --- a/OpenCL/m16513_a0-pure.cl +++ b/OpenCL/m16513_a0-pure.cl @@ -15,6 +15,15 @@ #include "inc_scalar.cl" #include "inc_hash_sha512.cl" +typedef struct jwt +{ + u32 salt_buf[1024]; + u32 salt_len; + + u32 signature_len; + +} jwt_t; + __kernel void m16513_mxx (KERN_ATTR_RULES_ESALT (jwt_t)) { /** diff --git a/OpenCL/m16513_a1-pure.cl b/OpenCL/m16513_a1-pure.cl index 8d1f7717d..09586726e 100644 --- a/OpenCL/m16513_a1-pure.cl +++ b/OpenCL/m16513_a1-pure.cl @@ -13,6 +13,15 @@ #include "inc_scalar.cl" #include "inc_hash_sha512.cl" +typedef struct jwt +{ + u32 salt_buf[1024]; + u32 salt_len; + + u32 signature_len; + +} jwt_t; + __kernel void m16513_mxx (KERN_ATTR_ESALT (jwt_t)) { /** diff --git a/OpenCL/m16513_a3-pure.cl b/OpenCL/m16513_a3-pure.cl index 5fb4445f1..590d41568 100644 --- a/OpenCL/m16513_a3-pure.cl +++ b/OpenCL/m16513_a3-pure.cl @@ -13,6 +13,15 @@ #include "inc_simd.cl" #include "inc_hash_sha512.cl" +typedef struct jwt +{ + u32 salt_buf[1024]; + u32 salt_len; + + u32 signature_len; + +} jwt_t; + __kernel void m16513_mxx (KERN_ATTR_VECTOR_ESALT (jwt_t)) { /** diff --git a/OpenCL/m16600_a0-optimized.cl b/OpenCL/m16600_a0-optimized.cl index 67c16956b..b054d1b4c 100644 --- a/OpenCL/m16600_a0-optimized.cl +++ b/OpenCL/m16600_a0-optimized.cl @@ -15,6 +15,14 @@ #include "inc_simd.cl" #include "inc_cipher_aes.cl" +typedef struct electrum_wallet +{ + u32 salt_type; + u32 iv[4]; + u32 encrypted[4]; + +} electrum_wallet_t; + __kernel void m16600_m04 (KERN_ATTR_RULES_ESALT (electrum_wallet_t)) { /** diff --git a/OpenCL/m16600_a0-pure.cl b/OpenCL/m16600_a0-pure.cl index 714fbd967..bc2d4614c 100644 --- a/OpenCL/m16600_a0-pure.cl +++ b/OpenCL/m16600_a0-pure.cl @@ -16,6 +16,14 @@ #include "inc_hash_sha256.cl" #include "inc_cipher_aes.cl" +typedef struct electrum_wallet +{ + u32 salt_type; + u32 iv[4]; + u32 encrypted[4]; + +} electrum_wallet_t; + __kernel void m16600_mxx (KERN_ATTR_RULES_ESALT (electrum_wallet_t)) { /** diff --git a/OpenCL/m16600_a1-optimized.cl b/OpenCL/m16600_a1-optimized.cl index 2f8e59e17..e6cafb3fc 100644 --- a/OpenCL/m16600_a1-optimized.cl +++ b/OpenCL/m16600_a1-optimized.cl @@ -13,6 +13,14 @@ #include "inc_simd.cl" #include "inc_cipher_aes.cl" +typedef struct electrum_wallet +{ + u32 salt_type; + u32 iv[4]; + u32 encrypted[4]; + +} electrum_wallet_t; + __kernel void m16600_m04 (KERN_ATTR_ESALT (electrum_wallet_t)) { /** diff --git a/OpenCL/m16600_a1-pure.cl b/OpenCL/m16600_a1-pure.cl index 58eb1e03a..d586de48b 100644 --- a/OpenCL/m16600_a1-pure.cl +++ b/OpenCL/m16600_a1-pure.cl @@ -14,6 +14,14 @@ #include "inc_hash_sha256.cl" #include "inc_cipher_aes.cl" +typedef struct electrum_wallet +{ + u32 salt_type; + u32 iv[4]; + u32 encrypted[4]; + +} electrum_wallet_t; + __kernel void m16600_mxx (KERN_ATTR_ESALT (electrum_wallet_t)) { /** diff --git a/OpenCL/m16600_a3-optimized.cl b/OpenCL/m16600_a3-optimized.cl index 4c65ca104..7f0546ed5 100644 --- a/OpenCL/m16600_a3-optimized.cl +++ b/OpenCL/m16600_a3-optimized.cl @@ -13,6 +13,14 @@ #include "inc_simd.cl" #include "inc_cipher_aes.cl" +typedef struct electrum_wallet +{ + u32 salt_type; + u32 iv[4]; + u32 encrypted[4]; + +} electrum_wallet_t; + DECLSPEC void m16600 (SHM_TYPE u32a *s_te0, SHM_TYPE u32a *s_te1, SHM_TYPE u32a *s_te2, SHM_TYPE u32a *s_te3, SHM_TYPE u32a *s_te4, SHM_TYPE u32a *s_td0, SHM_TYPE u32a *s_td1, SHM_TYPE u32a *s_td2, SHM_TYPE u32a *s_td3, SHM_TYPE u32a *s_td4, u32 *w, const u32 pw_len, KERN_ATTR_ESALT (electrum_wallet_t)) { /** diff --git a/OpenCL/m16600_a3-pure.cl b/OpenCL/m16600_a3-pure.cl index 215109afb..707474486 100644 --- a/OpenCL/m16600_a3-pure.cl +++ b/OpenCL/m16600_a3-pure.cl @@ -14,6 +14,14 @@ #include "inc_hash_sha256.cl" #include "inc_cipher_aes.cl" +typedef struct electrum_wallet +{ + u32 salt_type; + u32 iv[4]; + u32 encrypted[4]; + +} electrum_wallet_t; + __kernel void m16600_mxx (KERN_ATTR_VECTOR_ESALT (electrum_wallet_t)) { /** diff --git a/OpenCL/m16800-pure.cl b/OpenCL/m16800-pure.cl index 5526d4250..2415b78f8 100644 --- a/OpenCL/m16800-pure.cl +++ b/OpenCL/m16800-pure.cl @@ -16,6 +16,27 @@ #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" +typedef struct wpa_pbkdf2_tmp +{ + u32 ipad[5]; + u32 opad[5]; + + u32 dgst[10]; + u32 out[10]; + +} wpa_pbkdf2_tmp_t; + +typedef struct wpa_pmkid +{ + u32 pmkid[4]; + u32 pmkid_data[16]; + u8 orig_mac_ap[6]; + u8 orig_mac_sta[6]; + u8 essid_len; + u32 essid_buf[16]; + +} wpa_pmkid_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/m16801-pure.cl b/OpenCL/m16801-pure.cl index 87cdb7c21..b9c5cd513 100644 --- a/OpenCL/m16801-pure.cl +++ b/OpenCL/m16801-pure.cl @@ -16,6 +16,27 @@ #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" +typedef struct wpa_pbkdf2_tmp +{ + u32 ipad[5]; + u32 opad[5]; + + u32 dgst[10]; + u32 out[10]; + +} wpa_pbkdf2_tmp_t; + +typedef struct wpa_pmkid +{ + u32 pmkid[4]; + u32 pmkid_data[16]; + u8 orig_mac_ap[6]; + u8 orig_mac_sta[6]; + u8 essid_len; + u32 essid_buf[16]; + +} wpa_pmkid_t; + DECLSPEC u8 hex_convert (const u8 c) { return (c & 15) + (c >> 6) * 9; diff --git a/OpenCL/m16900-pure.cl b/OpenCL/m16900-pure.cl index 89322a47a..c017757d4 100644 --- a/OpenCL/m16900-pure.cl +++ b/OpenCL/m16900-pure.cl @@ -16,6 +16,25 @@ #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 ansible_vault +{ + u32 cipher; + u32 version; + u32 ct_data_buf[4096]; + u32 ct_data_len; + +} ansible_vault_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/m18200_a0-optimized.cl b/OpenCL/m18200_a0-optimized.cl index ac01f9c41..4937d5341 100644 --- a/OpenCL/m18200_a0-optimized.cl +++ b/OpenCL/m18200_a0-optimized.cl @@ -17,6 +17,15 @@ #include "inc_hash_md4.cl" #include "inc_hash_md5.cl" +typedef struct krb5asrep +{ + u32 account_info[512]; + u32 checksum[4]; + u32 edata2[5120]; + u32 edata2_len; + +} krb5asrep_t; + typedef struct { u8 S[256]; diff --git a/OpenCL/m18200_a0-pure.cl b/OpenCL/m18200_a0-pure.cl index 63012b5c3..6fb3061da 100644 --- a/OpenCL/m18200_a0-pure.cl +++ b/OpenCL/m18200_a0-pure.cl @@ -16,6 +16,15 @@ #include "inc_hash_md4.cl" #include "inc_hash_md5.cl" +typedef struct krb5asrep +{ + u32 account_info[512]; + u32 checksum[4]; + u32 edata2[5120]; + u32 edata2_len; + +} krb5asrep_t; + typedef struct { u8 S[256]; diff --git a/OpenCL/m18200_a1-optimized.cl b/OpenCL/m18200_a1-optimized.cl index 7de38dae1..546e7b857 100644 --- a/OpenCL/m18200_a1-optimized.cl +++ b/OpenCL/m18200_a1-optimized.cl @@ -15,6 +15,15 @@ #include "inc_hash_md4.cl" #include "inc_hash_md5.cl" +typedef struct krb5asrep +{ + u32 account_info[512]; + u32 checksum[4]; + u32 edata2[5120]; + u32 edata2_len; + +} krb5asrep_t; + typedef struct { u8 S[256]; diff --git a/OpenCL/m18200_a1-pure.cl b/OpenCL/m18200_a1-pure.cl index dc6d6d9f1..409de651a 100644 --- a/OpenCL/m18200_a1-pure.cl +++ b/OpenCL/m18200_a1-pure.cl @@ -14,6 +14,15 @@ #include "inc_hash_md4.cl" #include "inc_hash_md5.cl" +typedef struct krb5asrep +{ + u32 account_info[512]; + u32 checksum[4]; + u32 edata2[5120]; + u32 edata2_len; + +} krb5asrep_t; + typedef struct { u8 S[256]; diff --git a/OpenCL/m18200_a3-optimized.cl b/OpenCL/m18200_a3-optimized.cl index f429e33b5..1c0aa5368 100644 --- a/OpenCL/m18200_a3-optimized.cl +++ b/OpenCL/m18200_a3-optimized.cl @@ -15,6 +15,15 @@ #include "inc_hash_md4.cl" #include "inc_hash_md5.cl" +typedef struct krb5asrep +{ + u32 account_info[512]; + u32 checksum[4]; + u32 edata2[5120]; + u32 edata2_len; + +} krb5asrep_t; + typedef struct { u8 S[256]; diff --git a/OpenCL/m18200_a3-pure.cl b/OpenCL/m18200_a3-pure.cl index cb27d0b3f..40601a151 100644 --- a/OpenCL/m18200_a3-pure.cl +++ b/OpenCL/m18200_a3-pure.cl @@ -14,6 +14,15 @@ #include "inc_hash_md4.cl" #include "inc_hash_md5.cl" +typedef struct krb5asrep +{ + u32 account_info[512]; + u32 checksum[4]; + u32 edata2[5120]; + u32 edata2_len; + +} krb5asrep_t; + typedef struct { u8 S[256]; diff --git a/OpenCL/m18300-pure.cl b/OpenCL/m18300-pure.cl index 00bca9557..5cd3cfd72 100644 --- a/OpenCL/m18300-pure.cl +++ b/OpenCL/m18300-pure.cl @@ -17,6 +17,25 @@ #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" +typedef struct apple_secure_notes_tmp +{ + u32 ipad[8]; + u32 opad[8]; + + u32 dgst[8]; + u32 out[8]; + +} apple_secure_notes_tmp_t; + +typedef struct apple_secure_notes +{ + u32 Z_PK; + u32 ZCRYPTOITERATIONCOUNT; + u32 ZCRYPTOSALT[16]; + u32 ZCRYPTOWRAPPEDKEY[16]; + +} apple_secure_notes_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/m18400-pure.cl b/OpenCL/m18400-pure.cl index 9e2cf0975..5f1c08ee2 100644 --- a/OpenCL/m18400-pure.cl +++ b/OpenCL/m18400-pure.cl @@ -18,6 +18,25 @@ #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" +typedef struct odf12_tmp +{ + u32 ipad[5]; + u32 opad[5]; + + u32 dgst[10]; + u32 out[10]; + +} odf12_tmp_t; + +typedef struct odf12 +{ + u32 iterations; + u32 iv[4]; + u32 checksum[8]; + u32 encrypted_data[256]; + +} odf12_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/m18600-pure.cl b/OpenCL/m18600-pure.cl index 75ee593d1..d7ab79fe7 100644 --- a/OpenCL/m18600-pure.cl +++ b/OpenCL/m18600-pure.cl @@ -16,6 +16,25 @@ #define COMPARE_S "inc_comp_single.cl" #define COMPARE_M "inc_comp_multi.cl" +typedef struct odf11_tmp +{ + u32 ipad[5]; + u32 opad[5]; + + u32 dgst[5]; + u32 out[5]; + +} odf11_tmp_t; + +typedef struct odf11 +{ + u32 iterations; + u32 iv[2]; + u32 checksum[5]; + u32 encrypted_data[256]; + +} odf11_t; + // http://www.schneier.com/code/constants.txt __constant u32a c_sbox0[256] =