From 4bce25dd9dd15b42dd1fa0cc84e4e9cf514fbc80 Mon Sep 17 00:00:00 2001
From: Jens Steube <jens.steube@gmail.com>
Date: Fri, 8 Mar 2019 10:18:20 +0100
Subject: [PATCH] Move remaining module specific typedefs and structures from
 inc_types.cl to kernels

---
 OpenCL/inc_types.cl           | 958 +---------------------------------
 OpenCL/m00400-optimized.cl    |   6 +
 OpenCL/m00400-pure.cl         |   6 +
 OpenCL/m00500-optimized.cl    |   6 +
 OpenCL/m00500-pure.cl         |   6 +
 OpenCL/m00600_a0-optimized.cl |  10 +
 OpenCL/m00600_a1-optimized.cl |  10 +
 OpenCL/m00600_a3-optimized.cl |  10 +
 OpenCL/m01600-optimized.cl    |   6 +
 OpenCL/m01600-pure.cl         |   6 +
 OpenCL/m01800-optimized.cl    |  14 +
 OpenCL/m01800-pure.cl         |  14 +
 OpenCL/m02100-pure.cl         |  10 +
 OpenCL/m02500-pure.cl         |  33 ++
 OpenCL/m02501-pure.cl         |  33 ++
 OpenCL/m03200-pure.cl         |  13 +
 OpenCL/m05200-pure.cl         |   6 +
 OpenCL/m05300_a0-optimized.cl |  10 +
 OpenCL/m05300_a0-pure.cl      |  10 +
 OpenCL/m05300_a1-optimized.cl |  10 +
 OpenCL/m05300_a1-pure.cl      |  10 +
 OpenCL/m05300_a3-optimized.cl |  10 +
 OpenCL/m05300_a3-pure.cl      |  10 +
 OpenCL/m05400_a0-optimized.cl |  10 +
 OpenCL/m05400_a0-pure.cl      |  10 +
 OpenCL/m05400_a1-optimized.cl |  10 +
 OpenCL/m05400_a1-pure.cl      |  10 +
 OpenCL/m05400_a3-optimized.cl |  10 +
 OpenCL/m05400_a3-pure.cl      |  10 +
 OpenCL/m05500_a0-optimized.cl |  12 +
 OpenCL/m05500_a0-pure.cl      |  12 +
 OpenCL/m05500_a1-optimized.cl |  12 +
 OpenCL/m05500_a1-pure.cl      |  12 +
 OpenCL/m05500_a3-optimized.cl |  12 +
 OpenCL/m05500_a3-pure.cl      |  12 +
 OpenCL/m05600_a0-optimized.cl |  12 +
 OpenCL/m05600_a0-pure.cl      |  12 +
 OpenCL/m05600_a1-optimized.cl |  12 +
 OpenCL/m05600_a1-pure.cl      |  12 +
 OpenCL/m05600_a3-optimized.cl |  12 +
 OpenCL/m05600_a3-pure.cl      |  12 +
 OpenCL/m05800-optimized.cl    |   6 +
 OpenCL/m05800-pure.cl         |   6 +
 OpenCL/m06211-pure.cl         |  22 +
 OpenCL/m06212-pure.cl         |  22 +
 OpenCL/m06213-pure.cl         |  22 +
 OpenCL/m06221-pure.cl         |  22 +
 OpenCL/m06222-pure.cl         |  22 +
 OpenCL/m06223-pure.cl         |  22 +
 OpenCL/m06231-pure.cl         |  22 +
 OpenCL/m06232-pure.cl         |  22 +
 OpenCL/m06233-pure.cl         |  22 +
 OpenCL/m06300-optimized.cl    |   6 +
 OpenCL/m06300-pure.cl         |   6 +
 54 files changed, 682 insertions(+), 951 deletions(-)

diff --git a/OpenCL/inc_types.cl b/OpenCL/inc_types.cl
index dda8b416b..0d883871b 100644
--- a/OpenCL/inc_types.cl
+++ b/OpenCL/inc_types.cl
@@ -1237,6 +1237,13 @@ DECLSPEC u32 hc_lop_0x96_S (const u32 a, const u32 b, const u32 c)
 
 #endif
 
+typedef enum combinator_mode
+{
+  COMBINATOR_MODE_BASE_LEFT  = 10001,
+  COMBINATOR_MODE_BASE_RIGHT = 10002
+
+} combinator_mode_t;
+
 typedef struct digest
 {
   u32 digest_buf[DGST_ELEM];
@@ -1265,940 +1272,6 @@ typedef struct salt
 
 } salt_t;
 
