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

OpenCL/inc_types.cl

@@ -1329,3 +1329,12 @@ typedef struct plain
   u32  hash_pos;
 
 } plain_t;
+
+typedef struct keyboard_layout_mapping
+{
+  u32 src_char;
+  int src_len;
+  u32 dst_char;
+  int dst_len;
+
+} keyboard_layout_mapping_t;

OpenCL/m06400-pure.cl

@@ -16,6 +16,16 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
+typedef struct sha256aix_tmp
+{
+  u32 ipad[8];
+  u32 opad[8];
+
+  u32 dgst[8];
+  u32 out[8];
+
+} sha256aix_tmp_t;
+
 DECLSPEC void hmac_sha256_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest)
 {
   digest[0] = ipad[0];

OpenCL/m06500-pure.cl

@@ -16,6 +16,16 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
+typedef struct sha512aix_tmp
+{
+  u64  ipad[8];
+  u64  opad[8];
+
+  u64  dgst[8];
+  u64  out[8];
+
+} sha512aix_tmp_t;
+
 DECLSPEC void hmac_sha512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *w4, u32x *w5, u32x *w6, u32x *w7, u64x *ipad, u64x *opad, u64x *digest)
 {
   digest[0] = ipad[0];

OpenCL/m06600-pure.cl

@@ -17,6 +17,16 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
+typedef struct agilekey_tmp
+{
+  u32 ipad[5];
+  u32 opad[5];
+
+  u32 dgst[5];
+  u32 out[5];
+
+} agilekey_tmp_t;
+
 DECLSPEC void hmac_sha1_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest)
 {
   digest[0] = ipad[0];

OpenCL/m06700-pure.cl

@@ -16,6 +16,16 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
+typedef struct sha1aix_tmp
+{
+  u32 ipad[5];
+  u32 opad[5];
+
+  u32 dgst[5];
+  u32 out[5];
+
+} sha1aix_tmp_t;
+
 DECLSPEC void hmac_sha1_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest)
 {
   digest[0] = ipad[0];

OpenCL/m06800-pure.cl

@@ -17,6 +17,16 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
+typedef struct lastpass_tmp
+{
+  u32 ipad[8];
+  u32 opad[8];
+
+  u32 dgst[8];
+  u32 out[8];
+
+} lastpass_tmp_t;
+
 DECLSPEC void hmac_sha256_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest)
 {
   digest[0] = ipad[0];

OpenCL/m07100-pure.cl

@@ -16,6 +16,22 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
+typedef struct pbkdf2_sha512_tmp
+{
+  u64  ipad[8];
+  u64  opad[8];
+
+  u64  dgst[16];
+  u64  out[16];
+
+} pbkdf2_sha512_tmp_t;
+
+typedef struct pbkdf2_sha512
+{
+  u32 salt_buf[64];
+
+} pbkdf2_sha512_t;
+
 DECLSPEC void hmac_sha512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *w4, u32x *w5, u32x *w6, u32x *w7, u64x *ipad, u64x *opad, u64x *digest)
 {
   digest[0] = ipad[0];

OpenCL/m07300_a0-optimized.cl

@@ -15,6 +15,13 @@
 #include "inc_simd.cl"
 #include "inc_hash_sha1.cl"
 
+typedef struct rakp
+{
+  u32 salt_buf[128];
+  u32 salt_len;
+
+} rakp_t;
+
 DECLSPEC void hmac_sha1_pad (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad)
 {
   w0[0] = w0[0] ^ 0x36363636;

OpenCL/m07300_a0-pure.cl

@@ -15,6 +15,13 @@
 #include "inc_scalar.cl"
 #include "inc_hash_sha1.cl"
 
+typedef struct rakp
+{
+  u32 salt_buf[128];
+  u32 salt_len;
+
+} rakp_t;
+
 __kernel void m07300_mxx (KERN_ATTR_RULES_ESALT (rakp_t))
 {
   /**

OpenCL/m07300_a1-optimized.cl

@@ -13,6 +13,13 @@
 #include "inc_simd.cl"
 #include "inc_hash_sha1.cl"
 
+typedef struct rakp
+{
+  u32 salt_buf[128];
+  u32 salt_len;
+
+} rakp_t;
+
 DECLSPEC void hmac_sha1_pad (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad)
 {
   w0[0] = w0[0] ^ 0x36363636;

OpenCL/m07300_a1-pure.cl

@@ -13,6 +13,13 @@
 #include "inc_scalar.cl"
 #include "inc_hash_sha1.cl"
 
+typedef struct rakp
+{
+  u32 salt_buf[128];
+  u32 salt_len;
+
+} rakp_t;
+
 __kernel void m07300_mxx (KERN_ATTR_ESALT (rakp_t))
 {
   /**

OpenCL/m07300_a3-optimized.cl

@@ -13,6 +13,13 @@
 #include "inc_simd.cl"
 #include "inc_hash_sha1.cl"
 
+typedef struct rakp
+{
+  u32 salt_buf[128];
+  u32 salt_len;
+
+} rakp_t;
+
 DECLSPEC void hmac_sha1_pad (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad)
 {
   w0[0] = w0[0] ^ 0x36363636;

OpenCL/m07300_a3-pure.cl

@@ -13,6 +13,13 @@
 #include "inc_scalar.cl"
 #include "inc_hash_sha1.cl"
 
+typedef struct rakp
+{
+  u32 salt_buf[128];
+  u32 salt_len;
+
+} rakp_t;
+
 __kernel void m07300_mxx (KERN_ATTR_VECTOR_ESALT (rakp_t))
 {
   /**

OpenCL/m07400-optimized.cl

@@ -13,6 +13,16 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
+typedef struct sha256crypt_tmp
+{
+  // pure version
+
+  u32 alt_result[8];
+  u32 p_bytes[64];
+  u32 s_bytes[64];
+
+} sha256crypt_tmp_t;
+
 DECLSPEC void sha256_transform_transport (const u32 *w, u32 *digest)
 {
   u32 w0[4];

OpenCL/m07400-pure.cl

@@ -13,6 +13,16 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
+typedef struct sha256crypt_tmp
+{
+  // pure version
+
+  u32 alt_result[8];
+  u32 p_bytes[64];
+  u32 s_bytes[64];
+
+} sha256crypt_tmp_t;
+
 __kernel void m07400_init (KERN_ATTR_TMPS (sha256crypt_tmp_t))
 {
   /**

OpenCL/m07500_a0-optimized.cl

@@ -17,6 +17,16 @@
 #include "inc_hash_md4.cl"
 #include "inc_hash_md5.cl"
 
+typedef struct krb5pa
+{
+  u32 user[16];
+  u32 realm[16];
+  u32 salt[32];
+  u32 timestamp[16];
+  u32 checksum[4];
+
+} krb5pa_t;
+
 typedef struct
 {
   u8 S[256];

OpenCL/m07500_a0-pure.cl

@@ -16,6 +16,16 @@
 #include "inc_hash_md4.cl"
 #include "inc_hash_md5.cl"
 
+typedef struct krb5pa
+{
+  u32 user[16];
+  u32 realm[16];
+  u32 salt[32];
+  u32 timestamp[16];
+  u32 checksum[4];
+
+} krb5pa_t;
+
 typedef struct
 {
   u8 S[256];

OpenCL/m07500_a1-optimized.cl

@@ -15,6 +15,16 @@
 #include "inc_hash_md4.cl"
 #include "inc_hash_md5.cl"
 
+typedef struct krb5pa
+{
+  u32 user[16];
+  u32 realm[16];
+  u32 salt[32];
+  u32 timestamp[16];
+  u32 checksum[4];
+
+} krb5pa_t;
+
 typedef struct
 {
   u8 S[256];

OpenCL/m07500_a1-pure.cl

@@ -14,6 +14,16 @@
 #include "inc_hash_md4.cl"
 #include "inc_hash_md5.cl"
 
+typedef struct krb5pa
+{
+  u32 user[16];
+  u32 realm[16];
+  u32 salt[32];
+  u32 timestamp[16];
+  u32 checksum[4];
+
+} krb5pa_t;
+
 typedef struct
 {
   u8 S[256];

OpenCL/m07500_a3-optimized.cl

@@ -15,6 +15,16 @@
 #include "inc_hash_md4.cl"
 #include "inc_hash_md5.cl"
 
+typedef struct krb5pa
+{
+  u32 user[16];
+  u32 realm[16];
+  u32 salt[32];
+  u32 timestamp[16];
+  u32 checksum[4];
+
+} krb5pa_t;
+
 typedef struct
 {
   u8 S[256];

OpenCL/m07500_a3-pure.cl

@@ -14,6 +14,16 @@
 #include "inc_hash_md4.cl"
 #include "inc_hash_md5.cl"
 
+typedef struct krb5pa
+{
+  u32 user[16];
+  u32 realm[16];
+  u32 salt[32];
+  u32 timestamp[16];
+  u32 checksum[4];
+
+} krb5pa_t;
+
 typedef struct
 {
   u8 S[256];

OpenCL/m07900-pure.cl

@@ -13,6 +13,12 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
+typedef struct drupal7_tmp
+{
+  u64  digest_buf[8];
+
+} drupal7_tmp_t;
+
 __kernel void m07900_init (KERN_ATTR_TMPS (drupal7_tmp_t))
 {
   /**

OpenCL/m08200-pure.cl

@@ -17,6 +17,23 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
+typedef struct pbkdf2_sha512_tmp
+{
+  u64  ipad[8];
+  u64  opad[8];
+
+  u64  dgst[16];
+  u64  out[16];
+
+} pbkdf2_sha512_tmp_t;
+
+typedef struct cloudkey
+{
+  u32 data_len;
+  u32 data_buf[512];
+
+} cloudkey_t;
+
 DECLSPEC void hmac_sha512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *w4, u32x *w5, u32x *w6, u32x *w7, u64x *ipad, u64x *opad, u64x *digest)
 {
   digest[0] = ipad[0];

OpenCL/m08800-pure.cl

@@ -15,6 +15,22 @@
 #include "inc_hash_sha256.cl"
 #include "inc_cipher_aes.cl"
 
+typedef struct androidfde_tmp
+{
+  u32 ipad[5];
+  u32 opad[5];
+
+  u32 dgst[10];
+  u32 out[10];
+
+} androidfde_tmp_t;
+
+typedef struct androidfde
+{
+  u32 data[384];
+
+} androidfde_t;
+
 DECLSPEC void hmac_sha1_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest)
 {
   digest[0] = ipad[0];

OpenCL/m09000-pure.cl

@@ -12,6 +12,19 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
+typedef struct pwsafe2_tmp
+{
+  u32 digest[2];
+
+  u32 P[18];
+
+  u32 S0[256];
+  u32 S1[256];
+  u32 S2[256];
+  u32 S3[256];
+
+} pwsafe2_tmp_t;
+
 // http://www.schneier.com/code/constants.txt
 
 __constant u32a c_sbox0[256] =

OpenCL/m09100-pure.cl

@@ -16,6 +16,16 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
+typedef struct lotus8_tmp
+{
+  u32 ipad[5];
+  u32 opad[5];
+
+  u32 dgst[5];
+  u32 out[5];
+
+} lotus8_tmp_t;
+
 __constant u32 lotus64_table[64] =
 {
   '0', '1', '2', '3', '4', '5', '6', '7',

OpenCL/m09400-pure.cl

@@ -17,6 +17,21 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
+typedef struct office2007_tmp
+{
+  u32 out[5];
+
+} office2007_tmp_t;
+
+typedef struct office2007
+{
+  u32 encryptedVerifier[4];
+  u32 encryptedVerifierHash[5];
+
+  u32 keySize;
+
+} office2007_t;
+
 __kernel void m09400_init (KERN_ATTR_TMPS_ESALT (office2007_tmp_t, office2007_t))
 {
   /**

OpenCL/m09500-pure.cl

@@ -17,6 +17,19 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
+typedef struct office2010
+{
+  u32 encryptedVerifier[4];
+  u32 encryptedVerifierHash[8];
+
+} office2010_t;
+
+typedef struct office2010_tmp
+{
+  u32 out[5];
+
+} office2010_tmp_t;
+
 __kernel void m09500_init (KERN_ATTR_TMPS_ESALT (office2010_tmp_t, office2010_t))
 {
   /**

OpenCL/m09600-pure.cl

@@ -17,6 +17,19 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
+typedef struct office2013_tmp
+{
+  u64  out[8];
+
+} office2013_tmp_t;
+
+typedef struct office2013
+{
+  u32 encryptedVerifier[4];
+  u32 encryptedVerifierHash[8];
+
+} office2013_t;
+
 __kernel void m09600_init (KERN_ATTR_TMPS_ESALT (office2013_tmp_t, office2013_t))
 {
   /**

OpenCL/m09700_a0-optimized.cl

@@ -16,6 +16,15 @@
 #include "inc_simd.cl"
 #include "inc_hash_md5.cl"
 
+typedef struct oldoffice01
+{
+  u32 version;
+  u32 encryptedVerifier[4];
+  u32 encryptedVerifierHash[4];
+  u32 rc4key[2];
+
+} oldoffice01_t;
+
 typedef struct
 {
   u8 S[256];

OpenCL/m09700_a1-optimized.cl

@@ -14,6 +14,15 @@
 #include "inc_simd.cl"
 #include "inc_hash_md5.cl"
 
+typedef struct oldoffice01
+{
+  u32 version;
+  u32 encryptedVerifier[4];
+  u32 encryptedVerifierHash[4];
+  u32 rc4key[2];
+
+} oldoffice01_t;
+
 typedef struct
 {
   u8 S[256];

OpenCL/m09700_a3-optimized.cl

@@ -11,6 +11,15 @@
 #include "inc_simd.cl"
 #include "inc_hash_md5.cl"
 
+typedef struct oldoffice01
+{
+  u32 version;
+  u32 encryptedVerifier[4];
+  u32 encryptedVerifierHash[4];
+  u32 rc4key[2];
+
+} oldoffice01_t;
+
 typedef struct
 {
   u8 S[256];

OpenCL/m09710_a0-optimized.cl

@@ -16,6 +16,15 @@
 #include "inc_simd.cl"
 #include "inc_hash_md5.cl"
 
+typedef struct oldoffice01
+{
+  u32 version;
+  u32 encryptedVerifier[4];
+  u32 encryptedVerifierHash[4];
+  u32 rc4key[2];
+
+} oldoffice01_t;
+
 typedef struct
 {
   u8 S[256];

OpenCL/m09710_a1-optimized.cl

@@ -14,6 +14,15 @@
 #include "inc_simd.cl"
 #include "inc_hash_md5.cl"
 
+typedef struct oldoffice01
+{
+  u32 version;
+  u32 encryptedVerifier[4];
+  u32 encryptedVerifierHash[4];
+  u32 rc4key[2];
+
+} oldoffice01_t;
+
 typedef struct
 {
   u8 S[256];

OpenCL/m09710_a3-optimized.cl

@@ -14,6 +14,15 @@
 #include "inc_simd.cl"
 #include "inc_hash_md5.cl"
 
+typedef struct oldoffice01
+{
+  u32 version;
+  u32 encryptedVerifier[4];
+  u32 encryptedVerifierHash[4];
+  u32 rc4key[2];
+
+} oldoffice01_t;
+
 typedef struct
 {
   u8 S[256];

OpenCL/m09720_a0-optimized.cl

@@ -15,6 +15,15 @@
 #include "inc_simd.cl"
 #include "inc_hash_md5.cl"
 
+typedef struct oldoffice01
+{
+  u32 version;
+  u32 encryptedVerifier[4];
+  u32 encryptedVerifierHash[4];
+  u32 rc4key[2];
+
+} oldoffice01_t;
+
 DECLSPEC void gen336 (u32x *digest_pre, u32 *salt_buf, u32x *digest)
 {
   u32x digest_t0[2];

OpenCL/m09720_a1-optimized.cl

@@ -13,6 +13,15 @@
 #include "inc_simd.cl"
 #include "inc_hash_md5.cl"
 
+typedef struct oldoffice01
+{
+  u32 version;
+  u32 encryptedVerifier[4];
+  u32 encryptedVerifierHash[4];
+  u32 rc4key[2];
+
+} oldoffice01_t;
+
 DECLSPEC void gen336 (u32x *digest_pre, u32 *salt_buf, u32x *digest)
 {
   u32x digest_t0[2];

OpenCL/m09720_a3-optimized.cl

@@ -13,6 +13,15 @@
 #include "inc_simd.cl"
 #include "inc_hash_md5.cl"
 
+typedef struct oldoffice01
+{
+  u32 version;
+  u32 encryptedVerifier[4];
+  u32 encryptedVerifierHash[4];
+  u32 rc4key[2];
+
+} oldoffice01_t;
+
 DECLSPEC void gen336 (u32x *digest_pre, u32 *salt_buf, u32x *digest)
 {
   u32x digest_t0[2];

OpenCL/m09800_a0-optimized.cl

@@ -16,6 +16,15 @@
 #include "inc_simd.cl"
 #include "inc_hash_sha1.cl"
 
+typedef struct oldoffice34
+{
+  u32 version;
+  u32 encryptedVerifier[4];
+  u32 encryptedVerifierHash[5];
+  u32 rc4key[2];
+
+} oldoffice34_t;
+
 typedef struct
 {
   u8 S[256];

OpenCL/m09800_a1-optimized.cl

@@ -14,6 +14,15 @@
 #include "inc_simd.cl"
 #include "inc_hash_sha1.cl"
 
+typedef struct oldoffice34
+{
+  u32 version;
+  u32 encryptedVerifier[4];
+  u32 encryptedVerifierHash[5];
+  u32 rc4key[2];
+
+} oldoffice34_t;
+
 typedef struct
 {
   u8 S[256];

OpenCL/m09800_a3-optimized.cl

@@ -11,6 +11,15 @@
 #include "inc_simd.cl"
 #include "inc_hash_sha1.cl"
 
+typedef struct oldoffice34
+{
+  u32 version;
+  u32 encryptedVerifier[4];
+  u32 encryptedVerifierHash[5];
+  u32 rc4key[2];
+
+} oldoffice34_t;
+
 typedef struct
 {
   u8 S[256];

OpenCL/m09810_a0-optimized.cl

@@ -16,6 +16,15 @@
 #include "inc_simd.cl"
 #include "inc_hash_sha1.cl"
 
+typedef struct oldoffice34
+{
+  u32 version;
+  u32 encryptedVerifier[4];
+  u32 encryptedVerifierHash[5];
+  u32 rc4key[2];
+
+} oldoffice34_t;
+
 typedef struct
 {
   u8 S[256];

OpenCL/m09810_a1-optimized.cl

@@ -14,6 +14,15 @@
 #include "inc_simd.cl"
 #include "inc_hash_sha1.cl"
 
+typedef struct oldoffice34
+{
+  u32 version;
+  u32 encryptedVerifier[4];
+  u32 encryptedVerifierHash[5];
+  u32 rc4key[2];
+
+} oldoffice34_t;
+
 typedef struct
 {
   u8 S[256];

OpenCL/m09810_a3-optimized.cl

@@ -14,6 +14,15 @@
 #include "inc_simd.cl"
 #include "inc_hash_sha1.cl"
 
+typedef struct oldoffice34
+{
+  u32 version;
+  u32 encryptedVerifier[4];
+  u32 encryptedVerifierHash[5];
+  u32 rc4key[2];
+
+} oldoffice34_t;
+
 typedef struct
 {
   u8 S[256];

OpenCL/m09820_a0-optimized.cl

@@ -15,6 +15,15 @@
 #include "inc_simd.cl"
 #include "inc_hash_sha1.cl"
 
+typedef struct oldoffice34
+{
+  u32 version;
+  u32 encryptedVerifier[4];
+  u32 encryptedVerifierHash[5];
+  u32 rc4key[2];
+
+} oldoffice34_t;
+
 __kernel void m09820_m04 (KERN_ATTR_RULES_ESALT (oldoffice34_t))
 {
   /**

OpenCL/m09820_a1-optimized.cl

@@ -13,6 +13,15 @@
 #include "inc_simd.cl"
 #include "inc_hash_sha1.cl"
 
+typedef struct oldoffice34
+{
+  u32 version;
+  u32 encryptedVerifier[4];
+  u32 encryptedVerifierHash[5];
+  u32 rc4key[2];
+
+} oldoffice34_t;
+
 __kernel void m09820_m04 (KERN_ATTR_ESALT (oldoffice34_t))
 {
   /**

OpenCL/m09820_a3-optimized.cl

@@ -13,6 +13,15 @@
 #include "inc_simd.cl"
 #include "inc_hash_sha1.cl"
 
+typedef struct oldoffice34
+{
+  u32 version;
+  u32 encryptedVerifier[4];
+  u32 encryptedVerifierHash[5];
+  u32 rc4key[2];
+
+} oldoffice34_t;
+
 DECLSPEC void m09820m (u32 *w0, u32 *w1, u32 *w2, u32 *w3, const u32 pw_len, KERN_ATTR_ESALT (oldoffice34_t))
 {
   /**

OpenCL/m10300-pure.cl

@@ -13,6 +13,12 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
+typedef struct saph_sha1_tmp
+{
+  u32 digest_buf[5];
+
+} saph_sha1_tmp_t;
+
 __kernel void m10300_init (KERN_ATTR_TMPS (saph_sha1_tmp_t))
 {
   /**

OpenCL/m10400_a0-optimized.cl

@@ -28,6 +28,27 @@ __constant u32a padding[8] =
   0x7a695364
 };
 
+typedef struct pdf
+{
+  int  V;
+  int  R;
+  int  P;
+
+  int  enc_md;
+
+  u32  id_buf[8];
+  u32  u_buf[32];
+  u32  o_buf[32];
+
+  int  id_len;
+  int  o_len;
+  int  u_len;
+
+  u32  rc4key[2];
+  u32  rc4data[2];
+
+} pdf_t;
+
 typedef struct
 {
   u8 S[256];

OpenCL/m10400_a1-optimized.cl

@@ -26,6 +26,27 @@ __constant u32a padding[8] =
   0x7a695364
 };
 
+typedef struct pdf
+{
+  int  V;
+  int  R;
+  int  P;
+
+  int  enc_md;
+
+  u32  id_buf[8];
+  u32  u_buf[32];
+  u32  o_buf[32];
+
+  int  id_len;
+  int  o_len;
+  int  u_len;
+
+  u32  rc4key[2];
+  u32  rc4data[2];
+
+} pdf_t;
+
 typedef struct
 {
   u8 S[256];

OpenCL/m10400_a3-optimized.cl

@@ -26,6 +26,27 @@ __constant u32a padding[8] =
   0x7a695364
 };
 
+typedef struct pdf
+{
+  int  V;
+  int  R;
+  int  P;
+
+  int  enc_md;
+
+  u32  id_buf[8];
+  u32  u_buf[32];
+  u32  o_buf[32];
+
+  int  id_len;
+  int  o_len;
+  int  u_len;
+
+  u32  rc4key[2];
+  u32  rc4data[2];
+
+} pdf_t;
+
 typedef struct
 {
   u8 S[256];

OpenCL/m10410_a0-optimized.cl

@@ -28,6 +28,27 @@ __constant u32a padding[8] =
   0x7a695364
 };
 
+typedef struct pdf
+{
+  int  V;
+  int  R;
+  int  P;
+
+  int  enc_md;
+
+  u32  id_buf[8];
+  u32  u_buf[32];
+  u32  o_buf[32];
+
+  int  id_len;
+  int  o_len;
+  int  u_len;
+
+  u32  rc4key[2];
+  u32  rc4data[2];
+
+} pdf_t;
+
 typedef struct
 {
   u8 S[256];

OpenCL/m10410_a1-optimized.cl

@@ -26,6 +26,27 @@ __constant u32a padding[8] =
   0x7a695364
 };
 
+typedef struct pdf
+{
+  int  V;
+  int  R;
+  int  P;
+
+  int  enc_md;
+
+  u32  id_buf[8];
+  u32  u_buf[32];
+  u32  o_buf[32];
+
+  int  id_len;
+  int  o_len;
+  int  u_len;
+
+  u32  rc4key[2];
+  u32  rc4data[2];
+
+} pdf_t;
+
 typedef struct
 {
   u8 S[256];

OpenCL/m10410_a3-optimized.cl

@@ -26,6 +26,27 @@ __constant u32a padding[8] =
   0x7a695364
 };
 
+typedef struct pdf
+{
+  int  V;
+  int  R;
+  int  P;
+
+  int  enc_md;
+
+  u32  id_buf[8];
+  u32  u_buf[32];
+  u32  o_buf[32];
+
+  int  id_len;
+  int  o_len;
+  int  u_len;
+
+  u32  rc4key[2];
+  u32  rc4data[2];
+
+} pdf_t;
+
 typedef struct
 {
   u8 S[256];

OpenCL/m10420_a0-optimized.cl

@@ -27,6 +27,27 @@ __constant u32a padding[8] =
   0x7a695364
 };
 
+typedef struct pdf
+{
+  int  V;
+  int  R;
+  int  P;
+
+  int  enc_md;
+
+  u32  id_buf[8];
+  u32  u_buf[32];
+  u32  o_buf[32];
+
+  int  id_len;
+  int  o_len;
+  int  u_len;
+
+  u32  rc4key[2];
+  u32  rc4data[2];
+
+} pdf_t;
+
 __kernel void m10420_m04 (KERN_ATTR_RULES_ESALT (pdf_t))
 {
   /**

OpenCL/m10420_a1-optimized.cl

@@ -25,6 +25,27 @@ __constant u32a padding[8] =
   0x7a695364
 };
 
+typedef struct pdf
+{
+  int  V;
+  int  R;
+  int  P;
+
+  int  enc_md;
+
+  u32  id_buf[8];
+  u32  u_buf[32];
+  u32  o_buf[32];
+
+  int  id_len;
+  int  o_len;
+  int  u_len;
+
+  u32  rc4key[2];
+  u32  rc4data[2];
+
+} pdf_t;
+
 __kernel void m10420_m04 (KERN_ATTR_ESALT (pdf_t))
 {
   /**

OpenCL/m10420_a3-optimized.cl

@@ -25,6 +25,27 @@ __constant u32a padding[8] =
   0x7a695364
 };
 
+typedef struct pdf
+{
+  int  V;
+  int  R;
+  int  P;
+
+  int  enc_md;
+
+  u32  id_buf[8];
+  u32  u_buf[32];
+  u32  o_buf[32];
+
+  int  id_len;
+  int  o_len;
+  int  u_len;
+
+  u32  rc4key[2];
+  u32  rc4data[2];
+
+} pdf_t;
+
 DECLSPEC void m10420m (u32 *w0, u32 *w1, u32 *w2, u32 *w3, const u32 pw_len, KERN_ATTR_ESALT (pdf_t))
 {
   /**

OpenCL/m10500-pure.cl

@@ -25,6 +25,34 @@ __constant u32a padding[8] =
   0x7a695364
 };
 
+typedef struct pdf
+{
+  int  V;
+  int  R;
+  int  P;
+
+  int  enc_md;
+
+  u32  id_buf[8];
+  u32  u_buf[32];
+  u32  o_buf[32];
+
+  int  id_len;
+  int  o_len;
+  int  u_len;
+
+  u32  rc4key[2];
+  u32  rc4data[2];
+
+} pdf_t;
+
+typedef struct pdf14_tmp
+{
+  u32 digest[4];
+  u32 out[4];
+
+} pdf14_tmp_t;
+
 typedef struct
 {
   u8 S[256];

OpenCL/m10900-pure.cl

@@ -16,6 +16,22 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
+typedef struct pbkdf2_sha256_tmp
+{
+  u32  ipad[8];
+  u32  opad[8];
+
+  u32  dgst[32];
+  u32  out[32];
+
+} pbkdf2_sha256_tmp_t;
+
+typedef struct pbkdf2_sha256
+{
+  u32 salt_buf[64];
+
+} pbkdf2_sha256_t;
+
 DECLSPEC void hmac_sha256_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest)
 {
   digest[0] = ipad[0];

OpenCL/m11300-pure.cl

@@ -14,6 +14,19 @@
 #include "inc_hash_sha512.cl"
 #include "inc_cipher_aes.cl"
 
+typedef struct bitcoin_wallet_tmp
+{
+  u64  dgst[8];
+
+} bitcoin_wallet_tmp_t;
+
+typedef struct bitcoin_wallet
+{
+  u32 cry_master_buf[64];
+  u32 cry_master_len;
+
+} bitcoin_wallet_t;
+
 DECLSPEC void hmac_sha512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *w4, u32x *w5, u32x *w6, u32x *w7, u64x *ipad, u64x *opad, u64x *digest)
 {
   digest[0] = ipad[0];

OpenCL/m11400_a0-pure.cl

@@ -15,6 +15,16 @@
 #include "inc_scalar.cl"
 #include "inc_hash_md5.cl"
 
+typedef struct sip
+{
+  u32 salt_buf[32];
+  u32 salt_len;
+
+  u32 esalt_buf[256];
+  u32 esalt_len;
+
+} sip_t;
+
 #if   VECT_SIZE == 1
 #define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i)])
 #elif VECT_SIZE == 2

OpenCL/m11400_a1-pure.cl

@@ -13,6 +13,16 @@
 #include "inc_scalar.cl"
 #include "inc_hash_md5.cl"
 
+typedef struct sip
+{
+  u32 salt_buf[32];
+  u32 salt_len;
+
+  u32 esalt_buf[256];
+  u32 esalt_len;
+
+} sip_t;
+
 #if   VECT_SIZE == 1
 #define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i)])
 #elif VECT_SIZE == 2

OpenCL/m11400_a3-pure.cl

@@ -13,6 +13,16 @@
 #include "inc_simd.cl"
 #include "inc_hash_md5.cl"
 
+typedef struct sip
+{
+  u32 salt_buf[32];
+  u32 salt_len;
+
+  u32 esalt_buf[256];
+  u32 esalt_len;
+
+} sip_t;
+
 #if   VECT_SIZE == 1
 #define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i)])
 #elif VECT_SIZE == 2

OpenCL/m11600-pure.cl

@@ -10,6 +10,27 @@
 #include "inc_common.cl"
 #include "inc_hash_sha256.cl"
 
+typedef struct seven_zip_tmp
+{
+  u32 h[8];
+
+  u32 w0[4];
+  u32 w1[4];
+  u32 w2[4];
+  u32 w3[4];
+
+  int len;
+
+} seven_zip_tmp_t;
+
+typedef struct
+{
+  u32 ukey[8];
+
+  u32 hook_success;
+
+} seven_zip_hook_t;
+
 DECLSPEC void memcat8c_be (u32 *w0, u32 *w1, u32 *w2, u32 *w3, const u32 len, const u32 append, u32 *digest)
 {
   const u32 func_len = len & 63;

OpenCL/m11900-pure.cl

@@ -16,6 +16,22 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
+typedef struct pbkdf2_md5
+{
+  u32 salt_buf[64];
+
+} pbkdf2_md5_t;
+
+typedef struct pbkdf2_md5_tmp
+{
+  u32  ipad[4];
+  u32  opad[4];
+
+  u32  dgst[32];
+  u32  out[32];
+
+} pbkdf2_md5_tmp_t;
+
 DECLSPEC void hmac_md5_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest)
 {
   digest[0] = ipad[0];

OpenCL/m12000-pure.cl

@@ -16,6 +16,22 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
+typedef struct pbkdf2_sha1_tmp
+{
+  u32  ipad[5];
+  u32  opad[5];
+
+  u32  dgst[32];
+  u32  out[32];
+
+} pbkdf2_sha1_tmp_t;
+
+typedef struct pbkdf2_sha1
+{
+  u32 salt_buf[64];
+
+} pbkdf2_sha1_t;
+
 DECLSPEC void hmac_sha1_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest)
 {
   digest[0] = ipad[0];

OpenCL/m12200-pure.cl

@@ -16,6 +16,12 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
+typedef struct ecryptfs_tmp
+{
+  u64  out[8];
+
+} ecryptfs_tmp_t;
+
 __kernel void m12200_init (KERN_ATTR_TMPS (ecryptfs_tmp_t))
 {
   /**

OpenCL/m12300-pure.cl

@@ -16,6 +16,16 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
+typedef struct oraclet_tmp
+{
+  u64  ipad[8];
+  u64  opad[8];
+
+  u64  dgst[16];
+  u64  out[16];
+
+} oraclet_tmp_t;
+
 DECLSPEC void hmac_sha512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *w4, u32x *w5, u32x *w6, u32x *w7, u64x *ipad, u64x *opad, u64x *digest)
 {
   digest[0] = ipad[0];

OpenCL/m12500-pure.cl

@@ -23,6 +23,18 @@
 
 #define MIN(a,b) (((a) < (b)) ? (a) : (b))
 
+typedef struct pbkdf2_sha1
+{
+  u32 salt_buf[64];
+
+} pbkdf2_sha1_t;
+
+typedef struct rar3_tmp
+{
+  u32 dgst[17][5];
+
+} rar3_tmp_t;
+
 DECLSPEC void sha1_transform (const u32 *w, u32 *digest)
 {
   u32 A = digest[0];

OpenCL/m12700-pure.cl

@@ -17,6 +17,16 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
+typedef struct mywallet_tmp
+{
+  u32 ipad[5];
+  u32 opad[5];
+
+  u32 dgst[10];
+  u32 out[10];
+
+} mywallet_tmp_t;
+
 DECLSPEC void hmac_sha1_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest)
 {
   digest[0] = ipad[0];

OpenCL/m12800-pure.cl

@@ -17,6 +17,22 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
+typedef struct pbkdf2_sha256_tmp
+{
+  u32  ipad[8];
+  u32  opad[8];
+
+  u32  dgst[32];
+  u32  out[32];
+
+} pbkdf2_sha256_tmp_t;
+
+typedef struct pbkdf2_sha256
+{
+  u32 salt_buf[64];
+
+} pbkdf2_sha256_t;
+
 DECLSPEC void hmac_sha256_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest)
 {
   digest[0] = ipad[0];

OpenCL/m12900-pure.cl

@@ -16,6 +16,22 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
+typedef struct pbkdf2_sha256_tmp
+{
+  u32  ipad[8];
+  u32  opad[8];
+
+  u32  dgst[32];
+  u32  out[32];
+
+} pbkdf2_sha256_tmp_t;
+
+typedef struct pbkdf2_sha256
+{
+  u32 salt_buf[64];
+
+} pbkdf2_sha256_t;
+
 DECLSPEC void hmac_sha256_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest)
 {
   digest[0] = ipad[0];

OpenCL/m13000-pure.cl

@@ -16,6 +16,16 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
+typedef struct pbkdf2_sha256_tmp
+{
+  u32  ipad[8];
+  u32  opad[8];
+
+  u32  dgst[32];
+  u32  out[32];
+
+} pbkdf2_sha256_tmp_t;
+
 DECLSPEC void hmac_sha256_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest)
 {
   digest[0] = ipad[0];
@@ -58,7 +68,7 @@ DECLSPEC void hmac_sha256_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *i
   sha256_transform_vector (w0, w1, w2, w3, digest);
 }
 
-__kernel void m13000_init (KERN_ATTR_TMPS_ESALT (pbkdf2_sha256_tmp_t, pbkdf2_sha256_t))
+__kernel void m13000_init (KERN_ATTR_TMPS (pbkdf2_sha256_tmp_t))
 {
   /**
    * base
@@ -142,7 +152,7 @@ __kernel void m13000_init (KERN_ATTR_TMPS_ESALT (pbkdf2_sha256_tmp_t, pbkdf2_sha
   }
 }
 
-__kernel void m13000_loop (KERN_ATTR_TMPS_ESALT (pbkdf2_sha256_tmp_t, pbkdf2_sha256_t))
+__kernel void m13000_loop (KERN_ATTR_TMPS (pbkdf2_sha256_tmp_t))
 {
   const u64 gid = get_global_id (0);
 
@@ -248,7 +258,7 @@ __kernel void m13000_loop (KERN_ATTR_TMPS_ESALT (pbkdf2_sha256_tmp_t, pbkdf2_sha
   }
 }
 
-__kernel void m13000_comp (KERN_ATTR_TMPS_ESALT (pbkdf2_sha256_tmp_t, pbkdf2_sha256_t))
+__kernel void m13000_comp (KERN_ATTR_TMPS (pbkdf2_sha256_tmp_t))
 {
   /**
    * base

OpenCL/m13100_a0-optimized.cl

@@ -17,6 +17,15 @@
 #include "inc_hash_md4.cl"
 #include "inc_hash_md5.cl"
 
+typedef struct krb5tgs
+{
+  u32 account_info[512];
+  u32 checksum[4];
+  u32 edata2[5120];
+  u32 edata2_len;
+
+} krb5tgs_t;
+
 typedef struct
 {
   u8 S[256];

OpenCL/m13100_a0-pure.cl

@@ -16,6 +16,15 @@
 #include "inc_hash_md4.cl"
 #include "inc_hash_md5.cl"
 
+typedef struct krb5tgs
+{
+  u32 account_info[512];
+  u32 checksum[4];
+  u32 edata2[5120];
+  u32 edata2_len;
+
+} krb5tgs_t;
+
 typedef struct
 {
   u8 S[256];

OpenCL/m13100_a1-optimized.cl

@@ -15,6 +15,15 @@
 #include "inc_hash_md4.cl"
 #include "inc_hash_md5.cl"
 
+typedef struct krb5tgs
+{
+  u32 account_info[512];
+  u32 checksum[4];
+  u32 edata2[5120];
+  u32 edata2_len;
+
+} krb5tgs_t;
+
 typedef struct
 {
   u8 S[256];

OpenCL/m13100_a1-pure.cl

@@ -14,6 +14,15 @@
 #include "inc_hash_md4.cl"
 #include "inc_hash_md5.cl"
 
+typedef struct krb5tgs
+{
+  u32 account_info[512];
+  u32 checksum[4];
+  u32 edata2[5120];
+  u32 edata2_len;
+
+} krb5tgs_t;
+
 typedef struct
 {
   u8 S[256];

OpenCL/m13100_a3-optimized.cl

@@ -15,6 +15,15 @@
 #include "inc_hash_md4.cl"
 #include "inc_hash_md5.cl"
 
+typedef struct krb5tgs
+{
+  u32 account_info[512];
+  u32 checksum[4];
+  u32 edata2[5120];
+  u32 edata2_len;
+
+} krb5tgs_t;
+
 typedef struct
 {
   u8 S[256];

OpenCL/m13100_a3-pure.cl

@@ -14,6 +14,15 @@
 #include "inc_hash_md4.cl"
 #include "inc_hash_md5.cl"
 
+typedef struct krb5tgs
+{
+  u32 account_info[512];
+  u32 checksum[4];
+  u32 edata2[5120];
+  u32 edata2_len;
+
+} krb5tgs_t;
+
 typedef struct
 {
   u8 S[256];

OpenCL/m13200-pure.cl

@@ -11,6 +11,14 @@
 #include "inc_hash_sha1.cl"
 #include "inc_cipher_aes.cl"
 
+typedef struct axcrypt_tmp
+{
+  u32 KEK[4];
+  u32 lsb[4];
+  u32 cipher[4];
+
+} axcrypt_tmp_t;
+
 __kernel void m13200_init (KERN_ATTR_TMPS (axcrypt_tmp_t))
 {
   /**

OpenCL/m13400-pure.cl

@@ -15,6 +15,35 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
+typedef struct keepass_tmp
+{
+  u32 tmp_digest[8];
+
+} keepass_tmp_t;
+
+typedef struct keepass
+{
+  u32 version;
+  u32 algorithm;
+
+  /* key-file handling */
+  u32 keyfile_len;
+  u32 keyfile[8];
+
+  u32 final_random_seed[8];
+  u32 transf_random_seed[8];
+  u32 enc_iv[4];
+  u32 contents_hash[8];
+
+  /* specific to version 1 */
+  u32 contents_len;
+  u32 contents[75000];
+
+  /* specific to version 2 */
+  u32 expected_bytes[8];
+
+} keepass_t;
+
 __kernel void m13400_init (KERN_ATTR_TMPS_ESALT (keepass_tmp_t, keepass_t))
 {
   /**

OpenCL/m13500_a0-optimized.cl

@@ -15,6 +15,16 @@
 #include "inc_rp_optimized.cl"
 #include "inc_simd.cl"
 
+typedef struct pstoken
+{
+  u32 salt_buf[128];
+  u32 salt_len;
+
+  u32 pc_digest[5];
+  u32 pc_offset;
+
+} pstoken_t;
+
 __kernel void m13500_m04 (KERN_ATTR_RULES_ESALT (pstoken_t))
 {
   /**

OpenCL/m13500_a0-pure.cl

@@ -15,6 +15,16 @@
 #include "inc_scalar.cl"
 #include "inc_hash_sha1.cl"
 
+typedef struct pstoken
+{
+  u32 salt_buf[128];
+  u32 salt_len;
+
+  u32 pc_digest[5];
+  u32 pc_offset;
+
+} pstoken_t;
+
 __kernel void m13500_mxx (KERN_ATTR_RULES_ESALT (pstoken_t))
 {
   /**

OpenCL/m13500_a1-optimized.cl

@@ -13,6 +13,16 @@
 #include "inc_common.cl"
 #include "inc_simd.cl"
 
+typedef struct pstoken
+{
+  u32 salt_buf[128];
+  u32 salt_len;
+
+  u32 pc_digest[5];
+  u32 pc_offset;
+
+} pstoken_t;
+
 __kernel void m13500_m04 (KERN_ATTR_ESALT (pstoken_t))
 {
   /**

OpenCL/m13500_a1-pure.cl

@@ -13,6 +13,16 @@
 #include "inc_scalar.cl"
 #include "inc_hash_sha1.cl"
 
+typedef struct pstoken
+{
+  u32 salt_buf[128];
+  u32 salt_len;
+
+  u32 pc_digest[5];
+  u32 pc_offset;
+
+} pstoken_t;
+
 __kernel void m13500_mxx (KERN_ATTR_ESALT (pstoken_t))
 {
   /**

OpenCL/m13500_a3-optimized.cl

@@ -13,6 +13,16 @@
 #include "inc_common.cl"
 #include "inc_simd.cl"
 
+typedef struct pstoken
+{
+  u32 salt_buf[128];
+  u32 salt_len;
+
+  u32 pc_digest[5];
+  u32 pc_offset;
+
+} pstoken_t;
+
 DECLSPEC void m13500m (u32 *w0, u32 *w1, u32 *w2, u32 *w3, const u32 pw_len, KERN_ATTR_ESALT (pstoken_t))
 {
   /**

OpenCL/m13500_a3-pure.cl

@@ -13,6 +13,16 @@
 #include "inc_simd.cl"
 #include "inc_hash_sha1.cl"
 
+typedef struct pstoken
+{
+  u32 salt_buf[128];
+  u32 salt_len;
+
+  u32 pc_digest[5];
+  u32 pc_offset;
+
+} pstoken_t;
+
 __kernel void m13500_mxx (KERN_ATTR_VECTOR_ESALT (pstoken_t))
 {
   /**

OpenCL/m13600-pure.cl

@@ -16,6 +16,32 @@
 #define COMPARE_S "inc_comp_single.cl"
 #define COMPARE_M "inc_comp_multi.cl"
 
+typedef struct pbkdf2_sha1_tmp
+{
+  u32  ipad[5];
+  u32  opad[5];
+
+  u32  dgst[32];
+  u32  out[32];
+
+} pbkdf2_sha1_tmp_t;
+
+typedef struct zip2
+{
+  u32 type;
+  u32 mode;
+  u32 magic;
+  u32 salt_len;
+  u32 salt_buf[4];
+  u32 verify_bytes;
+  u32 compress_length;
+  u32 data_len;
+  u32 data_buf[2048];
+  u32 auth_len;
+  u32 auth_buf[4];
+
+} zip2_t;
+
 DECLSPEC void hmac_sha1_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest)
 {
   digest[0] = ipad[0];

OpenCL/m13751-pure.cl

@@ -24,6 +24,28 @@
 #include "inc_truecrypt_xts.cl"
 #include "inc_veracrypt_xts.cl"
 
+typedef struct tc
+{
+  u32 salt_buf[32];
+  u32 data_buf[112];
+  u32 keyfile_buf[16];
+  u32 signature;
+
+  keyboard_layout_mapping_t keyboard_layout_mapping_buf[256];
+  int                       keyboard_layout_mapping_cnt;
+
+} tc_t;
+
+typedef struct tc_tmp
+{
+  u32 ipad[16];
+  u32 opad[16];
+
+  u32 dgst[64];
+  u32 out[64];
+
+} tc_tmp_t;
+
 DECLSPEC void hmac_sha256_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest)
 {
   digest[0] = ipad[0];

OpenCL/m13752-pure.cl

@@ -24,6 +24,28 @@
 #include "inc_truecrypt_xts.cl"
 #include "inc_veracrypt_xts.cl"
 
+typedef struct tc
+{
+  u32 salt_buf[32];
+  u32 data_buf[112];
+  u32 keyfile_buf[16];
+  u32 signature;
+
+  keyboard_layout_mapping_t keyboard_layout_mapping_buf[256];
+  int                       keyboard_layout_mapping_cnt;
+
+} tc_t;
+
+typedef struct tc_tmp
+{
+  u32 ipad[16];
+  u32 opad[16];
+
+  u32 dgst[64];
+  u32 out[64];
+
+} tc_tmp_t;
+
 DECLSPEC void hmac_sha256_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest)
 {
   digest[0] = ipad[0];

OpenCL/m13753-pure.cl

@@ -24,6 +24,28 @@
 #include "inc_truecrypt_xts.cl"
 #include "inc_veracrypt_xts.cl"
 
+typedef struct tc
+{
+  u32 salt_buf[32];
+  u32 data_buf[112];
+  u32 keyfile_buf[16];
+  u32 signature;
+
+  keyboard_layout_mapping_t keyboard_layout_mapping_buf[256];
+  int                       keyboard_layout_mapping_cnt;
+
+} tc_t;
+
+typedef struct tc_tmp
+{
+  u32 ipad[16];
+  u32 opad[16];
+
+  u32 dgst[64];
+  u32 out[64];
+
+} tc_tmp_t;
+
 DECLSPEC void hmac_sha256_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest)
 {
   digest[0] = ipad[0];

OpenCL/m13771-pure.cl

@@ -24,6 +24,31 @@
 #include "inc_truecrypt_xts.cl"
 #include "inc_veracrypt_xts.cl"
 
+typedef struct tc
+{
+  u32 salt_buf[32];
+  u32 data_buf[112];
+  u32 keyfile_buf[16];
+  u32 signature;
+
+  keyboard_layout_mapping_t keyboard_layout_mapping_buf[256];
+  int                       keyboard_layout_mapping_cnt;
+
+} tc_t;
+
+typedef struct vc64_sbog_tmp
+{
+  u64  ipad_raw[8];
+  u64  opad_raw[8];
+
+  u64  ipad_hash[8];
+  u64  opad_hash[8];
+
+  u64  dgst[32];
+  u64  out[32];
+
+} vc64_sbog_tmp_t;
+
 DECLSPEC void hmac_streebog512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u64x *ipad_hash, u64x *opad_hash, u64x *ipad_raw, u64x *opad_raw, u64x *digest, SHM_TYPE u64a (*s_sbob_sl64)[256])
 {
   const u64x nullbuf[8] = { 0 };

OpenCL/m13772-pure.cl

@@ -24,6 +24,31 @@
 #include "inc_truecrypt_xts.cl"
 #include "inc_veracrypt_xts.cl"
 
+typedef struct tc
+{
+  u32 salt_buf[32];
+  u32 data_buf[112];
+  u32 keyfile_buf[16];
+  u32 signature;
+
+  keyboard_layout_mapping_t keyboard_layout_mapping_buf[256];
+  int                       keyboard_layout_mapping_cnt;
+
+} tc_t;
+
+typedef struct vc64_sbog_tmp
+{
+  u64  ipad_raw[8];
+  u64  opad_raw[8];
+
+  u64  ipad_hash[8];
+  u64  opad_hash[8];
+
+  u64  dgst[32];
+  u64  out[32];
+
+} vc64_sbog_tmp_t;
+
 DECLSPEC void hmac_streebog512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u64x *ipad_hash, u64x *opad_hash, u64x *ipad_raw, u64x *opad_raw, u64x *digest, SHM_TYPE u64a (*s_sbob_sl64)[256])
 {
   const u64x nullbuf[8] = { 0 };

OpenCL/m13773-pure.cl

@@ -24,6 +24,31 @@
 #include "inc_truecrypt_xts.cl"
 #include "inc_veracrypt_xts.cl"
 
+typedef struct tc
+{
+  u32 salt_buf[32];
+  u32 data_buf[112];
+  u32 keyfile_buf[16];
+  u32 signature;
+
+  keyboard_layout_mapping_t keyboard_layout_mapping_buf[256];
+  int                       keyboard_layout_mapping_cnt;
+
+} tc_t;
+
+typedef struct vc64_sbog_tmp
+{
+  u64  ipad_raw[8];
+  u64  opad_raw[8];
+
+  u64  ipad_hash[8];
+  u64  opad_hash[8];
+
+  u64  dgst[32];
+  u64  out[32];
+
+} vc64_sbog_tmp_t;
+
 DECLSPEC void hmac_streebog512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u64x *ipad_hash, u64x *opad_hash, u64x *ipad_raw, u64x *opad_raw, u64x *digest, SHM_TYPE u64a (*s_sbob_sl64)[256])
 {
   const u64x nullbuf[8] = { 0 };

OpenCL/m13800_a0-optimized.cl

@@ -16,6 +16,12 @@
 #include "inc_simd.cl"
 #include "inc_hash_sha256.cl"
 
+typedef struct win8phone
+{
+  u32 salt_buf[32];
+
+} win8phone_t;
+
 DECLSPEC void sha256_transform_transport_vector (const u32x *w, u32x *digest)
 {
   sha256_transform_vector (w + 0, w + 4, w + 8, w + 12, digest);

OpenCL/m13800_a0-pure.cl

@@ -15,6 +15,12 @@
 #include "inc_scalar.cl"
 #include "inc_hash_sha256.cl"
 
+typedef struct win8phone
+{
+  u32 salt_buf[32];
+
+} win8phone_t;
+
 __kernel void m13800_mxx (KERN_ATTR_RULES_ESALT (win8phone_t))
 {
   /**

OpenCL/m13800_a1-optimized.cl

@@ -14,6 +14,12 @@
 #include "inc_simd.cl"
 #include "inc_hash_sha256.cl"
 
+typedef struct win8phone
+{
+  u32 salt_buf[32];
+
+} win8phone_t;
+
 DECLSPEC void sha256_transform_transport_vector (const u32x *w, u32x *digest)
 {
   sha256_transform_vector (w + 0, w + 4, w + 8, w + 12, digest);

OpenCL/m13800_a1-pure.cl

@@ -13,6 +13,12 @@
 #include "inc_scalar.cl"
 #include "inc_hash_sha256.cl"
 
+typedef struct win8phone
+{
+  u32 salt_buf[32];
+
+} win8phone_t;
+
 __kernel void m13800_mxx (KERN_ATTR_ESALT (win8phone_t))
 {
   /**

OpenCL/m13800_a3-optimized.cl

@@ -13,6 +13,12 @@
 #include "inc_simd.cl"
 #include "inc_hash_sha256.cl"
 
+typedef struct win8phone
+{
+  u32 salt_buf[32];
+
+} win8phone_t;
+
 DECLSPEC void sha256_transform_transport_vector (const u32x *w, u32x *digest)
 {
   sha256_transform_vector (w + 0, w + 4, w + 8, w + 12, digest);

OpenCL/m13800_a3-pure.cl

@@ -13,6 +13,12 @@
 #include "inc_simd.cl"
 #include "inc_hash_sha256.cl"
 
+typedef struct win8phone
+{
+  u32 salt_buf[32];
+
+} win8phone_t;
+
 __kernel void m13800_mxx (KERN_ATTR_VECTOR_ESALT (win8phone_t))
 {
   /**