- Keymaps: Added hashcat keyboard mapping us.hckmap (can be used as template)
- Keymaps: Added hashcat keyboard mapping de.hckmap
This commit is contained in:
Jens Steube
@@ -0,0 +1,22 @@
|
||||
DECLSPEC void keyboard_map (u32 w[4], __local u32 *s_keyboard_layout)
|
||||
{
|
||||
w[0] = (s_keyboard_layout[(w[0] >> 0) & 0xff] << 0)
|
||||
| (s_keyboard_layout[(w[0] >> 8) & 0xff] << 8)
|
||||
| (s_keyboard_layout[(w[0] >> 16) & 0xff] << 16)
|
||||
| (s_keyboard_layout[(w[0] >> 24) & 0xff] << 24);
|
||||
|
||||
w[1] = (s_keyboard_layout[(w[1] >> 0) & 0xff] << 0)
|
||||
| (s_keyboard_layout[(w[1] >> 8) & 0xff] << 8)
|
||||
| (s_keyboard_layout[(w[1] >> 16) & 0xff] << 16)
|
||||
| (s_keyboard_layout[(w[1] >> 24) & 0xff] << 24);
|
||||
|
||||
w[2] = (s_keyboard_layout[(w[2] >> 0) & 0xff] << 0)
|
||||
| (s_keyboard_layout[(w[2] >> 8) & 0xff] << 8)
|
||||
| (s_keyboard_layout[(w[2] >> 16) & 0xff] << 16)
|
||||
| (s_keyboard_layout[(w[2] >> 24) & 0xff] << 24);
|
||||
|
||||
w[3] = (s_keyboard_layout[(w[3] >> 0) & 0xff] << 0)
|
||||
| (s_keyboard_layout[(w[3] >> 8) & 0xff] << 8)
|
||||
| (s_keyboard_layout[(w[3] >> 16) & 0xff] << 16)
|
||||
| (s_keyboard_layout[(w[3] >> 24) & 0xff] << 24);
|
||||
}
|
||||
@@ -0,0 +1,24 @@
|
||||
DECLSPEC u32 u8add (const u32 a, const u32 b)
|
||||
{
|
||||
const u32 a1 = (a >> 0) & 0xff;
|
||||
const u32 a2 = (a >> 8) & 0xff;
|
||||
const u32 a3 = (a >> 16) & 0xff;
|
||||
const u32 a4 = (a >> 24) & 0xff;
|
||||
|
||||
const u32 b1 = (b >> 0) & 0xff;
|
||||
const u32 b2 = (b >> 8) & 0xff;
|
||||
const u32 b3 = (b >> 16) & 0xff;
|
||||
const u32 b4 = (b >> 24) & 0xff;
|
||||
|
||||
const u32 r1 = (a1 + b1) & 0xff;
|
||||
const u32 r2 = (a2 + b2) & 0xff;
|
||||
const u32 r3 = (a3 + b3) & 0xff;
|
||||
const u32 r4 = (a4 + b4) & 0xff;
|
||||
|
||||
const u32 r = r1 << 0
|
||||
| r2 << 8
|
||||
| r3 << 16
|
||||
| r4 << 24;
|
||||
|
||||
return r;
|
||||
}
|
||||
+2
-48
@@ -17,57 +17,11 @@
|
||||
#include "inc_cipher_twofish.cl"
|
||||
#include "inc_cipher_serpent.cl"
|
||||
|
||||
#include "inc_truecrypt_keyfile.cl"
|
||||
#include "inc_truecrypt_keyboard.cl"
|
||||
#include "inc_truecrypt_crc32.cl"
|
||||
#include "inc_truecrypt_xts.cl"
|
||||
|
||||
DECLSPEC void keyboard_map (u32 w[4], __local u32 *s_keyboard_layout)
|
||||
{
|
||||
w[0] = (s_keyboard_layout[(w[0] >> 0) & 0xff] << 0)
|
||||
| (s_keyboard_layout[(w[0] >> 8) & 0xff] << 8)
|
||||
| (s_keyboard_layout[(w[0] >> 16) & 0xff] << 16)
|
||||
| (s_keyboard_layout[(w[0] >> 24) & 0xff] << 24);
|
||||
|
||||
w[1] = (s_keyboard_layout[(w[1] >> 0) & 0xff] << 0)
|
||||
| (s_keyboard_layout[(w[1] >> 8) & 0xff] << 8)
|
||||
| (s_keyboard_layout[(w[1] >> 16) & 0xff] << 16)
|
||||
| (s_keyboard_layout[(w[1] >> 24) & 0xff] << 24);
|
||||
|
||||
w[2] = (s_keyboard_layout[(w[2] >> 0) & 0xff] << 0)
|
||||
| (s_keyboard_layout[(w[2] >> 8) & 0xff] << 8)
|
||||
| (s_keyboard_layout[(w[2] >> 16) & 0xff] << 16)
|
||||
| (s_keyboard_layout[(w[2] >> 24) & 0xff] << 24);
|
||||
|
||||
w[3] = (s_keyboard_layout[(w[3] >> 0) & 0xff] << 0)
|
||||
| (s_keyboard_layout[(w[3] >> 8) & 0xff] << 8)
|
||||
| (s_keyboard_layout[(w[3] >> 16) & 0xff] << 16)
|
||||
| (s_keyboard_layout[(w[3] >> 24) & 0xff] << 24);
|
||||
}
|
||||
|
||||
DECLSPEC u32 u8add (const u32 a, const u32 b)
|
||||
{
|
||||
const u32 a1 = (a >> 0) & 0xff;
|
||||
const u32 a2 = (a >> 8) & 0xff;
|
||||
const u32 a3 = (a >> 16) & 0xff;
|
||||
const u32 a4 = (a >> 24) & 0xff;
|
||||
|
||||
const u32 b1 = (b >> 0) & 0xff;
|
||||
const u32 b2 = (b >> 8) & 0xff;
|
||||
const u32 b3 = (b >> 16) & 0xff;
|
||||
const u32 b4 = (b >> 24) & 0xff;
|
||||
|
||||
const u32 r1 = (a1 + b1) & 0xff;
|
||||
const u32 r2 = (a2 + b2) & 0xff;
|
||||
const u32 r3 = (a3 + b3) & 0xff;
|
||||
const u32 r4 = (a4 + b4) & 0xff;
|
||||
|
||||
const u32 r = r1 << 0
|
||||
| r2 << 8
|
||||
| r3 << 16
|
||||
| r4 << 24;
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
DECLSPEC void hmac_ripemd160_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest)
|
||||
{
|
||||
digest[0] = ipad[0];
|
||||
|
||||
+2
-48
@@ -17,57 +17,11 @@
|
||||
#include "inc_cipher_twofish.cl"
|
||||
#include "inc_cipher_serpent.cl"
|
||||
|
||||
#include "inc_truecrypt_keyfile.cl"
|
||||
#include "inc_truecrypt_keyboard.cl"
|
||||
#include "inc_truecrypt_crc32.cl"
|
||||
#include "inc_truecrypt_xts.cl"
|
||||
|
||||
DECLSPEC void keyboard_map (u32 w[4], __local u32 *s_keyboard_layout)
|
||||
{
|
||||
w[0] = (s_keyboard_layout[(w[0] >> 0) & 0xff] << 0)
|
||||
| (s_keyboard_layout[(w[0] >> 8) & 0xff] << 8)
|
||||
| (s_keyboard_layout[(w[0] >> 16) & 0xff] << 16)
|
||||
| (s_keyboard_layout[(w[0] >> 24) & 0xff] << 24);
|
||||
|
||||
w[1] = (s_keyboard_layout[(w[1] >> 0) & 0xff] << 0)
|
||||
| (s_keyboard_layout[(w[1] >> 8) & 0xff] << 8)
|
||||
| (s_keyboard_layout[(w[1] >> 16) & 0xff] << 16)
|
||||
| (s_keyboard_layout[(w[1] >> 24) & 0xff] << 24);
|
||||
|
||||
w[2] = (s_keyboard_layout[(w[2] >> 0) & 0xff] << 0)
|
||||
| (s_keyboard_layout[(w[2] >> 8) & 0xff] << 8)
|
||||
| (s_keyboard_layout[(w[2] >> 16) & 0xff] << 16)
|
||||
| (s_keyboard_layout[(w[2] >> 24) & 0xff] << 24);
|
||||
|
||||
w[3] = (s_keyboard_layout[(w[3] >> 0) & 0xff] << 0)
|
||||
| (s_keyboard_layout[(w[3] >> 8) & 0xff] << 8)
|
||||
| (s_keyboard_layout[(w[3] >> 16) & 0xff] << 16)
|
||||
| (s_keyboard_layout[(w[3] >> 24) & 0xff] << 24);
|
||||
}
|
||||
|
||||
DECLSPEC u32 u8add (const u32 a, const u32 b)
|
||||
{
|
||||
const u32 a1 = (a >> 0) & 0xff;
|
||||
const u32 a2 = (a >> 8) & 0xff;
|
||||
const u32 a3 = (a >> 16) & 0xff;
|
||||
const u32 a4 = (a >> 24) & 0xff;
|
||||
|
||||
const u32 b1 = (b >> 0) & 0xff;
|
||||
const u32 b2 = (b >> 8) & 0xff;
|
||||
const u32 b3 = (b >> 16) & 0xff;
|
||||
const u32 b4 = (b >> 24) & 0xff;
|
||||
|
||||
const u32 r1 = (a1 + b1) & 0xff;
|
||||
const u32 r2 = (a2 + b2) & 0xff;
|
||||
const u32 r3 = (a3 + b3) & 0xff;
|
||||
const u32 r4 = (a4 + b4) & 0xff;
|
||||
|
||||
const u32 r = r1 << 0
|
||||
| r2 << 8
|
||||
| r3 << 16
|
||||
| r4 << 24;
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
DECLSPEC void hmac_ripemd160_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest)
|
||||
{
|
||||
digest[0] = ipad[0];
|
||||
|
||||
+2
-48
@@ -17,57 +17,11 @@
|
||||
#include "inc_cipher_twofish.cl"
|
||||
#include "inc_cipher_serpent.cl"
|
||||
|
||||
#include "inc_truecrypt_keyfile.cl"
|
||||
#include "inc_truecrypt_keyboard.cl"
|
||||
#include "inc_truecrypt_crc32.cl"
|
||||
#include "inc_truecrypt_xts.cl"
|
||||
|
||||
DECLSPEC void keyboard_map (u32 w[4], __local u32 *s_keyboard_layout)
|
||||
{
|
||||
w[0] = (s_keyboard_layout[(w[0] >> 0) & 0xff] << 0)
|
||||
| (s_keyboard_layout[(w[0] >> 8) & 0xff] << 8)
|
||||
| (s_keyboard_layout[(w[0] >> 16) & 0xff] << 16)
|
||||
| (s_keyboard_layout[(w[0] >> 24) & 0xff] << 24);
|
||||
|
||||
w[1] = (s_keyboard_layout[(w[1] >> 0) & 0xff] << 0)
|
||||
| (s_keyboard_layout[(w[1] >> 8) & 0xff] << 8)
|
||||
| (s_keyboard_layout[(w[1] >> 16) & 0xff] << 16)
|
||||
| (s_keyboard_layout[(w[1] >> 24) & 0xff] << 24);
|
||||
|
||||
w[2] = (s_keyboard_layout[(w[2] >> 0) & 0xff] << 0)
|
||||
| (s_keyboard_layout[(w[2] >> 8) & 0xff] << 8)
|
||||
| (s_keyboard_layout[(w[2] >> 16) & 0xff] << 16)
|
||||
| (s_keyboard_layout[(w[2] >> 24) & 0xff] << 24);
|
||||
|
||||
w[3] = (s_keyboard_layout[(w[3] >> 0) & 0xff] << 0)
|
||||
| (s_keyboard_layout[(w[3] >> 8) & 0xff] << 8)
|
||||
| (s_keyboard_layout[(w[3] >> 16) & 0xff] << 16)
|
||||
| (s_keyboard_layout[(w[3] >> 24) & 0xff] << 24);
|
||||
}
|
||||
|
||||
DECLSPEC u32 u8add (const u32 a, const u32 b)
|
||||
{
|
||||
const u32 a1 = (a >> 0) & 0xff;
|
||||
const u32 a2 = (a >> 8) & 0xff;
|
||||
const u32 a3 = (a >> 16) & 0xff;
|
||||
const u32 a4 = (a >> 24) & 0xff;
|
||||
|
||||
const u32 b1 = (b >> 0) & 0xff;
|
||||
const u32 b2 = (b >> 8) & 0xff;
|
||||
const u32 b3 = (b >> 16) & 0xff;
|
||||
const u32 b4 = (b >> 24) & 0xff;
|
||||
|
||||
const u32 r1 = (a1 + b1) & 0xff;
|
||||
const u32 r2 = (a2 + b2) & 0xff;
|
||||
const u32 r3 = (a3 + b3) & 0xff;
|
||||
const u32 r4 = (a4 + b4) & 0xff;
|
||||
|
||||
const u32 r = r1 << 0
|
||||
| r2 << 8
|
||||
| r3 << 16
|
||||
| r4 << 24;
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
DECLSPEC void hmac_ripemd160_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest)
|
||||
{
|
||||
digest[0] = ipad[0];
|
||||
|
||||
+30
-29
@@ -17,34 +17,11 @@
|
||||
#include "inc_cipher_twofish.cl"
|
||||
#include "inc_cipher_serpent.cl"
|
||||
|
||||
#include "inc_truecrypt_keyfile.cl"
|
||||
#include "inc_truecrypt_keyboard.cl"
|
||||
#include "inc_truecrypt_crc32.cl"
|
||||
#include "inc_truecrypt_xts.cl"
|
||||
|
||||
DECLSPEC u32 u8add (const u32 a, const u32 b)
|
||||
{
|
||||
const u32 a1 = (a >> 0) & 0xff;
|
||||
const u32 a2 = (a >> 8) & 0xff;
|
||||
const u32 a3 = (a >> 16) & 0xff;
|
||||
const u32 a4 = (a >> 24) & 0xff;
|
||||
|
||||
const u32 b1 = (b >> 0) & 0xff;
|
||||
const u32 b2 = (b >> 8) & 0xff;
|
||||
const u32 b3 = (b >> 16) & 0xff;
|
||||
const u32 b4 = (b >> 24) & 0xff;
|
||||
|
||||
const u32 r1 = (a1 + b1) & 0xff;
|
||||
const u32 r2 = (a2 + b2) & 0xff;
|
||||
const u32 r3 = (a3 + b3) & 0xff;
|
||||
const u32 r4 = (a4 + b4) & 0xff;
|
||||
|
||||
const u32 r = r1 << 0
|
||||
| r2 << 8
|
||||
| r3 << 16
|
||||
| r4 << 24;
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
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];
|
||||
@@ -105,14 +82,29 @@ DECLSPEC void hmac_sha512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *w
|
||||
|
||||
__kernel void m06221_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global tc64_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const tc_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
|
||||
{
|
||||
const u64 gid = get_global_id (0);
|
||||
const u64 lid = get_local_id (0);
|
||||
const u64 lsz = get_local_size (0);
|
||||
|
||||
/**
|
||||
* keyboard layout shared
|
||||
*/
|
||||
|
||||
__local u32 s_keyboard_layout[256];
|
||||
|
||||
for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
|
||||
{
|
||||
s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
/**
|
||||
* base
|
||||
*/
|
||||
|
||||
const u64 gid = get_global_id (0);
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
u32 w0[4];
|
||||
u32 w1[4];
|
||||
u32 w2[4];
|
||||
@@ -155,6 +147,15 @@ __kernel void m06221_init (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
w7[2] = pws[gid].i[30];
|
||||
w7[3] = pws[gid].i[31];
|
||||
|
||||
keyboard_map (w0, s_keyboard_layout);
|
||||
keyboard_map (w1, s_keyboard_layout);
|
||||
keyboard_map (w2, s_keyboard_layout);
|
||||
keyboard_map (w3, s_keyboard_layout);
|
||||
keyboard_map (w4, s_keyboard_layout);
|
||||
keyboard_map (w5, s_keyboard_layout);
|
||||
keyboard_map (w6, s_keyboard_layout);
|
||||
keyboard_map (w7, s_keyboard_layout);
|
||||
|
||||
w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]);
|
||||
w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]);
|
||||
w0[2] = u8add (w0[2], esalt_bufs[digests_offset].keyfile_buf[ 2]);
|
||||
|
||||
+30
-29
@@ -17,34 +17,11 @@
|
||||
#include "inc_cipher_twofish.cl"
|
||||
#include "inc_cipher_serpent.cl"
|
||||
|
||||
#include "inc_truecrypt_keyfile.cl"
|
||||
#include "inc_truecrypt_keyboard.cl"
|
||||
#include "inc_truecrypt_crc32.cl"
|
||||
#include "inc_truecrypt_xts.cl"
|
||||
|
||||
DECLSPEC u32 u8add (const u32 a, const u32 b)
|
||||
{
|
||||
const u32 a1 = (a >> 0) & 0xff;
|
||||
const u32 a2 = (a >> 8) & 0xff;
|
||||
const u32 a3 = (a >> 16) & 0xff;
|
||||
const u32 a4 = (a >> 24) & 0xff;
|
||||
|
||||
const u32 b1 = (b >> 0) & 0xff;
|
||||
const u32 b2 = (b >> 8) & 0xff;
|
||||
const u32 b3 = (b >> 16) & 0xff;
|
||||
const u32 b4 = (b >> 24) & 0xff;
|
||||
|
||||
const u32 r1 = (a1 + b1) & 0xff;
|
||||
const u32 r2 = (a2 + b2) & 0xff;
|
||||
const u32 r3 = (a3 + b3) & 0xff;
|
||||
const u32 r4 = (a4 + b4) & 0xff;
|
||||
|
||||
const u32 r = r1 << 0
|
||||
| r2 << 8
|
||||
| r3 << 16
|
||||
| r4 << 24;
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
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];
|
||||
@@ -105,14 +82,29 @@ DECLSPEC void hmac_sha512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *w
|
||||
|
||||
__kernel void m06222_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global tc64_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const tc_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
|
||||
{
|
||||
const u64 gid = get_global_id (0);
|
||||
const u64 lid = get_local_id (0);
|
||||
const u64 lsz = get_local_size (0);
|
||||
|
||||
/**
|
||||
* keyboard layout shared
|
||||
*/
|
||||
|
||||
__local u32 s_keyboard_layout[256];
|
||||
|
||||
for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
|
||||
{
|
||||
s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
/**
|
||||
* base
|
||||
*/
|
||||
|
||||
const u64 gid = get_global_id (0);
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
u32 w0[4];
|
||||
u32 w1[4];
|
||||
u32 w2[4];
|
||||
@@ -155,6 +147,15 @@ __kernel void m06222_init (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
w7[2] = pws[gid].i[30];
|
||||
w7[3] = pws[gid].i[31];
|
||||
|
||||
keyboard_map (w0, s_keyboard_layout);
|
||||
keyboard_map (w1, s_keyboard_layout);
|
||||
keyboard_map (w2, s_keyboard_layout);
|
||||
keyboard_map (w3, s_keyboard_layout);
|
||||
keyboard_map (w4, s_keyboard_layout);
|
||||
keyboard_map (w5, s_keyboard_layout);
|
||||
keyboard_map (w6, s_keyboard_layout);
|
||||
keyboard_map (w7, s_keyboard_layout);
|
||||
|
||||
w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]);
|
||||
w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]);
|
||||
w0[2] = u8add (w0[2], esalt_bufs[digests_offset].keyfile_buf[ 2]);
|
||||
|
||||
+30
-29
@@ -17,34 +17,11 @@
|
||||
#include "inc_cipher_twofish.cl"
|
||||
#include "inc_cipher_serpent.cl"
|
||||
|
||||
#include "inc_truecrypt_keyfile.cl"
|
||||
#include "inc_truecrypt_keyboard.cl"
|
||||
#include "inc_truecrypt_crc32.cl"
|
||||
#include "inc_truecrypt_xts.cl"
|
||||
|
||||
DECLSPEC u32 u8add (const u32 a, const u32 b)
|
||||
{
|
||||
const u32 a1 = (a >> 0) & 0xff;
|
||||
const u32 a2 = (a >> 8) & 0xff;
|
||||
const u32 a3 = (a >> 16) & 0xff;
|
||||
const u32 a4 = (a >> 24) & 0xff;
|
||||
|
||||
const u32 b1 = (b >> 0) & 0xff;
|
||||
const u32 b2 = (b >> 8) & 0xff;
|
||||
const u32 b3 = (b >> 16) & 0xff;
|
||||
const u32 b4 = (b >> 24) & 0xff;
|
||||
|
||||
const u32 r1 = (a1 + b1) & 0xff;
|
||||
const u32 r2 = (a2 + b2) & 0xff;
|
||||
const u32 r3 = (a3 + b3) & 0xff;
|
||||
const u32 r4 = (a4 + b4) & 0xff;
|
||||
|
||||
const u32 r = r1 << 0
|
||||
| r2 << 8
|
||||
| r3 << 16
|
||||
| r4 << 24;
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
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];
|
||||
@@ -105,14 +82,29 @@ DECLSPEC void hmac_sha512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *w
|
||||
|
||||
__kernel void m06223_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global tc64_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const tc_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
|
||||
{
|
||||
const u64 gid = get_global_id (0);
|
||||
const u64 lid = get_local_id (0);
|
||||
const u64 lsz = get_local_size (0);
|
||||
|
||||
/**
|
||||
* keyboard layout shared
|
||||
*/
|
||||
|
||||
__local u32 s_keyboard_layout[256];
|
||||
|
||||
for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
|
||||
{
|
||||
s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
/**
|
||||
* base
|
||||
*/
|
||||
|
||||
const u64 gid = get_global_id (0);
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
u32 w0[4];
|
||||
u32 w1[4];
|
||||
u32 w2[4];
|
||||
@@ -155,6 +147,15 @@ __kernel void m06223_init (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
w7[2] = pws[gid].i[30];
|
||||
w7[3] = pws[gid].i[31];
|
||||
|
||||
keyboard_map (w0, s_keyboard_layout);
|
||||
keyboard_map (w1, s_keyboard_layout);
|
||||
keyboard_map (w2, s_keyboard_layout);
|
||||
keyboard_map (w3, s_keyboard_layout);
|
||||
keyboard_map (w4, s_keyboard_layout);
|
||||
keyboard_map (w5, s_keyboard_layout);
|
||||
keyboard_map (w6, s_keyboard_layout);
|
||||
keyboard_map (w7, s_keyboard_layout);
|
||||
|
||||
w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]);
|
||||
w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]);
|
||||
w0[2] = u8add (w0[2], esalt_bufs[digests_offset].keyfile_buf[ 2]);
|
||||
|
||||
+19
-30
@@ -17,34 +17,11 @@
|
||||
#include "inc_cipher_twofish.cl"
|
||||
#include "inc_cipher_serpent.cl"
|
||||
|
||||
#include "inc_truecrypt_keyfile.cl"
|
||||
#include "inc_truecrypt_keyboard.cl"
|
||||
#include "inc_truecrypt_crc32.cl"
|
||||
#include "inc_truecrypt_xts.cl"
|
||||
|
||||
DECLSPEC u32 u8add (const u32 a, const u32 b)
|
||||
{
|
||||
const u32 a1 = (a >> 0) & 0xff;
|
||||
const u32 a2 = (a >> 8) & 0xff;
|
||||
const u32 a3 = (a >> 16) & 0xff;
|
||||
const u32 a4 = (a >> 24) & 0xff;
|
||||
|
||||
const u32 b1 = (b >> 0) & 0xff;
|
||||
const u32 b2 = (b >> 8) & 0xff;
|
||||
const u32 b3 = (b >> 16) & 0xff;
|
||||
const u32 b4 = (b >> 24) & 0xff;
|
||||
|
||||
const u32 r1 = (a1 + b1) & 0xff;
|
||||
const u32 r2 = (a2 + b2) & 0xff;
|
||||
const u32 r3 = (a3 + b3) & 0xff;
|
||||
const u32 r4 = (a4 + b4) & 0xff;
|
||||
|
||||
const u32 r = r1 << 0
|
||||
| r2 << 8
|
||||
| r3 << 16
|
||||
| r4 << 24;
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
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];
|
||||
@@ -143,18 +120,21 @@ DECLSPEC void hmac_whirlpool_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x
|
||||
|
||||
__kernel void m06231_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global tc_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const tc_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
|
||||
{
|
||||
/**
|
||||
* modifier
|
||||
*/
|
||||
|
||||
const u64 gid = get_global_id (0);
|
||||
const u64 lid = get_local_id (0);
|
||||
const u64 lsz = get_local_size (0);
|
||||
|
||||
/**
|
||||
* shared
|
||||
* keyboard layout shared
|
||||
*/
|
||||
|
||||
__local u32 s_keyboard_layout[256];
|
||||
|
||||
for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
|
||||
{
|
||||
s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i];
|
||||
}
|
||||
|
||||
__local u32 s_Ch[8][256];
|
||||
__local u32 s_Cl[8][256];
|
||||
|
||||
@@ -183,6 +163,10 @@ __kernel void m06231_init (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
/**
|
||||
* base
|
||||
*/
|
||||
|
||||
u32 w0[4];
|
||||
u32 w1[4];
|
||||
u32 w2[4];
|
||||
@@ -205,6 +189,11 @@ __kernel void m06231_init (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
w3[2] = pws[gid].i[14];
|
||||
w3[3] = pws[gid].i[15];
|
||||
|
||||
keyboard_map (w0, s_keyboard_layout);
|
||||
keyboard_map (w1, s_keyboard_layout);
|
||||
keyboard_map (w2, s_keyboard_layout);
|
||||
keyboard_map (w3, s_keyboard_layout);
|
||||
|
||||
w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]);
|
||||
w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]);
|
||||
w0[2] = u8add (w0[2], esalt_bufs[digests_offset].keyfile_buf[ 2]);
|
||||
|
||||
+19
-30
@@ -17,34 +17,11 @@
|
||||
#include "inc_cipher_twofish.cl"
|
||||
#include "inc_cipher_serpent.cl"
|
||||
|
||||
#include "inc_truecrypt_keyfile.cl"
|
||||
#include "inc_truecrypt_keyboard.cl"
|
||||
#include "inc_truecrypt_crc32.cl"
|
||||
#include "inc_truecrypt_xts.cl"
|
||||
|
||||
DECLSPEC u32 u8add (const u32 a, const u32 b)
|
||||
{
|
||||
const u32 a1 = (a >> 0) & 0xff;
|
||||
const u32 a2 = (a >> 8) & 0xff;
|
||||
const u32 a3 = (a >> 16) & 0xff;
|
||||
const u32 a4 = (a >> 24) & 0xff;
|
||||
|
||||
const u32 b1 = (b >> 0) & 0xff;
|
||||
const u32 b2 = (b >> 8) & 0xff;
|
||||
const u32 b3 = (b >> 16) & 0xff;
|
||||
const u32 b4 = (b >> 24) & 0xff;
|
||||
|
||||
const u32 r1 = (a1 + b1) & 0xff;
|
||||
const u32 r2 = (a2 + b2) & 0xff;
|
||||
const u32 r3 = (a3 + b3) & 0xff;
|
||||
const u32 r4 = (a4 + b4) & 0xff;
|
||||
|
||||
const u32 r = r1 << 0
|
||||
| r2 << 8
|
||||
| r3 << 16
|
||||
| r4 << 24;
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
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];
|
||||
@@ -143,18 +120,21 @@ DECLSPEC void hmac_whirlpool_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x
|
||||
|
||||
__kernel void m06232_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global tc_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const tc_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
|
||||
{
|
||||
/**
|
||||
* modifier
|
||||
*/
|
||||
|
||||
const u64 gid = get_global_id (0);
|
||||
const u64 lid = get_local_id (0);
|
||||
const u64 lsz = get_local_size (0);
|
||||
|
||||
/**
|
||||
* shared
|
||||
* keyboard layout shared
|
||||
*/
|
||||
|
||||
__local u32 s_keyboard_layout[256];
|
||||
|
||||
for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
|
||||
{
|
||||
s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i];
|
||||
}
|
||||
|
||||
__local u32 s_Ch[8][256];
|
||||
__local u32 s_Cl[8][256];
|
||||
|
||||
@@ -183,6 +163,10 @@ __kernel void m06232_init (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
/**
|
||||
* base
|
||||
*/
|
||||
|
||||
u32 w0[4];
|
||||
u32 w1[4];
|
||||
u32 w2[4];
|
||||
@@ -205,6 +189,11 @@ __kernel void m06232_init (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
w3[2] = pws[gid].i[14];
|
||||
w3[3] = pws[gid].i[15];
|
||||
|
||||
keyboard_map (w0, s_keyboard_layout);
|
||||
keyboard_map (w1, s_keyboard_layout);
|
||||
keyboard_map (w2, s_keyboard_layout);
|
||||
keyboard_map (w3, s_keyboard_layout);
|
||||
|
||||
w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]);
|
||||
w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]);
|
||||
w0[2] = u8add (w0[2], esalt_bufs[digests_offset].keyfile_buf[ 2]);
|
||||
|
||||
+19
-30
@@ -17,34 +17,11 @@
|
||||
#include "inc_cipher_twofish.cl"
|
||||
#include "inc_cipher_serpent.cl"
|
||||
|
||||
#include "inc_truecrypt_keyfile.cl"
|
||||
#include "inc_truecrypt_keyboard.cl"
|
||||
#include "inc_truecrypt_crc32.cl"
|
||||
#include "inc_truecrypt_xts.cl"
|
||||
|
||||
DECLSPEC u32 u8add (const u32 a, const u32 b)
|
||||
{
|
||||
const u32 a1 = (a >> 0) & 0xff;
|
||||
const u32 a2 = (a >> 8) & 0xff;
|
||||
const u32 a3 = (a >> 16) & 0xff;
|
||||
const u32 a4 = (a >> 24) & 0xff;
|
||||
|
||||
const u32 b1 = (b >> 0) & 0xff;
|
||||
const u32 b2 = (b >> 8) & 0xff;
|
||||
const u32 b3 = (b >> 16) & 0xff;
|
||||
const u32 b4 = (b >> 24) & 0xff;
|
||||
|
||||
const u32 r1 = (a1 + b1) & 0xff;
|
||||
const u32 r2 = (a2 + b2) & 0xff;
|
||||
const u32 r3 = (a3 + b3) & 0xff;
|
||||
const u32 r4 = (a4 + b4) & 0xff;
|
||||
|
||||
const u32 r = r1 << 0
|
||||
| r2 << 8
|
||||
| r3 << 16
|
||||
| r4 << 24;
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
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];
|
||||
@@ -143,18 +120,21 @@ DECLSPEC void hmac_whirlpool_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x
|
||||
|
||||
__kernel void m06233_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global tc_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const tc_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
|
||||
{
|
||||
/**
|
||||
* modifier
|
||||
*/
|
||||
|
||||
const u64 gid = get_global_id (0);
|
||||
const u64 lid = get_local_id (0);
|
||||
const u64 lsz = get_local_size (0);
|
||||
|
||||
/**
|
||||
* shared
|
||||
* keyboard layout shared
|
||||
*/
|
||||
|
||||
__local u32 s_keyboard_layout[256];
|
||||
|
||||
for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
|
||||
{
|
||||
s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i];
|
||||
}
|
||||
|
||||
__local u32 s_Ch[8][256];
|
||||
__local u32 s_Cl[8][256];
|
||||
|
||||
@@ -183,6 +163,10 @@ __kernel void m06233_init (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
/**
|
||||
* base
|
||||
*/
|
||||
|
||||
u32 w0[4];
|
||||
u32 w1[4];
|
||||
u32 w2[4];
|
||||
@@ -205,6 +189,11 @@ __kernel void m06233_init (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
w3[2] = pws[gid].i[14];
|
||||
w3[3] = pws[gid].i[15];
|
||||
|
||||
keyboard_map (w0, s_keyboard_layout);
|
||||
keyboard_map (w1, s_keyboard_layout);
|
||||
keyboard_map (w2, s_keyboard_layout);
|
||||
keyboard_map (w3, s_keyboard_layout);
|
||||
|
||||
w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]);
|
||||
w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]);
|
||||
w0[2] = u8add (w0[2], esalt_bufs[digests_offset].keyfile_buf[ 2]);
|
||||
|
||||
+26
-29
@@ -17,34 +17,11 @@
|
||||
#include "inc_cipher_twofish.cl"
|
||||
#include "inc_cipher_serpent.cl"
|
||||
|
||||
#include "inc_truecrypt_keyfile.cl"
|
||||
#include "inc_truecrypt_keyboard.cl"
|
||||
#include "inc_truecrypt_crc32.cl"
|
||||
#include "inc_truecrypt_xts.cl"
|
||||
|
||||
DECLSPEC u32 u8add (const u32 a, const u32 b)
|
||||
{
|
||||
const u32 a1 = (a >> 0) & 0xff;
|
||||
const u32 a2 = (a >> 8) & 0xff;
|
||||
const u32 a3 = (a >> 16) & 0xff;
|
||||
const u32 a4 = (a >> 24) & 0xff;
|
||||
|
||||
const u32 b1 = (b >> 0) & 0xff;
|
||||
const u32 b2 = (b >> 8) & 0xff;
|
||||
const u32 b3 = (b >> 16) & 0xff;
|
||||
const u32 b4 = (b >> 24) & 0xff;
|
||||
|
||||
const u32 r1 = (a1 + b1) & 0xff;
|
||||
const u32 r2 = (a2 + b2) & 0xff;
|
||||
const u32 r3 = (a3 + b3) & 0xff;
|
||||
const u32 r4 = (a4 + b4) & 0xff;
|
||||
|
||||
const u32 r = r1 << 0
|
||||
| r2 << 8
|
||||
| r3 << 16
|
||||
| r4 << 24;
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
DECLSPEC void hmac_sha256_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest)
|
||||
{
|
||||
digest[0] = ipad[0];
|
||||
@@ -89,14 +66,29 @@ DECLSPEC void hmac_sha256_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *i
|
||||
|
||||
__kernel void m13751_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global tc_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global tc_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
|
||||
{
|
||||
const u64 gid = get_global_id (0);
|
||||
const u64 lid = get_local_id (0);
|
||||
const u64 lsz = get_local_size (0);
|
||||
|
||||
/**
|
||||
* keyboard layout shared
|
||||
*/
|
||||
|
||||
__local u32 s_keyboard_layout[256];
|
||||
|
||||
for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
|
||||
{
|
||||
s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
/**
|
||||
* base
|
||||
*/
|
||||
|
||||
const u64 gid = get_global_id (0);
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
u32 w0[4];
|
||||
u32 w1[4];
|
||||
u32 w2[4];
|
||||
@@ -119,6 +111,11 @@ __kernel void m13751_init (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
w3[2] = pws[gid].i[14];
|
||||
w3[3] = pws[gid].i[15];
|
||||
|
||||
keyboard_map (w0, s_keyboard_layout);
|
||||
keyboard_map (w1, s_keyboard_layout);
|
||||
keyboard_map (w2, s_keyboard_layout);
|
||||
keyboard_map (w3, s_keyboard_layout);
|
||||
|
||||
w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]);
|
||||
w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]);
|
||||
w0[2] = u8add (w0[2], esalt_bufs[digests_offset].keyfile_buf[ 2]);
|
||||
|
||||
+26
-29
@@ -17,34 +17,11 @@
|
||||
#include "inc_cipher_twofish.cl"
|
||||
#include "inc_cipher_serpent.cl"
|
||||
|
||||
#include "inc_truecrypt_keyfile.cl"
|
||||
#include "inc_truecrypt_keyboard.cl"
|
||||
#include "inc_truecrypt_crc32.cl"
|
||||
#include "inc_truecrypt_xts.cl"
|
||||
|
||||
DECLSPEC u32 u8add (const u32 a, const u32 b)
|
||||
{
|
||||
const u32 a1 = (a >> 0) & 0xff;
|
||||
const u32 a2 = (a >> 8) & 0xff;
|
||||
const u32 a3 = (a >> 16) & 0xff;
|
||||
const u32 a4 = (a >> 24) & 0xff;
|
||||
|
||||
const u32 b1 = (b >> 0) & 0xff;
|
||||
const u32 b2 = (b >> 8) & 0xff;
|
||||
const u32 b3 = (b >> 16) & 0xff;
|
||||
const u32 b4 = (b >> 24) & 0xff;
|
||||
|
||||
const u32 r1 = (a1 + b1) & 0xff;
|
||||
const u32 r2 = (a2 + b2) & 0xff;
|
||||
const u32 r3 = (a3 + b3) & 0xff;
|
||||
const u32 r4 = (a4 + b4) & 0xff;
|
||||
|
||||
const u32 r = r1 << 0
|
||||
| r2 << 8
|
||||
| r3 << 16
|
||||
| r4 << 24;
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
DECLSPEC void hmac_sha256_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest)
|
||||
{
|
||||
digest[0] = ipad[0];
|
||||
@@ -89,14 +66,29 @@ DECLSPEC void hmac_sha256_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *i
|
||||
|
||||
__kernel void m13752_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global tc_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global tc_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
|
||||
{
|
||||
const u64 gid = get_global_id (0);
|
||||
const u64 lid = get_local_id (0);
|
||||
const u64 lsz = get_local_size (0);
|
||||
|
||||
/**
|
||||
* keyboard layout shared
|
||||
*/
|
||||
|
||||
__local u32 s_keyboard_layout[256];
|
||||
|
||||
for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
|
||||
{
|
||||
s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
/**
|
||||
* base
|
||||
*/
|
||||
|
||||
const u64 gid = get_global_id (0);
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
u32 w0[4];
|
||||
u32 w1[4];
|
||||
u32 w2[4];
|
||||
@@ -119,6 +111,11 @@ __kernel void m13752_init (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
w3[2] = pws[gid].i[14];
|
||||
w3[3] = pws[gid].i[15];
|
||||
|
||||
keyboard_map (w0, s_keyboard_layout);
|
||||
keyboard_map (w1, s_keyboard_layout);
|
||||
keyboard_map (w2, s_keyboard_layout);
|
||||
keyboard_map (w3, s_keyboard_layout);
|
||||
|
||||
w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]);
|
||||
w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]);
|
||||
w0[2] = u8add (w0[2], esalt_bufs[digests_offset].keyfile_buf[ 2]);
|
||||
|
||||
+26
-29
@@ -17,34 +17,11 @@
|
||||
#include "inc_cipher_twofish.cl"
|
||||
#include "inc_cipher_serpent.cl"
|
||||
|
||||
#include "inc_truecrypt_keyfile.cl"
|
||||
#include "inc_truecrypt_keyboard.cl"
|
||||
#include "inc_truecrypt_crc32.cl"
|
||||
#include "inc_truecrypt_xts.cl"
|
||||
|
||||
DECLSPEC u32 u8add (const u32 a, const u32 b)
|
||||
{
|
||||
const u32 a1 = (a >> 0) & 0xff;
|
||||
const u32 a2 = (a >> 8) & 0xff;
|
||||
const u32 a3 = (a >> 16) & 0xff;
|
||||
const u32 a4 = (a >> 24) & 0xff;
|
||||
|
||||
const u32 b1 = (b >> 0) & 0xff;
|
||||
const u32 b2 = (b >> 8) & 0xff;
|
||||
const u32 b3 = (b >> 16) & 0xff;
|
||||
const u32 b4 = (b >> 24) & 0xff;
|
||||
|
||||
const u32 r1 = (a1 + b1) & 0xff;
|
||||
const u32 r2 = (a2 + b2) & 0xff;
|
||||
const u32 r3 = (a3 + b3) & 0xff;
|
||||
const u32 r4 = (a4 + b4) & 0xff;
|
||||
|
||||
const u32 r = r1 << 0
|
||||
| r2 << 8
|
||||
| r3 << 16
|
||||
| r4 << 24;
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
DECLSPEC void hmac_sha256_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest)
|
||||
{
|
||||
digest[0] = ipad[0];
|
||||
@@ -89,14 +66,29 @@ DECLSPEC void hmac_sha256_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *i
|
||||
|
||||
__kernel void m13753_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global tc_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global tc_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
|
||||
{
|
||||
const u64 gid = get_global_id (0);
|
||||
const u64 lid = get_local_id (0);
|
||||
const u64 lsz = get_local_size (0);
|
||||
|
||||
/**
|
||||
* keyboard layout shared
|
||||
*/
|
||||
|
||||
__local u32 s_keyboard_layout[256];
|
||||
|
||||
for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
|
||||
{
|
||||
s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i];
|
||||
}
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
/**
|
||||
* base
|
||||
*/
|
||||
|
||||
const u64 gid = get_global_id (0);
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
u32 w0[4];
|
||||
u32 w1[4];
|
||||
u32 w2[4];
|
||||
@@ -119,6 +111,11 @@ __kernel void m13753_init (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
w3[2] = pws[gid].i[14];
|
||||
w3[3] = pws[gid].i[15];
|
||||
|
||||
keyboard_map (w0, s_keyboard_layout);
|
||||
keyboard_map (w1, s_keyboard_layout);
|
||||
keyboard_map (w2, s_keyboard_layout);
|
||||
keyboard_map (w3, s_keyboard_layout);
|
||||
|
||||
w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]);
|
||||
w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]);
|
||||
w0[2] = u8add (w0[2], esalt_bufs[digests_offset].keyfile_buf[ 2]);
|
||||
|
||||
+17
-32
@@ -17,34 +17,11 @@
|
||||
#include "inc_cipher_twofish.cl"
|
||||
#include "inc_cipher_serpent.cl"
|
||||
|
||||
#include "inc_truecrypt_keyfile.cl"
|
||||
#include "inc_truecrypt_keyboard.cl"
|
||||
#include "inc_truecrypt_crc32.cl"
|
||||
#include "inc_truecrypt_xts.cl"
|
||||
|
||||
DECLSPEC u32 u8add (const u32 a, const u32 b)
|
||||
{
|
||||
const u32 a1 = (a >> 0) & 0xff;
|
||||
const u32 a2 = (a >> 8) & 0xff;
|
||||
const u32 a3 = (a >> 16) & 0xff;
|
||||
const u32 a4 = (a >> 24) & 0xff;
|
||||
|
||||
const u32 b1 = (b >> 0) & 0xff;
|
||||
const u32 b2 = (b >> 8) & 0xff;
|
||||
const u32 b3 = (b >> 16) & 0xff;
|
||||
const u32 b4 = (b >> 24) & 0xff;
|
||||
|
||||
const u32 r1 = (a1 + b1) & 0xff;
|
||||
const u32 r2 = (a2 + b2) & 0xff;
|
||||
const u32 r3 = (a3 + b3) & 0xff;
|
||||
const u32 r4 = (a4 + b4) & 0xff;
|
||||
|
||||
const u32 r = r1 << 0
|
||||
| r2 << 8
|
||||
| r3 << 16
|
||||
| r4 << 24;
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
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 };
|
||||
@@ -132,17 +109,16 @@ DECLSPEC void hmac_streebog512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u6
|
||||
|
||||
__kernel void m13771_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global vc64_sbog_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const tc_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
|
||||
{
|
||||
/**
|
||||
* base
|
||||
*/
|
||||
|
||||
const u64 gid = get_global_id (0);
|
||||
const u64 lid = get_local_id (0);
|
||||
const u64 lsz = get_local_size (0);
|
||||
|
||||
/**
|
||||
* shared lookup table
|
||||
*/
|
||||
__local u32 s_keyboard_layout[256];
|
||||
|
||||
for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
|
||||
{
|
||||
s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i];
|
||||
}
|
||||
|
||||
#ifdef REAL_SHM
|
||||
|
||||
@@ -170,6 +146,10 @@ __kernel void m13771_init (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
/**
|
||||
* base
|
||||
*/
|
||||
|
||||
u32 w0[4];
|
||||
u32 w1[4];
|
||||
u32 w2[4];
|
||||
@@ -192,6 +172,11 @@ __kernel void m13771_init (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
w3[2] = pws[gid].i[14];
|
||||
w3[3] = pws[gid].i[15];
|
||||
|
||||
keyboard_map (w0, s_keyboard_layout);
|
||||
keyboard_map (w1, s_keyboard_layout);
|
||||
keyboard_map (w2, s_keyboard_layout);
|
||||
keyboard_map (w3, s_keyboard_layout);
|
||||
|
||||
w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]);
|
||||
w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]);
|
||||
w0[2] = u8add (w0[2], esalt_bufs[digests_offset].keyfile_buf[ 2]);
|
||||
|
||||
+17
-32
@@ -17,34 +17,11 @@
|
||||
#include "inc_cipher_twofish.cl"
|
||||
#include "inc_cipher_serpent.cl"
|
||||
|
||||
#include "inc_truecrypt_keyfile.cl"
|
||||
#include "inc_truecrypt_keyboard.cl"
|
||||
#include "inc_truecrypt_crc32.cl"
|
||||
#include "inc_truecrypt_xts.cl"
|
||||
|
||||
DECLSPEC u32 u8add (const u32 a, const u32 b)
|
||||
{
|
||||
const u32 a1 = (a >> 0) & 0xff;
|
||||
const u32 a2 = (a >> 8) & 0xff;
|
||||
const u32 a3 = (a >> 16) & 0xff;
|
||||
const u32 a4 = (a >> 24) & 0xff;
|
||||
|
||||
const u32 b1 = (b >> 0) & 0xff;
|
||||
const u32 b2 = (b >> 8) & 0xff;
|
||||
const u32 b3 = (b >> 16) & 0xff;
|
||||
const u32 b4 = (b >> 24) & 0xff;
|
||||
|
||||
const u32 r1 = (a1 + b1) & 0xff;
|
||||
const u32 r2 = (a2 + b2) & 0xff;
|
||||
const u32 r3 = (a3 + b3) & 0xff;
|
||||
const u32 r4 = (a4 + b4) & 0xff;
|
||||
|
||||
const u32 r = r1 << 0
|
||||
| r2 << 8
|
||||
| r3 << 16
|
||||
| r4 << 24;
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
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 };
|
||||
@@ -132,17 +109,16 @@ DECLSPEC void hmac_streebog512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u6
|
||||
|
||||
__kernel void m13772_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global vc64_sbog_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const tc_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
|
||||
{
|
||||
/**
|
||||
* base
|
||||
*/
|
||||
|
||||
const u64 gid = get_global_id (0);
|
||||
const u64 lid = get_local_id (0);
|
||||
const u64 lsz = get_local_size (0);
|
||||
|
||||
/**
|
||||
* shared lookup table
|
||||
*/
|
||||
__local u32 s_keyboard_layout[256];
|
||||
|
||||
for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
|
||||
{
|
||||
s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i];
|
||||
}
|
||||
|
||||
#ifdef REAL_SHM
|
||||
|
||||
@@ -170,6 +146,10 @@ __kernel void m13772_init (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
/**
|
||||
* base
|
||||
*/
|
||||
|
||||
u32 w0[4];
|
||||
u32 w1[4];
|
||||
u32 w2[4];
|
||||
@@ -192,6 +172,11 @@ __kernel void m13772_init (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
w3[2] = pws[gid].i[14];
|
||||
w3[3] = pws[gid].i[15];
|
||||
|
||||
keyboard_map (w0, s_keyboard_layout);
|
||||
keyboard_map (w1, s_keyboard_layout);
|
||||
keyboard_map (w2, s_keyboard_layout);
|
||||
keyboard_map (w3, s_keyboard_layout);
|
||||
|
||||
w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]);
|
||||
w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]);
|
||||
w0[2] = u8add (w0[2], esalt_bufs[digests_offset].keyfile_buf[ 2]);
|
||||
|
||||
+17
-32
@@ -17,34 +17,11 @@
|
||||
#include "inc_cipher_twofish.cl"
|
||||
#include "inc_cipher_serpent.cl"
|
||||
|
||||
#include "inc_truecrypt_keyfile.cl"
|
||||
#include "inc_truecrypt_keyboard.cl"
|
||||
#include "inc_truecrypt_crc32.cl"
|
||||
#include "inc_truecrypt_xts.cl"
|
||||
|
||||
DECLSPEC u32 u8add (const u32 a, const u32 b)
|
||||
{
|
||||
const u32 a1 = (a >> 0) & 0xff;
|
||||
const u32 a2 = (a >> 8) & 0xff;
|
||||
const u32 a3 = (a >> 16) & 0xff;
|
||||
const u32 a4 = (a >> 24) & 0xff;
|
||||
|
||||
const u32 b1 = (b >> 0) & 0xff;
|
||||
const u32 b2 = (b >> 8) & 0xff;
|
||||
const u32 b3 = (b >> 16) & 0xff;
|
||||
const u32 b4 = (b >> 24) & 0xff;
|
||||
|
||||
const u32 r1 = (a1 + b1) & 0xff;
|
||||
const u32 r2 = (a2 + b2) & 0xff;
|
||||
const u32 r3 = (a3 + b3) & 0xff;
|
||||
const u32 r4 = (a4 + b4) & 0xff;
|
||||
|
||||
const u32 r = r1 << 0
|
||||
| r2 << 8
|
||||
| r3 << 16
|
||||
| r4 << 24;
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
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 };
|
||||
@@ -132,17 +109,16 @@ DECLSPEC void hmac_streebog512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u6
|
||||
|
||||
__kernel void m13773_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global vc64_sbog_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const tc_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
|
||||
{
|
||||
/**
|
||||
* base
|
||||
*/
|
||||
|
||||
const u64 gid = get_global_id (0);
|
||||
const u64 lid = get_local_id (0);
|
||||
const u64 lsz = get_local_size (0);
|
||||
|
||||
/**
|
||||
* shared lookup table
|
||||
*/
|
||||
__local u32 s_keyboard_layout[256];
|
||||
|
||||
for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
|
||||
{
|
||||
s_keyboard_layout[i] = esalt_bufs[digests_offset].keyboard_layout[i];
|
||||
}
|
||||
|
||||
#ifdef REAL_SHM
|
||||
|
||||
@@ -170,6 +146,10 @@ __kernel void m13773_init (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
/**
|
||||
* base
|
||||
*/
|
||||
|
||||
u32 w0[4];
|
||||
u32 w1[4];
|
||||
u32 w2[4];
|
||||
@@ -192,6 +172,11 @@ __kernel void m13773_init (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
w3[2] = pws[gid].i[14];
|
||||
w3[3] = pws[gid].i[15];
|
||||
|
||||
keyboard_map (w0, s_keyboard_layout);
|
||||
keyboard_map (w1, s_keyboard_layout);
|
||||
keyboard_map (w2, s_keyboard_layout);
|
||||
keyboard_map (w3, s_keyboard_layout);
|
||||
|
||||
w0[0] = u8add (w0[0], esalt_bufs[digests_offset].keyfile_buf[ 0]);
|
||||
w0[1] = u8add (w0[1], esalt_bufs[digests_offset].keyfile_buf[ 1]);
|
||||
w0[2] = u8add (w0[2], esalt_bufs[digests_offset].keyfile_buf[ 2]);
|
||||
|
||||
Reference in New Issue
Block a user