-#define LUKS_STRIPES 4000
-
-typedef enum hc_luks_hash_type
-{
-  HC_LUKS_HASH_TYPE_SHA1      = 1,
-  HC_LUKS_HASH_TYPE_SHA256    = 2,
-  HC_LUKS_HASH_TYPE_SHA512    = 3,
-  HC_LUKS_HASH_TYPE_RIPEMD160 = 4,
-  HC_LUKS_HASH_TYPE_WHIRLPOOL = 5,
-
-} hc_luks_hash_type_t;
-
-typedef enum hc_luks_key_size
-{
-  HC_LUKS_KEY_SIZE_128 = 128,
-  HC_LUKS_KEY_SIZE_256 = 256,
-  HC_LUKS_KEY_SIZE_512 = 512,
-
-} hc_luks_key_size_t;
-
-typedef enum hc_luks_cipher_type
-{
-  HC_LUKS_CIPHER_TYPE_AES     = 1,
-  HC_LUKS_CIPHER_TYPE_SERPENT = 2,
-  HC_LUKS_CIPHER_TYPE_TWOFISH = 3,
-
-} hc_luks_cipher_type_t;
-
-typedef enum hc_luks_cipher_mode
-{
-  HC_LUKS_CIPHER_MODE_CBC_ESSIV = 1,
-  HC_LUKS_CIPHER_MODE_CBC_PLAIN = 2,
-  HC_LUKS_CIPHER_MODE_XTS_PLAIN = 3,
-
-} hc_luks_cipher_mode_t;
-
-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 itunes_backup
-{
-  u32 wpky[10];
-  u32 dpsl[5];
-
-} itunes_backup_t;
-
-typedef struct blake2
-{
-  u64 h[8];
-  u64 t[2];
-  u64 f[2];
-  u32 buflen;
-  u32 outlen;
-
-} blake2_t;
-
-typedef struct chacha20
-{
-  u32 iv[2];
-  u32 plain[2];
-  u32 position[2];
-  u32 offset;
-
-} chacha20_t;
-
-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 wpa_eapol
-{
-  u32  pke[32];
-  u32  eapol[64 + 16];
-  u16  eapol_len;
-  u8   message_pair;
-  int  message_pair_chgd;
-  u8   keyver;
-  u8   orig_mac_ap[6];
-  u8   orig_mac_sta[6];
-  u8   orig_nonce_ap[32];
-  u8   orig_nonce_sta[32];
-  u8   essid_len;
-  u8   essid[32];
-  u32  keymic[4];
-  u32  hash[4];
-  int  nonce_compare;
-  int  nonce_error_corrections;
-  int  detected_le;
-  int  detected_be;
-
-} wpa_eapol_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;
-
-typedef struct bitcoin_wallet
-{
-  u32 cry_master_buf[64];
-  u32 cry_master_len;
-
-} bitcoin_wallet_t;
-
-typedef struct sip
-{
-  u32 salt_buf[32];
-  u32 salt_len;
-
-  u32 esalt_buf[256];
-  u32 esalt_len;
-
-} sip_t;
-
-typedef struct androidfde
-{
-  u32 data[384];
-
-} androidfde_t;
-
-typedef struct ikepsk
-{
-  u32 nr_buf[16];
-  u32 nr_len;
-
-  u32 msg_buf[128];
-  u32 msg_len[6];
-
-} ikepsk_t;
-
-typedef struct netntlm
-{
-  u32 user_len;
-  u32 domain_len;
-  u32 srvchall_len;
-  u32 clichall_len;
-
-  u32 userdomain_buf[64];
-  u32 chall_buf[256];
-
-} netntlm_t;
-
-typedef struct krb5pa
-{
-  u32 user[16];
-  u32 realm[16];
-  u32 salt[32];
-  u32 timestamp[16];
-  u32 checksum[4];
-
-} krb5pa_t;
-
-typedef struct krb5tgs
-{
-  u32 account_info[512];
-  u32 checksum[4];
-  u32 edata2[5120];
-  u32 edata2_len;
-
-} krb5tgs_t;
-
-typedef struct krb5asrep
-{
-  u32 account_info[512];
-  u32 checksum[4];
-  u32 edata2[5120];
-  u32 edata2_len;
-
-} krb5asrep_t;
-
-typedef struct keyboard_layout_mapping
-{
-  u32 src_char;
-  int src_len;
-  u32 dst_char;
-  int dst_len;
-
-} keyboard_layout_mapping_t;
-
-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 pbkdf2_md5
-{
-  u32 salt_buf[64];
-
-} pbkdf2_md5_t;
-
-typedef struct pbkdf2_sha1
-{
-  u32 salt_buf[64];
-
-} pbkdf2_sha1_t;
-
-typedef struct pbkdf2_sha256
-{
-  u32 salt_buf[64];
-
-} pbkdf2_sha256_t;
-
-typedef struct pbkdf2_sha512
-{
-  u32 salt_buf[64];
-
-} pbkdf2_sha512_t;
-
-typedef struct rakp
-{
-  u32 salt_buf[128];
-  u32 salt_len;
-
-} rakp_t;
-
-typedef struct cloudkey
-{
-  u32 data_len;
-  u32 data_buf[512];
-
-} cloudkey_t;
-
-typedef struct office2007
-{
-  u32 encryptedVerifier[4];
-  u32 encryptedVerifierHash[5];
-
-  u32 keySize;
-
-} office2007_t;
-
-typedef struct office2010
-{
-  u32 encryptedVerifier[4];
-  u32 encryptedVerifierHash[8];
-
-} office2010_t;
-
-typedef struct office2013
-{
-  u32 encryptedVerifier[4];
-  u32 encryptedVerifierHash[8];
-
-} office2013_t;
-
-typedef struct oldoffice01
-{
-  u32 version;
-  u32 encryptedVerifier[4];
-  u32 encryptedVerifierHash[4];
-  u32 rc4key[2];
-
-} oldoffice01_t;
-
-typedef struct oldoffice34
-{
-  u32 version;
-  u32 encryptedVerifier[4];
-  u32 encryptedVerifierHash[5];
-  u32 rc4key[2];
-
-} oldoffice34_t;
-
-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;
-
-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;
-
-typedef struct pstoken
-{
-  u32 salt_buf[128];
-  u32 salt_len;
-
-  u32 pc_digest[5];
-  u32 pc_offset;
-
-} pstoken_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;
-
-typedef struct win8phone
-{
-  u32 salt_buf[32];
-
-} win8phone_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;
-
-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 jks_sha1
-{
-  u32 checksum[5];
-  u32 iv[5];
-  u32 enc_key_buf[4096];
-  u32 enc_key_len;
-  u32 der[5];
-  u32 alias[16];
-
-} jks_sha1_t;
-
-typedef struct ethereum_pbkdf2
-{
-  u32 salt_buf[16];
-  u32 ciphertext[8];
-
-} ethereum_pbkdf2_t;
-
-typedef struct ethereum_scrypt
-{
-  u32 salt_buf[16];
-  u32 ciphertext[8];
-
-} ethereum_scrypt_t;
-
-typedef struct ethereum_presale
-{
-  u32 iv[4];
-  u32 enc_seed[152];
-  u32 enc_seed_len;
-
-} ethereum_presale_t;
-
-typedef struct tacacs_plus
-{
-  u32 session_buf[16];
-
-  u32 ct_data_buf[64];
-  u32 ct_data_len;
-
-  u32 sequence_buf[16];
-
-} tacacs_plus_t;
-
-typedef struct apple_secure_notes
-{
-  u32 Z_PK;
-  u32 ZCRYPTOITERATIONCOUNT;
-  u32 ZCRYPTOSALT[16];
-  u32 ZCRYPTOWRAPPEDKEY[16];
-
-} apple_secure_notes_t;
-
-typedef struct jwt
-{
-  u32 salt_buf[1024];
-  u32 salt_len;
-
-  u32 signature_len;
-
-} jwt_t;
-
-typedef struct electrum_wallet
-{
-  u32 salt_type;
-  u32 iv[4];
-  u32 encrypted[4];
-
-} electrum_wallet_t;
-
-typedef struct ansible_vault
-{
-  u32 cipher;
-  u32 version;
-  u32 ct_data_buf[4096];
-  u32 ct_data_len;
-} ansible_vault_t;
-
-typedef struct pdf14_tmp
-{
-  u32 digest[4];
-  u32 out[4];
-
-} pdf14_tmp_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;
-
-typedef struct pdf17l8_tmp
-{
-  union
-  {
-    u32 dgst32[16];
-    u64 dgst64[8];
-  };
-
-  u32 dgst_len;
-  u32 W_len;
-
-} pdf17l8_tmp_t;
-
-typedef struct phpass_tmp
-{
-  u32 digest_buf[4];
-
-} phpass_tmp_t;
-
-typedef struct md5crypt_tmp
-{
-  u32 digest_buf[4];
-
-} md5crypt_tmp_t;
-
-typedef struct sha256crypt_tmp
-{
-  // pure version
-
-  u32 alt_result[8];
-  u32 p_bytes[64];
-  u32 s_bytes[64];
-
-} sha256crypt_tmp_t;
-
-typedef struct sha512crypt_tmp
-{
-  u64 l_alt_result[8];
-  u64 l_p_bytes[2];
-  u64 l_s_bytes[2];
-
-  // pure version
-
-  u32 alt_result[16];
-  u32 p_bytes[64];
-  u32 s_bytes[64];
-
-} sha512crypt_tmp_t;
-
-typedef struct wpa_pbkdf2_tmp
-{
-  u32 ipad[5];
-  u32 opad[5];
-
-  u32 dgst[10];
-  u32 out[10];
-
-} wpa_pbkdf2_tmp_t;
-
-typedef struct wpa_pmk_tmp
-{
-  u32 out[8];
-
-} wpa_pmk_tmp_t;
-
-typedef struct bitcoin_wallet_tmp
-{
-  u64  dgst[8];
-
-} bitcoin_wallet_tmp_t;
-
-typedef struct dcc2_tmp
-{
-  u32 ipad[5];
-  u32 opad[5];
-
-  u32 dgst[5];
-  u32 out[4];
-
-} dcc2_tmp_t;
-
-typedef struct bcrypt_tmp
-{
-  u32 E[18];
-
-  u32 P[18];
-
-  u32 S0[256];
-  u32 S1[256];
-  u32 S2[256];
-  u32 S3[256];
-
-} bcrypt_tmp_t;
-
-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;
-
-typedef struct pwsafe3_tmp
-{
-  u32 digest_buf[8];
-
-} pwsafe3_tmp_t;
-
-typedef struct androidpin_tmp
-{
-  u32 digest_buf[5];
-
-} androidpin_tmp_t;
-
-typedef struct androidfde_tmp
-{
-  u32 ipad[5];
-  u32 opad[5];
-
-  u32 dgst[10];
-  u32 out[10];
-
-} androidfde_tmp_t;
-
-typedef struct tc_tmp
-{
-  u32 ipad[16];
-  u32 opad[16];
-
-  u32 dgst[64];
-  u32 out[64];
-
-} tc_tmp_t;
-
-typedef struct tc64_tmp
-{
-  u64  ipad[8];
-  u64  opad[8];
-
-  u64  dgst[32];
-  u64  out[32];
-
-} tc64_tmp_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;
-
-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;
-
-typedef struct pbkdf2_md5_tmp
-{
-  u32  ipad[4];
-  u32  opad[4];
-
-  u32  dgst[32];
-  u32  out[32];
-
-} pbkdf2_md5_tmp_t;
-
-typedef struct pbkdf2_sha1_tmp
-{
-  u32  ipad[5];
-  u32  opad[5];
-
-  u32  dgst[32];
-  u32  out[32];
-
-} pbkdf2_sha1_tmp_t;
-
-typedef struct pbkdf2_sha256_tmp
-{
-  u32  ipad[8];
-  u32  opad[8];
-
-  u32  dgst[32];
-  u32  out[32];
-
-} pbkdf2_sha256_tmp_t;
-
-typedef struct pbkdf2_sha512_tmp
-{
-  u64  ipad[8];
-  u64  opad[8];
-
-  u64  dgst[16];
-  u64  out[16];
-
-} pbkdf2_sha512_tmp_t;
-
-typedef struct ecryptfs_tmp
-{
-  u64  out[8];
-
-} ecryptfs_tmp_t;
-
-typedef struct oraclet_tmp
-{
-  u64  ipad[8];
-  u64  opad[8];
-
-  u64  dgst[16];
-  u64  out[16];
-
-} oraclet_tmp_t;
-
-typedef struct agilekey_tmp
-{
-  u32 ipad[5];
-  u32 opad[5];
-
-  u32 dgst[5];
-  u32 out[5];
-
-} agilekey_tmp_t;
-
-typedef struct mywallet_tmp
-{
-  u32 ipad[5];
-  u32 opad[5];
-
-  u32 dgst[10];
-  u32 out[10];
-
-} mywallet_tmp_t;
-
-typedef struct sha1aix_tmp
-{
-  u32 ipad[5];
-  u32 opad[5];
-
-  u32 dgst[5];
-  u32 out[5];
-
-} sha1aix_tmp_t;
-
-typedef struct sha256aix_tmp
-{
-  u32 ipad[8];
-  u32 opad[8];
-
-  u32 dgst[8];
-  u32 out[8];
-
-} sha256aix_tmp_t;
-
-typedef struct sha512aix_tmp
-{
-  u64  ipad[8];
-  u64  opad[8];
-
-  u64  dgst[8];
-  u64  out[8];
-
-} sha512aix_tmp_t;
-
-typedef struct lastpass_tmp
-{
-  u32 ipad[8];
-  u32 opad[8];
-
-  u32 dgst[8];
-  u32 out[8];
-
-} lastpass_tmp_t;
-
-typedef struct drupal7_tmp
-{
-  u64  digest_buf[8];
-
-} drupal7_tmp_t;
-
-typedef struct lotus8_tmp
-{
-  u32 ipad[5];
-  u32 opad[5];
-
-  u32 dgst[5];
-  u32 out[5];
-
-} lotus8_tmp_t;
-
-typedef struct office2007_tmp
-{
-  u32 out[5];
-
-} office2007_tmp_t;
-
-typedef struct office2010_tmp
-{
-  u32 out[5];
-
-} office2010_tmp_t;
-
-typedef struct office2013_tmp
-{
-  u64  out[8];
-
-} office2013_tmp_t;
-
-typedef struct saph_sha1_tmp
-{
-  u32 digest_buf[5];
-
-} saph_sha1_tmp_t;
-
-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 axcrypt_tmp
-{
-  u32 KEK[4];
-  u32 lsb[4];
-  u32 cipher[4];
-
-} axcrypt_tmp_t;
-
-typedef struct keepass_tmp
-{
-  u32 tmp_digest[8];
-
-} keepass_tmp_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;
-
-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 apple_secure_notes_tmp
-{
-  u32 ipad[8];
-  u32 opad[8];
-
-  u32 dgst[8];
-  u32 out[8];
-
-} apple_secure_notes_tmp_t;
-
-typedef struct bsdicrypt_tmp
-{
-  u32 Kc[16];
-  u32 Kd[16];
-
-  u32 iv[2];
-
-} bsdicrypt_tmp_t;
-
-typedef struct rar3_tmp
-{
-  u32 dgst[17][5];
-
-} rar3_tmp_t;
-
-typedef struct
-{
-  u32 ukey[8];
-
-  u32 hook_success;
-
-} seven_zip_hook_t;
-
-typedef struct cram_md5
-{
-  u32 user[16];
-
-} cram_md5_t;
-
 typedef struct
 {
   u32 key;
@@ -2256,20 +1329,3 @@ typedef struct plain
   u32  hash_pos;
 
 } plain_t;
