arno/hashcat
1
0
Fork 0
mirror of https://github.com/hashcat/hashcat.git synced 2025-02-16 17:42:04 +00:00
Code Issues Releases Wiki Activity

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

Browse Source
This commit is contained in:
Jens Steube 2019-03-08 10:18:20 +01:00
parent b585f25c8c
commit 4bce25dd9d
54 changed files with 682 additions and 951 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

958
OpenCL/inc_types.cl
View File

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

6
OpenCL/m00400-optimized.cl
View File

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

6
OpenCL/m00400-pure.cl
View File

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

6
OpenCL/m00500-optimized.cl
View File

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

6
OpenCL/m00500-pure.cl
View File

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

10
OpenCL/m00600_a0-optimized.cl
View File

@ -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

10
OpenCL/m00600_a1-optimized.cl
View File

@ -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

10
OpenCL/m00600_a3-optimized.cl
View File

@ -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

6
OpenCL/m01600-optimized.cl
View File

@ -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

6
OpenCL/m01600-pure.cl
View File

@ -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

14
OpenCL/m01800-optimized.cl
View File

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

14
OpenCL/m01800-pure.cl
View File

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

10
OpenCL/m02100-pure.cl
View File

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

33
OpenCL/m02500-pure.cl
View File

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

33
OpenCL/m02501-pure.cl
View File

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

13
OpenCL/m03200-pure.cl
View File

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

6
OpenCL/m05200-pure.cl
View File

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

10
OpenCL/m05300_a0-optimized.cl
View File

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

10
OpenCL/m05300_a0-pure.cl
View File

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

10
OpenCL/m05300_a1-optimized.cl
View File

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

10
OpenCL/m05300_a1-pure.cl
View File

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

10
OpenCL/m05300_a3-optimized.cl
View File

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

10
OpenCL/m05300_a3-pure.cl
View File

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

10
OpenCL/m05400_a0-optimized.cl
View File

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

10
OpenCL/m05400_a0-pure.cl
View File

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

10
OpenCL/m05400_a1-optimized.cl
View File

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

10
OpenCL/m05400_a1-pure.cl
View File

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

10
OpenCL/m05400_a3-optimized.cl
View File

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

10
OpenCL/m05400_a3-pure.cl
View File

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

12
OpenCL/m05500_a0-optimized.cl
View File

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

12
OpenCL/m05500_a0-pure.cl
View File

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

12
OpenCL/m05500_a1-optimized.cl
View File

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

12
OpenCL/m05500_a1-pure.cl
View File

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

12
OpenCL/m05500_a3-optimized.cl
View File

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

12
OpenCL/m05500_a3-pure.cl
View File

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

12
OpenCL/m05600_a0-optimized.cl
View File

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

12
OpenCL/m05600_a0-pure.cl
View File

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

12
OpenCL/m05600_a1-optimized.cl
View File

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

12
OpenCL/m05600_a1-pure.cl
View File

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

12
OpenCL/m05600_a3-optimized.cl
View File

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

12
OpenCL/m05600_a3-pure.cl
View File

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

6
OpenCL/m05800-optimized.cl
View File

@ -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,

6
OpenCL/m05800-pure.cl
View File

@ -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,

22
OpenCL/m06211-pure.cl
View File

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

22
OpenCL/m06212-pure.cl
View File

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

22
OpenCL/m06213-pure.cl
View File

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

22
OpenCL/m06221-pure.cl
View File

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

22
OpenCL/m06222-pure.cl
View File

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

22
OpenCL/m06223-pure.cl
View File

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

22
OpenCL/m06231-pure.cl
View File

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

22
OpenCL/m06232-pure.cl
View File

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

22
OpenCL/m06233-pure.cl
View File

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

6
OpenCL/m06300-optimized.cl
View File

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

6
OpenCL/m06300-pure.cl
View File

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