Move remaining module specific typedefs and structures from inc_types.cl to kernels

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));

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];

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];

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];

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];

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];

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];

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];

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];

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];

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];

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];

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];

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];

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];

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];

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;              \

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

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

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

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"
 

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));

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];

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))
 {
   /**

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))
 {
   /**

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))
 {
   /**

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))
 {
   /**

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))
 {
   /**

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))
 {
   /**

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];

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,

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))
 {
   /**

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))
 {
   /**

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))
 {
   /**

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))
 {
   /**

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))
 {
   /**

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))
 {
   /**

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))
 {
   /**

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))
 {
   /**

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))
 {
   /**

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))
 {
   /**

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))
 {
   /**

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))
 {
   /**

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))
 {
   /**

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))
 {
   /**

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))
 {
   /**

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];

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;

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];

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];

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];

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];

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];

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];

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];

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];

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];

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] =