-
-typedef struct
-{
-  #ifndef SCRYPT_TMP_ELEM
-  #define SCRYPT_TMP_ELEM 1
-  #endif
-
-  uint4 P[SCRYPT_TMP_ELEM];
-
-} scrypt_tmp_t;
-
-typedef enum combinator_mode
-{
-  COMBINATOR_MODE_BASE_LEFT  = 10001,
-  COMBINATOR_MODE_BASE_RIGHT = 10002
-
-} combinator_mode_t;
diff --git a/OpenCL/m00400-optimized.cl b/OpenCL/m00400-optimized.cl
index ffb51984b..62970e2b1 100644
--- a/OpenCL/m00400-optimized.cl
+++ b/OpenCL/m00400-optimized.cl
@@ -16,6 +16,12 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
+typedef struct phpass_tmp
+{
+  u32 digest_buf[4];
+
+} phpass_tmp_t;
+
 __kernel void m00400_init (KERN_ATTR_TMPS (phpass_tmp_t))
 {
   /**
diff --git a/OpenCL/m00400-pure.cl b/OpenCL/m00400-pure.cl
index 65430e0bd..c945ae3c3 100644
--- a/OpenCL/m00400-pure.cl
+++ b/OpenCL/m00400-pure.cl
@@ -16,6 +16,12 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
+typedef struct phpass_tmp
+{
+  u32 digest_buf[4];
+
+} phpass_tmp_t;
+
 __kernel void m00400_init (KERN_ATTR_TMPS (phpass_tmp_t))
 {
   /**
diff --git a/OpenCL/m00500-optimized.cl b/OpenCL/m00500-optimized.cl
index e2a39b401..eddc33a0a 100644
--- a/OpenCL/m00500-optimized.cl
+++ b/OpenCL/m00500-optimized.cl
@@ -15,6 +15,12 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
+typedef struct md5crypt_tmp
+{
+  u32 digest_buf[4];
+
+} md5crypt_tmp_t;
+
 #define md5crypt_magic 0x00243124u
 
 DECLSPEC void memcat16 (u32 *block0, u32 *block1, u32 *block2, u32 *block3, const u32 offset, const u32 *append)
diff --git a/OpenCL/m00500-pure.cl b/OpenCL/m00500-pure.cl
index bc4816999..f4ea85e76 100644
--- a/OpenCL/m00500-pure.cl
+++ b/OpenCL/m00500-pure.cl
@@ -15,6 +15,12 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
+typedef struct md5crypt_tmp
+{
+  u32 digest_buf[4];
+
+} md5crypt_tmp_t;
+
 #define md5crypt_magic 0x00243124u
 
 __kernel void m00500_init (KERN_ATTR_TMPS (md5crypt_tmp_t))
diff --git a/OpenCL/m00600_a0-optimized.cl b/OpenCL/m00600_a0-optimized.cl
index f5991d1b2..8a64b977b 100644
--- a/OpenCL/m00600_a0-optimized.cl
+++ b/OpenCL/m00600_a0-optimized.cl
@@ -14,6 +14,16 @@
 #include "inc_rp_optimized.cl"
 #include "inc_simd.cl"
 
+typedef struct blake2
+{
+  u64 h[8];
+  u64 t[2];
+  u64 f[2];
+  u32 buflen;
+  u32 outlen;
+
+} blake2_t;
+
 #define BLAKE2B_FINAL   1
 #define BLAKE2B_UPDATE  0
 
diff --git a/OpenCL/m00600_a1-optimized.cl b/OpenCL/m00600_a1-optimized.cl
index be4658f9b..7d4d837bb 100644
--- a/OpenCL/m00600_a1-optimized.cl
+++ b/OpenCL/m00600_a1-optimized.cl
@@ -14,6 +14,16 @@
 #include "inc_rp_optimized.cl"
 #include "inc_simd.cl"
 
+typedef struct blake2
+{
+  u64 h[8];
+  u64 t[2];
+  u64 f[2];
+  u32 buflen;
+  u32 outlen;
+
+} blake2_t;
+
 #define BLAKE2B_FINAL   1
 #define BLAKE2B_UPDATE  0
 
diff --git a/OpenCL/m00600_a3-optimized.cl b/OpenCL/m00600_a3-optimized.cl
index d95f8a5c4..bb3796ade 100644
--- a/OpenCL/m00600_a3-optimized.cl
+++ b/OpenCL/m00600_a3-optimized.cl
@@ -12,6 +12,16 @@
 #include "inc_common.cl"
 #include "inc_simd.cl"
 
+typedef struct blake2
+{
+  u64 h[8];
+  u64 t[2];
+  u64 f[2];
+  u32 buflen;
+  u32 outlen;
+
+} blake2_t;
+
 #define BLAKE2B_FINAL   1
 #define BLAKE2B_UPDATE  0
 
diff --git a/OpenCL/m01600-optimized.cl b/OpenCL/m01600-optimized.cl
index e5a09c5b3..0281c8b87 100644
--- a/OpenCL/m01600-optimized.cl
+++ b/OpenCL/m01600-optimized.cl
@@ -13,6 +13,12 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
+typedef struct md5crypt_tmp
+{
+  u32 digest_buf[4];
+
+} md5crypt_tmp_t;
+
 #define md5apr1_magic0 0x72706124u
 #define md5apr1_magic1 0x00002431u
 
diff --git a/OpenCL/m01600-pure.cl b/OpenCL/m01600-pure.cl
index 0337b1ad1..741a35022 100644
--- a/OpenCL/m01600-pure.cl
+++ b/OpenCL/m01600-pure.cl
@@ -18,6 +18,12 @@
 #define PUTCHAR_LE(a,p,c) ((u8 *)(a))[(p)] = (u8) (c)
 #define GETCHAR_LE(a,p)   ((u8 *)(a))[(p)]
 
+typedef struct md5crypt_tmp
+{
+  u32 digest_buf[4];
+
+} md5crypt_tmp_t;
+
 #define md5apr1_magic0 0x72706124u
 #define md5apr1_magic1 0x00002431u
 
diff --git a/OpenCL/m01800-optimized.cl b/OpenCL/m01800-optimized.cl
index 15a9d6fd9..087c64ef0 100644
--- a/OpenCL/m01800-optimized.cl
+++ b/OpenCL/m01800-optimized.cl
@@ -16,6 +16,20 @@
 #define PUTCHAR64_BE(a,p,c) ((u8 *)(a))[(p) ^ 7] = (u8) (c)
 #define GETCHAR64_BE(a,p)   ((u8 *)(a))[(p) ^ 7]
 
+typedef struct sha512crypt_tmp
+{
+  u64 l_alt_result[8];
+  u64 l_p_bytes[2];
+  u64 l_s_bytes[2];
+
+  // pure version
+
+  u32 alt_result[16];
+  u32 p_bytes[64];
+  u32 s_bytes[64];
+
+} sha512crypt_tmp_t;
+
 typedef struct
 {
   u64 state[8];
diff --git a/OpenCL/m01800-pure.cl b/OpenCL/m01800-pure.cl
index c31be393f..0884f84a3 100644
--- a/OpenCL/m01800-pure.cl
+++ b/OpenCL/m01800-pure.cl
@@ -13,6 +13,20 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
+typedef struct sha512crypt_tmp
+{
+  u64 l_alt_result[8];
+  u64 l_p_bytes[2];
+  u64 l_s_bytes[2];
+
+  // pure version
+
+  u32 alt_result[16];
+  u32 p_bytes[64];
+  u32 s_bytes[64];
+
+} sha512crypt_tmp_t;
+
 __kernel void m01800_init (KERN_ATTR_TMPS (sha512crypt_tmp_t))
 {
   /**
diff --git a/OpenCL/m02100-pure.cl b/OpenCL/m02100-pure.cl
index ecaa9b75b..6060a4a7e 100644
--- a/OpenCL/m02100-pure.cl
+++ b/OpenCL/m02100-pure.cl
@@ -17,6 +17,16 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
+typedef struct dcc2_tmp
+{
+  u32 ipad[5];
+  u32 opad[5];
+
+  u32 dgst[5];
+  u32 out[4];
+
+} dcc2_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/m02500-pure.cl b/OpenCL/m02500-pure.cl
index 907e1566f..c90cd9472 100644
--- a/OpenCL/m02500-pure.cl
+++ b/OpenCL/m02500-pure.cl
@@ -19,6 +19,39 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
+typedef struct wpa_eapol
+{
+  u32  pke[32];
+  u32  eapol[64 + 16];
+  u16  eapol_len;
+  u8   message_pair;
+  int  message_pair_chgd;
+  u8   keyver;
+  u8   orig_mac_ap[6];
+  u8   orig_mac_sta[6];
+  u8   orig_nonce_ap[32];
+  u8   orig_nonce_sta[32];
+  u8   essid_len;
+  u8   essid[32];
+  u32  keymic[4];
+  u32  hash[4];
+  int  nonce_compare;
+  int  nonce_error_corrections;
+  int  detected_le;
+  int  detected_be;
+
+} wpa_eapol_t;
+
+typedef struct wpa_pbkdf2_tmp
+{
+  u32 ipad[5];
+  u32 opad[5];
+
+  u32 dgst[10];
+  u32 out[10];
+
+} wpa_pbkdf2_tmp_t;
+
 DECLSPEC void make_kn (u32 *k)
 {
   u32 kl[4];
diff --git a/OpenCL/m02501-pure.cl b/OpenCL/m02501-pure.cl
index 13d66301c..b86f0bb41 100644
--- a/OpenCL/m02501-pure.cl
+++ b/OpenCL/m02501-pure.cl
@@ -19,6 +19,39 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
+typedef struct wpa_eapol
+{
+  u32  pke[32];
+  u32  eapol[64 + 16];
+  u16  eapol_len;
+  u8   message_pair;
+  int  message_pair_chgd;
+  u8   keyver;
+  u8   orig_mac_ap[6];
+  u8   orig_mac_sta[6];
+  u8   orig_nonce_ap[32];
+  u8   orig_nonce_sta[32];
+  u8   essid_len;
+  u8   essid[32];
+  u32  keymic[4];
+  u32  hash[4];
+  int  nonce_compare;
+  int  nonce_error_corrections;
+  int  detected_le;
+  int  detected_be;
+
+} wpa_eapol_t;
+
+typedef struct wpa_pbkdf2_tmp
+{
+  u32 ipad[5];
+  u32 opad[5];
+
+  u32 dgst[10];
+  u32 out[10];
+
+} wpa_pbkdf2_tmp_t;
+
 DECLSPEC u8 hex_convert (const u8 c)
 {
   return (c & 15) + (c >> 6) * 9;
diff --git a/OpenCL/m03200-pure.cl b/OpenCL/m03200-pure.cl
index 92b2dabfb..6918f2c36 100644
--- a/OpenCL/m03200-pure.cl
+++ b/OpenCL/m03200-pure.cl
@@ -12,6 +12,19 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
+typedef struct bcrypt_tmp
+{
+  u32 E[18];
+
+  u32 P[18];
+
+  u32 S0[256];
+  u32 S1[256];
+  u32 S2[256];
+  u32 S3[256];
+
+} bcrypt_tmp_t;
+
 // http://www.schneier.com/code/constants.txt
 
 __constant u32a c_sbox0[256] =
diff --git a/OpenCL/m05200-pure.cl b/OpenCL/m05200-pure.cl
index c659110cf..23a2d0d0a 100644
--- a/OpenCL/m05200-pure.cl
+++ b/OpenCL/m05200-pure.cl
@@ -16,6 +16,12 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
+typedef struct pwsafe3_tmp
+{
+  u32 digest_buf[8];
+
+} pwsafe3_tmp_t;
+
 __kernel void m05200_init (KERN_ATTR_TMPS (pwsafe3_tmp_t))
 {
   /**
diff --git a/OpenCL/m05300_a0-optimized.cl b/OpenCL/m05300_a0-optimized.cl
index 3e8aed4b4..17ebfe988 100644
--- a/OpenCL/m05300_a0-optimized.cl
+++ b/OpenCL/m05300_a0-optimized.cl
@@ -15,6 +15,16 @@
 #include "inc_simd.cl"
 #include "inc_hash_md5.cl"
 
+typedef struct ikepsk
+{
+  u32 nr_buf[16];
+  u32 nr_len;
+
+  u32 msg_buf[128];
+  u32 msg_len[6];
+
+} ikepsk_t;
+
 DECLSPEC void hmac_md5_pad (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad)
 {
   w0[0] = w0[0] ^ 0x36363636;
diff --git a/OpenCL/m05300_a0-pure.cl b/OpenCL/m05300_a0-pure.cl
index 51e56415a..86e2ffc8b 100644
--- a/OpenCL/m05300_a0-pure.cl
+++ b/OpenCL/m05300_a0-pure.cl
@@ -15,6 +15,16 @@
 #include "inc_scalar.cl"
 #include "inc_hash_md5.cl"
 
+typedef struct ikepsk
+{
+  u32 nr_buf[16];
+  u32 nr_len;
+
+  u32 msg_buf[128];
+  u32 msg_len[6];
+
+} ikepsk_t;
+
 __kernel void m05300_mxx (KERN_ATTR_RULES_ESALT (ikepsk_t))
 {
   /**
diff --git a/OpenCL/m05300_a1-optimized.cl b/OpenCL/m05300_a1-optimized.cl
index db4bf40d7..ac717373f 100644
--- a/OpenCL/m05300_a1-optimized.cl
+++ b/OpenCL/m05300_a1-optimized.cl
@@ -13,6 +13,16 @@
 #include "inc_simd.cl"
 #include "inc_hash_md5.cl"
 
+typedef struct ikepsk
+{
+  u32 nr_buf[16];
+  u32 nr_len;
+
+  u32 msg_buf[128];
+  u32 msg_len[6];
+
+} ikepsk_t;
+
 DECLSPEC void hmac_md5_pad (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad)
 {
   w0[0] = w0[0] ^ 0x36363636;
diff --git a/OpenCL/m05300_a1-pure.cl b/OpenCL/m05300_a1-pure.cl
index 39d8019e2..3b6d61142 100644
--- a/OpenCL/m05300_a1-pure.cl
+++ b/OpenCL/m05300_a1-pure.cl
@@ -13,6 +13,16 @@
 #include "inc_scalar.cl"
 #include "inc_hash_md5.cl"
 
+typedef struct ikepsk
+{
+  u32 nr_buf[16];
+  u32 nr_len;
+
+  u32 msg_buf[128];
+  u32 msg_len[6];
+
+} ikepsk_t;
+
 __kernel void m05300_mxx (KERN_ATTR_ESALT (ikepsk_t))
 {
   /**
diff --git a/OpenCL/m05300_a3-optimized.cl b/OpenCL/m05300_a3-optimized.cl
index 50f35935a..fe953a59e 100644
--- a/OpenCL/m05300_a3-optimized.cl
+++ b/OpenCL/m05300_a3-optimized.cl
@@ -13,6 +13,16 @@
 #include "inc_simd.cl"
 #include "inc_hash_md5.cl"
 
+typedef struct ikepsk
+{
+  u32 nr_buf[16];
+  u32 nr_len;
+
+  u32 msg_buf[128];
+  u32 msg_len[6];
+
+} ikepsk_t;
+
 DECLSPEC void hmac_md5_pad (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad)
 {
   w0[0] = w0[0] ^ 0x36363636;
diff --git a/OpenCL/m05300_a3-pure.cl b/OpenCL/m05300_a3-pure.cl
index f9b4b3bca..aa387c52c 100644
--- a/OpenCL/m05300_a3-pure.cl
+++ b/OpenCL/m05300_a3-pure.cl
@@ -13,6 +13,16 @@
 #include "inc_scalar.cl"
 #include "inc_hash_md5.cl"
 
+typedef struct ikepsk
+{
+  u32 nr_buf[16];
+  u32 nr_len;
+
+  u32 msg_buf[128];
+  u32 msg_len[6];
+
+} ikepsk_t;
+
 __kernel void m05300_mxx (KERN_ATTR_VECTOR_ESALT (ikepsk_t))
 {
   /**
diff --git a/OpenCL/m05400_a0-optimized.cl b/OpenCL/m05400_a0-optimized.cl
index 9ec74d7cb..c37980d8c 100644
--- a/OpenCL/m05400_a0-optimized.cl
+++ b/OpenCL/m05400_a0-optimized.cl
@@ -15,6 +15,16 @@
 #include "inc_simd.cl"
 #include "inc_hash_sha1.cl"
 
+typedef struct ikepsk
+{
+  u32 nr_buf[16];
+  u32 nr_len;
+
+  u32 msg_buf[128];
+  u32 msg_len[6];
+
+} ikepsk_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/m05400_a0-pure.cl b/OpenCL/m05400_a0-pure.cl
index 8a0e1f80b..3180e0011 100644
--- a/OpenCL/m05400_a0-pure.cl
+++ b/OpenCL/m05400_a0-pure.cl
@@ -15,6 +15,16 @@
 #include "inc_scalar.cl"
 #include "inc_hash_sha1.cl"
 
+typedef struct ikepsk
+{
+  u32 nr_buf[16];
+  u32 nr_len;
+
+  u32 msg_buf[128];
+  u32 msg_len[6];
+
+} ikepsk_t;
+
 __kernel void m05400_mxx (KERN_ATTR_RULES_ESALT (ikepsk_t))
 {
   /**
diff --git a/OpenCL/m05400_a1-optimized.cl b/OpenCL/m05400_a1-optimized.cl
index 67aa1fba7..08c2b6884 100644
--- a/OpenCL/m05400_a1-optimized.cl
+++ b/OpenCL/m05400_a1-optimized.cl
@@ -13,6 +13,16 @@
 #include "inc_simd.cl"
 #include "inc_hash_sha1.cl"
 
+typedef struct ikepsk
+{
+  u32 nr_buf[16];
+  u32 nr_len;
+
+  u32 msg_buf[128];
+  u32 msg_len[6];
+
+} ikepsk_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/m05400_a1-pure.cl b/OpenCL/m05400_a1-pure.cl
index ff33647d2..12a2dcf52 100644
--- a/OpenCL/m05400_a1-pure.cl
+++ b/OpenCL/m05400_a1-pure.cl
@@ -13,6 +13,16 @@
 #include "inc_scalar.cl"
 #include "inc_hash_sha1.cl"
 
+typedef struct ikepsk
+{
+  u32 nr_buf[16];
+  u32 nr_len;
+
+  u32 msg_buf[128];
+  u32 msg_len[6];
+
+} ikepsk_t;
+
 __kernel void m05400_mxx (KERN_ATTR_ESALT (ikepsk_t))
 {
   /**
diff --git a/OpenCL/m05400_a3-optimized.cl b/OpenCL/m05400_a3-optimized.cl
index 7cab76cf7..88b9d1932 100644
--- a/OpenCL/m05400_a3-optimized.cl
+++ b/OpenCL/m05400_a3-optimized.cl
@@ -13,6 +13,16 @@
 #include "inc_simd.cl"
 #include "inc_hash_sha1.cl"
 
+typedef struct ikepsk
+{
+  u32 nr_buf[16];
+  u32 nr_len;
+
+  u32 msg_buf[128];
+  u32 msg_len[6];
+
+} ikepsk_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/m05400_a3-pure.cl b/OpenCL/m05400_a3-pure.cl
index 92992936f..0d9815a34 100644
--- a/OpenCL/m05400_a3-pure.cl
+++ b/OpenCL/m05400_a3-pure.cl
@@ -13,6 +13,16 @@
 #include "inc_scalar.cl"
 #include "inc_hash_sha1.cl"
 
+typedef struct ikepsk
+{
+  u32 nr_buf[16];
+  u32 nr_len;
+
+  u32 msg_buf[128];
+  u32 msg_len[6];
+
+} ikepsk_t;
+
 __kernel void m05400_mxx (KERN_ATTR_VECTOR_ESALT (ikepsk_t))
 {
   /**
diff --git a/OpenCL/m05500_a0-optimized.cl b/OpenCL/m05500_a0-optimized.cl
index 4cf1f0007..597f79314 100644
--- a/OpenCL/m05500_a0-optimized.cl
+++ b/OpenCL/m05500_a0-optimized.cl
@@ -14,6 +14,18 @@
 #include "inc_rp_optimized.cl"
 #include "inc_simd.cl"
 
+typedef struct netntlm
+{
+  u32 user_len;
+  u32 domain_len;
+  u32 srvchall_len;
+  u32 clichall_len;
+
+  u32 userdomain_buf[64];
+  u32 chall_buf[256];
+
+} netntlm_t;
+
 #define PERM_OP(a,b,tt,n,m) \
 {                           \
   tt = a >> n;              \
diff --git a/OpenCL/m05500_a0-pure.cl b/OpenCL/m05500_a0-pure.cl
index 6e8302af4..caf590513 100644
--- a/OpenCL/m05500_a0-pure.cl
+++ b/OpenCL/m05500_a0-pure.cl
@@ -15,6 +15,18 @@
 #include "inc_scalar.cl"
 #include "inc_hash_md4.cl"
 
+typedef struct netntlm
+{
+  u32 user_len;
+  u32 domain_len;
+  u32 srvchall_len;
+  u32 clichall_len;
+
+  u32 userdomain_buf[64];
+  u32 chall_buf[256];
+
+} netntlm_t;
+
 #define PERM_OP(a,b,tt,n,m) \
 {                           \
   tt = a >> n;              \
diff --git a/OpenCL/m05500_a1-optimized.cl b/OpenCL/m05500_a1-optimized.cl
index ed4a04033..8e7bd8370 100644
--- a/OpenCL/m05500_a1-optimized.cl
+++ b/OpenCL/m05500_a1-optimized.cl
@@ -12,6 +12,18 @@
 #include "inc_common.cl"
 #include "inc_simd.cl"
 
+typedef struct netntlm
+{
+  u32 user_len;
+  u32 domain_len;
+  u32 srvchall_len;
+  u32 clichall_len;
+
+  u32 userdomain_buf[64];
+  u32 chall_buf[256];
+
+} netntlm_t;
+
 #define PERM_OP(a,b,tt,n,m) \
 {                           \
   tt = a >> n;              \
diff --git a/OpenCL/m05500_a1-pure.cl b/OpenCL/m05500_a1-pure.cl
index e774bea51..89ee6bcef 100644
--- a/OpenCL/m05500_a1-pure.cl
+++ b/OpenCL/m05500_a1-pure.cl
@@ -13,6 +13,18 @@
 #include "inc_scalar.cl"
 #include "inc_hash_md4.cl"
 
+typedef struct netntlm
+{
+  u32 user_len;
+  u32 domain_len;
+  u32 srvchall_len;
+  u32 clichall_len;
+
+  u32 userdomain_buf[64];
+  u32 chall_buf[256];
+
+} netntlm_t;
+
 #define PERM_OP(a,b,tt,n,m) \
 {                           \
   tt = a >> n;              \
diff --git a/OpenCL/m05500_a3-optimized.cl b/OpenCL/m05500_a3-optimized.cl
index 5fb88b555..ca0aae8ca 100644
--- a/OpenCL/m05500_a3-optimized.cl
+++ b/OpenCL/m05500_a3-optimized.cl
@@ -12,6 +12,18 @@
 #include "inc_common.cl"
 #include "inc_simd.cl"
 
+typedef struct netntlm
+{
+  u32 user_len;
+  u32 domain_len;
+  u32 srvchall_len;
+  u32 clichall_len;
+
+  u32 userdomain_buf[64];
+  u32 chall_buf[256];
+
+} netntlm_t;
+
 #define PERM_OP(a,b,tt,n,m) \
 {                           \
   tt = a >> n;              \
diff --git a/OpenCL/m05500_a3-pure.cl b/OpenCL/m05500_a3-pure.cl
index 04ce63171..7d5fa0529 100644
--- a/OpenCL/m05500_a3-pure.cl
+++ b/OpenCL/m05500_a3-pure.cl
@@ -13,6 +13,18 @@
 #include "inc_simd.cl"
 #include "inc_hash_md4.cl"
 
+typedef struct netntlm
+{
+  u32 user_len;
+  u32 domain_len;
+  u32 srvchall_len;
+  u32 clichall_len;
+
+  u32 userdomain_buf[64];
+  u32 chall_buf[256];
+
+} netntlm_t;
+
 #define PERM_OP(a,b,tt,n,m) \
 {                           \
   tt = a >> n;              \
diff --git a/OpenCL/m05600_a0-optimized.cl b/OpenCL/m05600_a0-optimized.cl
index 7fc122ca8..9b088441a 100644
--- a/OpenCL/m05600_a0-optimized.cl
+++ b/OpenCL/m05600_a0-optimized.cl
@@ -16,6 +16,18 @@
 #include "inc_hash_md4.cl"
 #include "inc_hash_md5.cl"
 
+typedef struct netntlm
+{
+  u32 user_len;
+  u32 domain_len;
+  u32 srvchall_len;
+  u32 clichall_len;
+
+  u32 userdomain_buf[64];
+  u32 chall_buf[256];
+
+} netntlm_t;
+
 DECLSPEC void hmac_md5_pad (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad)
 {
   w0[0] = w0[0] ^ 0x36363636;
diff --git a/OpenCL/m05600_a0-pure.cl b/OpenCL/m05600_a0-pure.cl
index 4440ed7de..4fbfded94 100644
--- a/OpenCL/m05600_a0-pure.cl
+++ b/OpenCL/m05600_a0-pure.cl
@@ -16,6 +16,18 @@
 #include "inc_hash_md4.cl"
 #include "inc_hash_md5.cl"
 
+typedef struct netntlm
+{
+  u32 user_len;
+  u32 domain_len;
+  u32 srvchall_len;
+  u32 clichall_len;
+
+  u32 userdomain_buf[64];
+  u32 chall_buf[256];
+
+} netntlm_t;
+
 __kernel void m05600_mxx (KERN_ATTR_RULES_ESALT (netntlm_t))
 {
   /**
diff --git a/OpenCL/m05600_a1-optimized.cl b/OpenCL/m05600_a1-optimized.cl
index 62a424ee7..86978d520 100644
--- a/OpenCL/m05600_a1-optimized.cl
+++ b/OpenCL/m05600_a1-optimized.cl
@@ -14,6 +14,18 @@
 #include "inc_hash_md4.cl"
 #include "inc_hash_md5.cl"
 
+typedef struct netntlm
+{
+  u32 user_len;
+  u32 domain_len;
+  u32 srvchall_len;
+  u32 clichall_len;
+
+  u32 userdomain_buf[64];
+  u32 chall_buf[256];
+
+} netntlm_t;
+
 DECLSPEC void hmac_md5_pad (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad)
 {
   w0[0] = w0[0] ^ 0x36363636;
diff --git a/OpenCL/m05600_a1-pure.cl b/OpenCL/m05600_a1-pure.cl
index 2f1204932..644fde5c6 100644
--- a/OpenCL/m05600_a1-pure.cl
+++ b/OpenCL/m05600_a1-pure.cl
@@ -14,6 +14,18 @@
 #include "inc_hash_md4.cl"
 #include "inc_hash_md5.cl"
 
+typedef struct netntlm
+{
+  u32 user_len;
+  u32 domain_len;
+  u32 srvchall_len;
+  u32 clichall_len;
+
+  u32 userdomain_buf[64];
+  u32 chall_buf[256];
+
+} netntlm_t;
+
 __kernel void m05600_mxx (KERN_ATTR_ESALT (netntlm_t))
 {
   /**
diff --git a/OpenCL/m05600_a3-optimized.cl b/OpenCL/m05600_a3-optimized.cl
index 2f9904e1a..eb30a65dd 100644
--- a/OpenCL/m05600_a3-optimized.cl
+++ b/OpenCL/m05600_a3-optimized.cl
@@ -14,6 +14,18 @@
 #include "inc_hash_md4.cl"
 #include "inc_hash_md5.cl"
 
+typedef struct netntlm
+{
+  u32 user_len;
+  u32 domain_len;
+  u32 srvchall_len;
+  u32 clichall_len;
+
+  u32 userdomain_buf[64];
+  u32 chall_buf[256];
+
+} netntlm_t;
+
 DECLSPEC void hmac_md5_pad (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad)
 {
   w0[0] = w0[0] ^ 0x36363636;
diff --git a/OpenCL/m05600_a3-pure.cl b/OpenCL/m05600_a3-pure.cl
index 5b9515b14..d90b6551f 100644
--- a/OpenCL/m05600_a3-pure.cl
+++ b/OpenCL/m05600_a3-pure.cl
@@ -14,6 +14,18 @@
 #include "inc_hash_md4.cl"
 #include "inc_hash_md5.cl"
 
+typedef struct netntlm
+{
+  u32 user_len;
+  u32 domain_len;
+  u32 srvchall_len;
+  u32 clichall_len;
+
+  u32 userdomain_buf[64];
+  u32 chall_buf[256];
+
+} netntlm_t;
+
 __kernel void m05600_mxx (KERN_ATTR_VECTOR_ESALT (netntlm_t))
 {
   /**
diff --git a/OpenCL/m05800-optimized.cl b/OpenCL/m05800-optimized.cl
index ecf699a7a..661ed576d 100644
--- a/OpenCL/m05800-optimized.cl
+++ b/OpenCL/m05800-optimized.cl
@@ -13,6 +13,12 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
+typedef struct androidpin_tmp
+{
+  u32 digest_buf[5];
+
+} androidpin_tmp_t;
+
 __constant u32a c_pc_dec[1024] =
 {
   0x00000030,
diff --git a/OpenCL/m05800-pure.cl b/OpenCL/m05800-pure.cl
index a757c3ef7..07547afee 100644
--- a/OpenCL/m05800-pure.cl
+++ b/OpenCL/m05800-pure.cl
@@ -13,6 +13,12 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
+typedef struct androidpin_tmp
+{
+  u32 digest_buf[5];
+
+} androidpin_tmp_t;
+
 __constant u32a c_pc_dec[1024] =
 {
   0x00000030,
diff --git a/OpenCL/m06211-pure.cl b/OpenCL/m06211-pure.cl
index 6a43054fc..8545dd466 100644
--- a/OpenCL/m06211-pure.cl
+++ b/OpenCL/m06211-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_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/m06212-pure.cl b/OpenCL/m06212-pure.cl
index b1c675fa3..b7e799807 100644
--- a/OpenCL/m06212-pure.cl
+++ b/OpenCL/m06212-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_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/m06213-pure.cl b/OpenCL/m06213-pure.cl
index a15ef1298..5027283be 100644
--- a/OpenCL/m06213-pure.cl
+++ b/OpenCL/m06213-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_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/m06221-pure.cl b/OpenCL/m06221-pure.cl
index 1d25cd197..3daf5d7a9 100644
--- a/OpenCL/m06221-pure.cl
+++ b/OpenCL/m06221-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 tc64_tmp
+{
+  u64  ipad[8];
+  u64  opad[8];
+
+  u64  dgst[32];
+  u64  out[32];
+
+} tc64_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/m06222-pure.cl b/OpenCL/m06222-pure.cl
index 8d4879b22..c3a8eb62b 100644
--- a/OpenCL/m06222-pure.cl
+++ b/OpenCL/m06222-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 tc64_tmp
+{
+  u64  ipad[8];
+  u64  opad[8];
+
+  u64  dgst[32];
+  u64  out[32];
+
+} tc64_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/m06223-pure.cl b/OpenCL/m06223-pure.cl
index d6bbdec1e..df71f8908 100644
--- a/OpenCL/m06223-pure.cl
+++ b/OpenCL/m06223-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 tc64_tmp
+{
+  u64  ipad[8];
+  u64  opad[8];
+
+  u64  dgst[32];
+  u64  out[32];
+
+} tc64_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/m06231-pure.cl b/OpenCL/m06231-pure.cl
index 76b689a31..5ed0b0d96 100644
--- a/OpenCL/m06231-pure.cl
+++ b/OpenCL/m06231-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_whirlpool_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest, __local u32 (*s_Ch)[256], __local u32 (*s_Cl)[256])
 {
   digest[ 0] = ipad[ 0];
diff --git a/OpenCL/m06232-pure.cl b/OpenCL/m06232-pure.cl
index 8b0e40aae..6efb5318d 100644
--- a/OpenCL/m06232-pure.cl
+++ b/OpenCL/m06232-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_whirlpool_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest, __local u32 (*s_Ch)[256], __local u32 (*s_Cl)[256])
 {
   digest[ 0] = ipad[ 0];
diff --git a/OpenCL/m06233-pure.cl b/OpenCL/m06233-pure.cl
index f4992e362..f44255221 100644
--- a/OpenCL/m06233-pure.cl
+++ b/OpenCL/m06233-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_whirlpool_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest, __local u32 (*s_Ch)[256], __local u32 (*s_Cl)[256])
 {
   digest[ 0] = ipad[ 0];
diff --git a/OpenCL/m06300-optimized.cl b/OpenCL/m06300-optimized.cl
index 1f5aceedc..72442a94b 100644
--- a/OpenCL/m06300-optimized.cl
+++ b/OpenCL/m06300-optimized.cl
@@ -13,6 +13,12 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
+typedef struct md5crypt_tmp
+{
+  u32 digest_buf[4];
+
+} md5crypt_tmp_t;
+
 DECLSPEC void memcat16 (u32 *block0, u32 *block1, u32 *block2, u32 *block3, const u32 offset, const u32 *append)
 {
   u32 tmp0;
diff --git a/OpenCL/m06300-pure.cl b/OpenCL/m06300-pure.cl
index a3e39ca90..d81b7dfe1 100644
--- a/OpenCL/m06300-pure.cl
+++ b/OpenCL/m06300-pure.cl
@@ -18,6 +18,12 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
+typedef struct md5crypt_tmp
+{
+  u32 digest_buf[4];
+
+} md5crypt_tmp_t;
+
 __kernel void m06300_init (KERN_ATTR_TMPS (md5crypt_tmp_t))
 {
   /**