__constant u32 shifts3s0[16] = { 1, 1, 2, 2, 2, 2, 2, 2, 1, 2, 2, 2, 2, 2, 2, 1 };
__constant u32 shifts3s1[16] = { 27, 27, 26, 26, 26, 26, 26, 26, 27, 26, 26, 26, 26, 26, 26, 27 };
-#ifdef VECT_SIZE1
-#define BOX(i,n,S) (u32x) ((S)[(n)][(i)])
-#endif
-
-#ifdef VECT_SIZE2
-#define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3])
-#endif
+#define BOX(i,n,S) (S)[(n)][(i)]
static void _des_crypt_keysetup (u32 c, u32 d, u32 Kc[16], u32 Kd[16], __local u32 s_skb[8][64])
{
__constant u32 shifts3s0[16] = { 1, 1, 2, 2, 2, 2, 2, 2, 1, 2, 2, 2, 2, 2, 2, 1 };
__constant u32 shifts3s1[16] = { 27, 27, 26, 26, 26, 26, 26, 26, 27, 26, 26, 26, 26, 26, 26, 27 };
-#ifdef VECT_SIZE1
-#define BOX(i,n,S) (u32x) ((S)[(n)][(i)])
-#endif
-
-#ifdef VECT_SIZE2
-#define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3])
-#endif
+#define BOX(i,n,S) (S)[(n)][(i)]
static void _des_crypt_keysetup (u32 c, u32 d, u32 Kc[16], u32 Kd[16], __local u32 s_skb[8][64])
{
#include "types_ocl.c"
#include "common.c"
-#ifdef VECT_SIZE1
-#define COMPARE_S "check_single_vect1_comp4_bs.c"
-#define COMPARE_M "check_multi_vect1_comp4_bs.c"
-#endif
-
-#ifdef VECT_SIZE2
-#define COMPARE_S "check_single_vect2_comp4_bs.c"
-#define COMPARE_M "check_multi_vect2_comp4_bs.c"
-#endif
-
-#ifdef VECT_SIZE4
-#define COMPARE_S "check_single_vect4_comp4_bs.c"
-#define COMPARE_M "check_multi_vect4_comp4_bs.c"
-#endif
+#define COMPARE_S "check_single_comp4_bs.c"
+#define COMPARE_M "check_multi_comp4_bs.c"
#define KXX_DECL volatile
#define sXXX_DECL volatile
#include "types_ocl.c"
#include "common.c"
-#ifdef VECT_SIZE1
-#define COMPARE_M "check_multi_vect1_comp4.c"
-#endif
+#define COMPARE_S "check_single_comp4.c"
+#define COMPARE_M "check_multi_comp4.c"
#define PUTCHAR64_BE(a,p,c) ((u8 *)(a))[(p) ^ 7] = (u8) (c)
#define GETCHAR64_BE(a,p) ((u8 *)(a))[(p) ^ 7]
u64 pw[2];
- pw[0] = swap32 (hl32_to_64 (w0[1], w0[0]));
- pw[1] = swap32 (hl32_to_64 (w0[3], w0[2]));
+ pw[0] = swap64 (hl32_to_64 (w0[1], w0[0]));
+ pw[1] = swap64 (hl32_to_64 (w0[3], w0[2]));
u64 salt[2];
- salt[0] = swap32 (hl32_to_64 (salt_buf[1], salt_buf[0]));
- salt[1] = swap32 (hl32_to_64 (salt_buf[3], salt_buf[2]));
+ salt[0] = swap64 (hl32_to_64 (salt_buf[1], salt_buf[0]));
+ salt[1] = swap64 (hl32_to_64 (salt_buf[3], salt_buf[2]));
/**
* begin
const u32 lid = get_local_id (0);
- const u64 a = swap32 (tmps[gid].l_alt_result[0]);
- const u64 b = swap32 (tmps[gid].l_alt_result[1]);
+ const u64 a = swap64 (tmps[gid].l_alt_result[0]);
+ const u64 b = swap64 (tmps[gid].l_alt_result[1]);
const u32 r0 = l32_from_64 (a);
const u32 r1 = h32_from_64 (a);
#include "types_ocl.c"
#include "common.c"
-#ifdef VECT_SIZE1
-#define COMPARE_M "check_multi_vect1_comp4.c"
-#endif
-
-#ifdef VECT_SIZE2
-#define COMPARE_M "check_multi_vect2_comp4.c"
-#endif
+#define COMPARE_S "check_single_comp4.c"
+#define COMPARE_M "check_multi_comp4.c"
static void md4_transform (const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], u32 digest[4])
{
#define COMPARE_S "check_single_comp4.c"
#define COMPARE_M "check_multi_comp4.c"
-#ifdef VECT_SIZE1
#define uint_to_hex_lower8(i) l_bin2asc[(i)]
-#endif
-
-#ifdef VECT_SIZE2
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3])
-#endif
__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m02610_m04 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
#define COMPARE_S "check_single_comp4.c"
#define COMPARE_M "check_multi_comp4.c"
-#ifdef VECT_SIZE1
#define uint_to_hex_lower8(i) l_bin2asc[(i)]
-#endif
-
-#ifdef VECT_SIZE2
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3])
-#endif
__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m02610_m04 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
#define COMPARE_S "check_single_comp4.c"
#define COMPARE_M "check_multi_comp4.c"
-#ifdef VECT_SIZE1
#define uint_to_hex_lower8(i) l_bin2asc[(i)]
-#endif
-
-#ifdef VECT_SIZE2
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3])
-#endif
static void m02610m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, __local u32 l_bin2asc[256])
{
#define COMPARE_S "check_single_comp4.c"
#define COMPARE_M "check_multi_comp4.c"
-#ifdef VECT_SIZE1
#define uint_to_hex_lower8(i) l_bin2asc[(i)]
-#endif
-
-#ifdef VECT_SIZE2
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3])
-#endif
__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m02710_m04 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
#define COMPARE_S "check_single_comp4.c"
#define COMPARE_M "check_multi_comp4.c"
-#ifdef VECT_SIZE1
#define uint_to_hex_lower8(i) l_bin2asc[(i)]
-#endif
-
-#ifdef VECT_SIZE1
-#define uint_to_hex_lower8(i) l_bin2asc[(i)]
-#endif
-
-#ifdef VECT_SIZE2
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3])
-#endif
__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m02710_m04 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
#define COMPARE_S "check_single_comp4.c"
#define COMPARE_M "check_multi_comp4.c"
-#ifdef VECT_SIZE1
#define uint_to_hex_lower8(i) l_bin2asc[(i)]
-#endif
-
-#ifdef VECT_SIZE2
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3])
-#endif
static void m02710m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, __local u32 l_bin2asc[256])
{
#define COMPARE_S "check_single_comp4.c"
#define COMPARE_M "check_multi_comp4.c"
-#ifdef VECT_SIZE1
#define uint_to_hex_lower8(i) l_bin2asc[(i)]
-#endif
-
-#ifdef VECT_SIZE2
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3])
-#endif
__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m02810_m04 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
#define COMPARE_S "check_single_comp4.c"
#define COMPARE_M "check_multi_comp4.c"
-#ifdef VECT_SIZE1
#define uint_to_hex_lower8(i) l_bin2asc[(i)]
-#endif
-
-#ifdef VECT_SIZE2
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3])
-#endif
__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m02810_m04 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
#define COMPARE_S "check_single_comp4.c"
#define COMPARE_M "check_multi_comp4.c"
-#ifdef VECT_SIZE1
#define uint_to_hex_lower8(i) l_bin2asc[(i)]
-#endif
-
-#ifdef VECT_SIZE2
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3])
-#endif
static void m02810m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, __local u32 l_bin2asc[256])
{
#define LM_IV_0_IP_RR3 0x2400b807
#define LM_IV_1_IP_RR3 0xaa190747
-#ifdef VECT_SIZE1
-#define BOX(i,n,S) u32 ((S)[(n)][(i)])
-#endif
-
-#ifdef VECT_SIZE2
-#define BOX(i,n,S) u32 ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define BOX(i,n,S) u32 ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3])
-#endif
+#define BOX(i,n,S) (S)[(n)][(i)]
static void _des_crypt_encrypt (u32 iv[2], u32 data[2], u32 Kc[16], u32 Kd[16], __local u32 s_SPtrans[8][64])
{
static void transform_netntlmv1_key (const u32 w0, const u32 w1, u32 out[2])
{
- #ifdef VECT_SIZE1
const uchar4 t0 = as_uchar4 (w0);
const uchar4 t1 = as_uchar4 (w1);
out[0] = as_uint (k0);
out[1] = as_uint (k1);
- #endif
-
- #ifdef VECT_SIZE2
- const uchar8 t0 = as_uchar8 (w0);
- const uchar8 t1 = as_uchar8 (w1);
-
- uchar8 k0;
- uchar8 k1;
-
- k0.s0 = (t0.s0 >> 0);
- k0.s1 = (t0.s0 << 7) | (t0.s1 >> 1);
- k0.s2 = (t0.s1 << 6) | (t0.s2 >> 2);
- k0.s3 = (t0.s2 << 5) | (t0.s3 >> 3);
- k1.s0 = (t0.s3 << 4) | (t1.s0 >> 4);
- k1.s1 = (t1.s0 << 3) | (t1.s1 >> 5);
- k1.s2 = (t1.s1 << 2) | (t1.s2 >> 6);
- k1.s3 = (t1.s2 << 1);
-
- k0.s4 = (t0.s4 >> 0);
- k0.s5 = (t0.s4 << 7) | (t0.s5 >> 1);
- k0.s6 = (t0.s5 << 6) | (t0.s6 >> 2);
- k0.s7 = (t0.s6 << 5) | (t0.s7 >> 3);
- k1.s4 = (t0.s7 << 4) | (t1.s4 >> 4);
- k1.s5 = (t1.s4 << 3) | (t1.s5 >> 5);
- k1.s6 = (t1.s5 << 2) | (t1.s6 >> 6);
- k1.s7 = (t1.s6 << 1);
-
- out[0] = as_uint2 (k0);
- out[1] = as_uint2 (k1);
- #endif
-
- #ifdef VECT_SIZE4
- const uchar16 t0 = as_uchar16 (w0);
- const uchar16 t1 = as_uchar16 (w1);
-
- uchar16 k0;
- uchar16 k1;
-
- k0.s0 = (t0.s0 >> 0);
- k0.s1 = (t0.s0 << 7) | (t0.s1 >> 1);
- k0.s2 = (t0.s1 << 6) | (t0.s2 >> 2);
- k0.s3 = (t0.s2 << 5) | (t0.s3 >> 3);
- k1.s0 = (t0.s3 << 4) | (t1.s0 >> 4);
- k1.s1 = (t1.s0 << 3) | (t1.s1 >> 5);
- k1.s2 = (t1.s1 << 2) | (t1.s2 >> 6);
- k1.s3 = (t1.s2 << 1);
-
- k0.s4 = (t0.s4 >> 0);
- k0.s5 = (t0.s4 << 7) | (t0.s5 >> 1);
- k0.s6 = (t0.s5 << 6) | (t0.s6 >> 2);
- k0.s7 = (t0.s6 << 5) | (t0.s7 >> 3);
- k1.s4 = (t0.s7 << 4) | (t1.s4 >> 4);
- k1.s5 = (t1.s4 << 3) | (t1.s5 >> 5);
- k1.s6 = (t1.s5 << 2) | (t1.s6 >> 6);
- k1.s7 = (t1.s6 << 1);
-
- k0.s8 = (t0.s8 >> 0);
- k0.s9 = (t0.s8 << 7) | (t0.s9 >> 1);
- k0.sa = (t0.s9 << 6) | (t0.sa >> 2);
- k0.sb = (t0.sa << 5) | (t0.sb >> 3);
- k1.s8 = (t0.sb << 4) | (t1.s8 >> 4);
- k1.s9 = (t1.s8 << 3) | (t1.s9 >> 5);
- k1.sa = (t1.s9 << 2) | (t1.sa >> 6);
- k1.sb = (t1.sa << 1);
-
- k0.sc = (t0.sc >> 0);
- k0.sd = (t0.sc << 7) | (t0.sd >> 1);
- k0.se = (t0.sd << 6) | (t0.se >> 2);
- k0.sf = (t0.se << 5) | (t0.sf >> 3);
- k1.sc = (t0.sf << 4) | (t1.sc >> 4);
- k1.sd = (t1.sc << 3) | (t1.sd >> 5);
- k1.se = (t1.sd << 2) | (t1.se >> 6);
- k1.sf = (t1.se << 1);
-
- out[0] = as_uint4 (k0);
- out[1] = as_uint4 (k1);
- #endif
}
__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m03000_m04 (__global pw_t *pws, __global gpu_rule_t * rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
#define LM_IV_0_IP_RR3 0x2400b807
#define LM_IV_1_IP_RR3 0xaa190747
-#ifdef VECT_SIZE1
-#define BOX(i,n,S) u32 ((S)[(n)][(i)])
-#endif
-
-#ifdef VECT_SIZE2
-#define BOX(i,n,S) u32 ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define BOX(i,n,S) u32 ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3])
-#endif
+#define BOX(i,n,S) (S)[(n)][(i)]
static void _des_crypt_encrypt (u32 iv[2], u32 data[2], u32 Kc[16], u32 Kd[16], __local u32 s_SPtrans[8][64])
{
static void transform_netntlmv1_key (const u32 w0, const u32 w1, u32 out[2])
{
- #ifdef VECT_SIZE1
const uchar4 t0 = as_uchar4 (w0);
const uchar4 t1 = as_uchar4 (w1);
out[0] = as_uint (k0);
out[1] = as_uint (k1);
- #endif
-
- #ifdef VECT_SIZE2
- const uchar8 t0 = as_uchar8 (w0);
- const uchar8 t1 = as_uchar8 (w1);
-
- uchar8 k0;
- uchar8 k1;
-
- k0.s0 = (t0.s0 >> 0);
- k0.s1 = (t0.s0 << 7) | (t0.s1 >> 1);
- k0.s2 = (t0.s1 << 6) | (t0.s2 >> 2);
- k0.s3 = (t0.s2 << 5) | (t0.s3 >> 3);
- k1.s0 = (t0.s3 << 4) | (t1.s0 >> 4);
- k1.s1 = (t1.s0 << 3) | (t1.s1 >> 5);
- k1.s2 = (t1.s1 << 2) | (t1.s2 >> 6);
- k1.s3 = (t1.s2 << 1);
-
- k0.s4 = (t0.s4 >> 0);
- k0.s5 = (t0.s4 << 7) | (t0.s5 >> 1);
- k0.s6 = (t0.s5 << 6) | (t0.s6 >> 2);
- k0.s7 = (t0.s6 << 5) | (t0.s7 >> 3);
- k1.s4 = (t0.s7 << 4) | (t1.s4 >> 4);
- k1.s5 = (t1.s4 << 3) | (t1.s5 >> 5);
- k1.s6 = (t1.s5 << 2) | (t1.s6 >> 6);
- k1.s7 = (t1.s6 << 1);
-
- out[0] = as_uint2 (k0);
- out[1] = as_uint2 (k1);
- #endif
-
- #ifdef VECT_SIZE4
- const uchar16 t0 = as_uchar16 (w0);
- const uchar16 t1 = as_uchar16 (w1);
-
- uchar16 k0;
- uchar16 k1;
-
- k0.s0 = (t0.s0 >> 0);
- k0.s1 = (t0.s0 << 7) | (t0.s1 >> 1);
- k0.s2 = (t0.s1 << 6) | (t0.s2 >> 2);
- k0.s3 = (t0.s2 << 5) | (t0.s3 >> 3);
- k1.s0 = (t0.s3 << 4) | (t1.s0 >> 4);
- k1.s1 = (t1.s0 << 3) | (t1.s1 >> 5);
- k1.s2 = (t1.s1 << 2) | (t1.s2 >> 6);
- k1.s3 = (t1.s2 << 1);
-
- k0.s4 = (t0.s4 >> 0);
- k0.s5 = (t0.s4 << 7) | (t0.s5 >> 1);
- k0.s6 = (t0.s5 << 6) | (t0.s6 >> 2);
- k0.s7 = (t0.s6 << 5) | (t0.s7 >> 3);
- k1.s4 = (t0.s7 << 4) | (t1.s4 >> 4);
- k1.s5 = (t1.s4 << 3) | (t1.s5 >> 5);
- k1.s6 = (t1.s5 << 2) | (t1.s6 >> 6);
- k1.s7 = (t1.s6 << 1);
-
- k0.s8 = (t0.s8 >> 0);
- k0.s9 = (t0.s8 << 7) | (t0.s9 >> 1);
- k0.sa = (t0.s9 << 6) | (t0.sa >> 2);
- k0.sb = (t0.sa << 5) | (t0.sb >> 3);
- k1.s8 = (t0.sb << 4) | (t1.s8 >> 4);
- k1.s9 = (t1.s8 << 3) | (t1.s9 >> 5);
- k1.sa = (t1.s9 << 2) | (t1.sa >> 6);
- k1.sb = (t1.sa << 1);
-
- k0.sc = (t0.sc >> 0);
- k0.sd = (t0.sc << 7) | (t0.sd >> 1);
- k0.se = (t0.sd << 6) | (t0.se >> 2);
- k0.sf = (t0.se << 5) | (t0.sf >> 3);
- k1.sc = (t0.sf << 4) | (t1.sc >> 4);
- k1.sd = (t1.sc << 3) | (t1.sd >> 5);
- k1.se = (t1.sd << 2) | (t1.se >> 6);
- k1.sf = (t1.se << 1);
-
- out[0] = as_uint4 (k0);
- out[1] = as_uint4 (k1);
- #endif
}
__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m03000_m04 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
#include "types_ocl.c"
#include "common.c"
-#ifdef VECT_SIZE1
-#define COMPARE_S "check_single_vect1_comp4_bs.c"
-#define COMPARE_M "check_multi_vect1_comp4_bs.c"
-#endif
-
-#ifdef VECT_SIZE2
-#define COMPARE_S "check_single_vect2_comp4_bs.c"
-#define COMPARE_M "check_multi_vect2_comp4_bs.c"
-#endif
-
-#ifdef VECT_SIZE4
-#define COMPARE_S "check_single_vect4_comp4_bs.c"
-#define COMPARE_M "check_multi_vect4_comp4_bs.c"
-#endif
+#define COMPARE_S "check_single_comp4_bs.c"
+#define COMPARE_M "check_multi_comp4_bs.c"
#define KXX_DECL
__constant u32 shifts3s0[16] = { 1, 1, 2, 2, 2, 2, 2, 2, 1, 2, 2, 2, 2, 2, 2, 1 };
__constant u32 shifts3s1[16] = { 27, 27, 26, 26, 26, 26, 26, 26, 27, 26, 26, 26, 26, 26, 26, 27 };
-#ifdef VECT_SIZE1
-#define BOX(i,n,S) u32 ((S)[(n)][(i)])
-#endif
-
-#ifdef VECT_SIZE2
-#define BOX(i,n,S) u32 ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define BOX(i,n,S) u32 ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3])
-#endif
+#define BOX(i,n,S) (S)[(n)][(i)]
static void _des_crypt_encrypt (u32 iv[2], u32 data[2], u32 Kc[16], u32 Kd[16], __local u32 s_SPtrans[8][64])
{
__constant u32 shifts3s0[16] = { 1, 1, 2, 2, 2, 2, 2, 2, 1, 2, 2, 2, 2, 2, 2, 1 };
__constant u32 shifts3s1[16] = { 27, 27, 26, 26, 26, 26, 26, 26, 27, 26, 26, 26, 26, 26, 26, 27 };
-#ifdef VECT_SIZE1
-#define BOX(i,n,S) u32 ((S)[(n)][(i)])
-#endif
-
-#ifdef VECT_SIZE2
-#define BOX(i,n,S) u32 ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define BOX(i,n,S) u32 ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3])
-#endif
+#define BOX(i,n,S) (S)[(n)][(i)]
static void _des_crypt_encrypt (u32 iv[2], u32 data[2], u32 Kc[16], u32 Kd[16], __local u32 s_SPtrans[8][64])
{
__constant u32 shifts3s0[16] = { 1, 1, 2, 2, 2, 2, 2, 2, 1, 2, 2, 2, 2, 2, 2, 1 };
__constant u32 shifts3s1[16] = { 27, 27, 26, 26, 26, 26, 26, 26, 27, 26, 26, 26, 26, 26, 26, 27 };
-#ifdef VECT_SIZE1
-#define BOX(i,n,S) u32 ((S)[(n)][(i)])
-#endif
-
-#ifdef VECT_SIZE2
-#define BOX(i,n,S) u32 ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define BOX(i,n,S) u32 ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3])
-#endif
+#define BOX(i,n,S) (S)[(n)][(i)]
static void _des_crypt_encrypt (u32 iv[2], u32 data[2], u32 Kc[16], u32 Kd[16], __local u32 s_SPtrans[8][64])
{
#include "types_ocl.c"
#include "common.c"
-#ifdef VECT_SIZE1
-#define COMPARE_M "check_multi_vect1_comp4.c"
-#endif
+#define COMPARE_S "check_single_comp4.c"
+#define COMPARE_M "check_multi_comp4.c"
// http://www.schneier.com/code/constants.txt
#define COMPARE_S "check_single_comp4.c"
#define COMPARE_M "check_multi_comp4.c"
-#ifdef VECT_SIZE1
#define uint_to_hex_lower8(i) l_bin2asc[(i)]
-#endif
-
-#ifdef VECT_SIZE2
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3])
-#endif
__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m03710_m04 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
#define COMPARE_S "check_single_comp4.c"
#define COMPARE_M "check_multi_comp4.c"
-#ifdef VECT_SIZE1
#define uint_to_hex_lower8(i) l_bin2asc[(i)]
-#endif
-
-#ifdef VECT_SIZE2
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3])
-#endif
__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m03710_m04 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
#define COMPARE_S "check_single_comp4.c"
#define COMPARE_M "check_multi_comp4.c"
-#ifdef VECT_SIZE1
#define uint_to_hex_lower8(i) l_bin2asc[(i)]
-#endif
-
-#ifdef VECT_SIZE2
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3])
-#endif
static void m03710m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, __local u32 l_bin2asc[256])
{
#define COMPARE_S "check_single_comp4.c"
#define COMPARE_M "check_multi_comp4.c"
-#ifdef VECT_SIZE1
#define uint_to_hex_lower8(i) l_bin2asc[(i)]
-#endif
-
-#ifdef VECT_SIZE2
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3])
-#endif
__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m04310_m04 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
#define COMPARE_S "check_single_comp4.c"
#define COMPARE_M "check_multi_comp4.c"
-#ifdef VECT_SIZE1
#define uint_to_hex_lower8(i) l_bin2asc[(i)]
-#endif
-
-#ifdef VECT_SIZE2
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3])
-#endif
__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m04310_m04 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
#define COMPARE_S "check_single_comp4.c"
#define COMPARE_M "check_multi_comp4.c"
-#ifdef VECT_SIZE1
#define uint_to_hex_lower8(i) l_bin2asc[(i)]
-#endif
-
-#ifdef VECT_SIZE2
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3])
-#endif
static void m04310m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, __local u32 l_bin2asc[256])
{
#define COMPARE_S "check_single_comp4.c"
#define COMPARE_M "check_multi_comp4.c"
-#ifdef VECT_SIZE1
#define uint_to_hex_lower8(i) l_bin2asc[(i)]
-#endif
-
-#ifdef VECT_SIZE2
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3])
-#endif
__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m04400_m04 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
#define COMPARE_S "check_single_comp4.c"
#define COMPARE_M "check_multi_comp4.c"
-#ifdef VECT_SIZE1
#define uint_to_hex_lower8(i) l_bin2asc[(i)]
-#endif
-
-#ifdef VECT_SIZE2
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3])
-#endif
__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m04400_m04 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
#define COMPARE_S "check_single_comp4.c"
#define COMPARE_M "check_multi_comp4.c"
-#ifdef VECT_SIZE1
#define uint_to_hex_lower8(i) l_bin2asc[(i)]
-#endif
-
-#ifdef VECT_SIZE2
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3])
-#endif
static void m04400m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, __local u32 l_bin2asc[256])
{
#define COMPARE_S "check_single_comp4.c"
#define COMPARE_M "check_multi_comp4.c"
-#ifdef VECT_SIZE1
#define uint_to_hex_lower8_le(i) l_bin2asc[(i)]
-#endif
-
-#ifdef VECT_SIZE2
-#define uint_to_hex_lower8_le(i) u32 (l_bin2asc[(i).s0], l_bin2asc[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define uint_to_hex_lower8_le(i) u32 (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3])
-#endif
__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m04500_m04 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
#define COMPARE_S "check_single_comp4.c"
#define COMPARE_M "check_multi_comp4.c"
-#ifdef VECT_SIZE1
#define uint_to_hex_lower8_le(i) l_bin2asc[(i)]
-#endif
-
-#ifdef VECT_SIZE2
-#define uint_to_hex_lower8_le(i) u32 (l_bin2asc[(i).s0], l_bin2asc[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define uint_to_hex_lower8_le(i) u32 (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3])
-#endif
__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m04500_m04 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
#define COMPARE_S "check_single_comp4.c"
#define COMPARE_M "check_multi_comp4.c"
-#ifdef VECT_SIZE1
#define uint_to_hex_lower8_le(i) l_bin2asc[(i)]
-#endif
-
-#ifdef VECT_SIZE2
-#define uint_to_hex_lower8_le(i) u32 (l_bin2asc[(i).s0], l_bin2asc[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define uint_to_hex_lower8_le(i) u32 (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3])
-#endif
static void m04500m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, __local u32 l_bin2asc[256])
{
#define COMPARE_S "check_single_comp4.c"
#define COMPARE_M "check_multi_comp4.c"
-#ifdef VECT_SIZE1
#define uint_to_hex_lower8_le(i) l_bin2asc[(i)]
-#endif
-
-#ifdef VECT_SIZE2
-#define uint_to_hex_lower8_le(i) u32 (l_bin2asc[(i).s0], l_bin2asc[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define uint_to_hex_lower8_le(i) u32 (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3])
-#endif
__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m04700_m04 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
#define COMPARE_S "check_single_comp4.c"
#define COMPARE_M "check_multi_comp4.c"
-#ifdef VECT_SIZE1
#define uint_to_hex_lower8_le(i) l_bin2asc[(i)]
-#endif
-
-#ifdef VECT_SIZE2
-#define uint_to_hex_lower8_le(i) u32 (l_bin2asc[(i).s0], l_bin2asc[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define uint_to_hex_lower8_le(i) u32 (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3])
-#endif
__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m04700_m04 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
#define COMPARE_S "check_single_comp4.c"
#define COMPARE_M "check_multi_comp4.c"
-#ifdef VECT_SIZE1
#define uint_to_hex_lower8_le(i) l_bin2asc[(i)]
-#endif
-
-#ifdef VECT_SIZE2
-#define uint_to_hex_lower8_le(i) u32 (l_bin2asc[(i).s0], l_bin2asc[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define uint_to_hex_lower8_le(i) u32 (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3])
-#endif
static void m04700m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, __local u32 l_bin2asc[256])
{
#define Rho_Pi(s) \
{ \
- u32 j = keccakf_piln[s]; \
- u32 k = keccakf_rotc[s]; \
+ u32 j = keccakf_piln[s]; \
+ u32 k = keccakf_rotc[s]; \
bc0 = st[j]; \
st[j] = rotl64 (t, k); \
t = bc0; \
u64 st[25];
- #ifdef VECT_SIZE1
- st[ 0] = (u64x) (w0[0]) | (u64x) (w0[1]) << 32;
- st[ 1] = (u64x) (w0[2]) | (u64x) (w0[3]) << 32;
- st[ 2] = (u64x) (w1[0]) | (u64x) (w1[1]) << 32;
- st[ 3] = (u64x) (w1[2]) | (u64x) (w1[3]) << 32;
- #endif
-
- #ifdef VECT_SIZE2
- st[ 0] = (u64x) (w0[0].s0, w0[0].s1) | (u64x) (w0[1].s0, w0[1].s1) << 32;
- st[ 1] = (u64x) (w0[2].s0, w0[2].s1) | (u64x) (w0[3].s0, w0[3].s1) << 32;
- st[ 2] = (u64x) (w1[0].s0, w1[0].s1) | (u64x) (w1[1].s0, w1[1].s1) << 32;
- st[ 3] = (u64x) (w1[2].s0, w1[2].s1) | (u64x) (w1[3].s0, w1[3].s1) << 32;
- #endif
-
+ st[ 0] = (u64) (w0[0]) | (u64) (w0[1]) << 32;
+ st[ 1] = (u64) (w0[2]) | (u64) (w0[3]) << 32;
+ st[ 2] = (u64) (w1[0]) | (u64) (w1[1]) << 32;
+ st[ 3] = (u64) (w1[2]) | (u64) (w1[3]) << 32;
st[ 4] = 0;
st[ 5] = 0;
st[ 6] = 0;
u64 st[25];
- #ifdef VECT_SIZE1
- st[ 0] = (u64x) (w0[0]) | (u64x) (w0[1]) << 32;
- st[ 1] = (u64x) (w0[2]) | (u64x) (w0[3]) << 32;
- st[ 2] = (u64x) (w1[0]) | (u64x) (w1[1]) << 32;
- st[ 3] = (u64x) (w1[2]) | (u64x) (w1[3]) << 32;
- #endif
-
- #ifdef VECT_SIZE2
- st[ 0] = (u64x) (w0[0].s0, w0[0].s1) | (u64x) (w0[1].s0, w0[1].s1) << 32;
- st[ 1] = (u64x) (w0[2].s0, w0[2].s1) | (u64x) (w0[3].s0, w0[3].s1) << 32;
- st[ 2] = (u64x) (w1[0].s0, w1[0].s1) | (u64x) (w1[1].s0, w1[1].s1) << 32;
- st[ 3] = (u64x) (w1[2].s0, w1[2].s1) | (u64x) (w1[3].s0, w1[3].s1) << 32;
- #endif
-
+ st[ 0] = (u64) (w0[0]) | (u64) (w0[1]) << 32;
+ st[ 1] = (u64) (w0[2]) | (u64) (w0[3]) << 32;
+ st[ 2] = (u64) (w1[0]) | (u64) (w1[1]) << 32;
+ st[ 3] = (u64) (w1[2]) | (u64) (w1[3]) << 32;
st[ 4] = 0;
st[ 5] = 0;
st[ 6] = 0;
u64 st[25];
- #ifdef VECT_SIZE1
- st[ 0] = (u64x) (w0[0]) | (u64x) (w0[1]) << 32;
- st[ 1] = (u64x) (w0[2]) | (u64x) (w0[3]) << 32;
- st[ 2] = (u64x) (w1[0]) | (u64x) (w1[1]) << 32;
- st[ 3] = (u64x) (w1[2]) | (u64x) (w1[3]) << 32;
- #endif
-
- #ifdef VECT_SIZE2
- st[ 0] = (u64x) (w0[0].s0, w0[0].s1) | (u64x) (w0[1].s0, w0[1].s1) << 32;
- st[ 1] = (u64x) (w0[2].s0, w0[2].s1) | (u64x) (w0[3].s0, w0[3].s1) << 32;
- st[ 2] = (u64x) (w1[0].s0, w1[0].s1) | (u64x) (w1[1].s0, w1[1].s1) << 32;
- st[ 3] = (u64x) (w1[2].s0, w1[2].s1) | (u64x) (w1[3].s0, w1[3].s1) << 32;
- #endif
-
+ st[ 0] = (u64) (w0[0]) | (u64) (w0[1]) << 32;
+ st[ 1] = (u64) (w0[2]) | (u64) (w0[3]) << 32;
+ st[ 2] = (u64) (w1[0]) | (u64) (w1[1]) << 32;
+ st[ 3] = (u64) (w1[2]) | (u64) (w1[3]) << 32;
st[ 4] = 0;
st[ 5] = 0;
st[ 6] = 0;
u64 st[25];
- #ifdef VECT_SIZE1
- st[ 0] = (u64x) (w0[0]) | (u64x) (w0[1]) << 32;
- st[ 1] = (u64x) (w0[2]) | (u64x) (w0[3]) << 32;
- st[ 2] = (u64x) (w1[0]) | (u64x) (w1[1]) << 32;
- st[ 3] = (u64x) (w1[2]) | (u64x) (w1[3]) << 32;
- #endif
-
- #ifdef VECT_SIZE2
- st[ 0] = (u64x) (w0[0].s0, w0[0].s1) | (u64x) (w0[1].s0, w0[1].s1) << 32;
- st[ 1] = (u64x) (w0[2].s0, w0[2].s1) | (u64x) (w0[3].s0, w0[3].s1) << 32;
- st[ 2] = (u64x) (w1[0].s0, w1[0].s1) | (u64x) (w1[1].s0, w1[1].s1) << 32;
- st[ 3] = (u64x) (w1[2].s0, w1[2].s1) | (u64x) (w1[3].s0, w1[3].s1) << 32;
- #endif
-
+ st[ 0] = (u64) (w0[0]) | (u64) (w0[1]) << 32;
+ st[ 1] = (u64) (w0[2]) | (u64) (w0[3]) << 32;
+ st[ 2] = (u64) (w1[0]) | (u64) (w1[1]) << 32;
+ st[ 3] = (u64) (w1[2]) | (u64) (w1[3]) << 32;
st[ 4] = 0;
st[ 5] = 0;
st[ 6] = 0;
u64 st[25];
- #ifdef VECT_SIZE1
- st[ 0] = (u64x) (w0[0]) | (u64x) (w0[1]) << 32;
- st[ 1] = (u64x) (w0[2]) | (u64x) (w0[3]) << 32;
- st[ 2] = (u64x) (w1[0]) | (u64x) (w1[1]) << 32;
- st[ 3] = (u64x) (w1[2]) | (u64x) (w1[3]) << 32;
- st[ 4] = (u64x) (w2[0]) | (u64x) (w2[1]) << 32;
- st[ 5] = (u64x) (w2[2]) | (u64x) (w2[3]) << 32;
- st[ 6] = (u64x) (w3[0]) | (u64x) (w3[1]) << 32;
- st[ 7] = (u64x) (w3[2]) | (u64x) (w3[3]) << 32;
- #endif
-
- #ifdef VECT_SIZE2
- st[ 0] = (u64x) (w0[0].s0, w0[0].s1) | (u64x) (w0[1].s0, w0[1].s1) << 32;
- st[ 1] = (u64x) (w0[2].s0, w0[2].s1) | (u64x) (w0[3].s0, w0[3].s1) << 32;
- st[ 2] = (u64x) (w1[0].s0, w1[0].s1) | (u64x) (w1[1].s0, w1[1].s1) << 32;
- st[ 3] = (u64x) (w1[2].s0, w1[2].s1) | (u64x) (w1[3].s0, w1[3].s1) << 32;
- st[ 4] = (u64x) (w2[0].s0, w2[0].s1) | (u64x) (w2[1].s0, w2[1].s1) << 32;
- st[ 5] = (u64x) (w2[2].s0, w2[2].s1) | (u64x) (w2[3].s0, w2[3].s1) << 32;
- st[ 6] = (u64x) (w3[0].s0, w3[0].s1) | (u64x) (w3[1].s0, w3[1].s1) << 32;
- st[ 7] = (u64x) (w3[2].s0, w3[2].s1) | (u64x) (w3[3].s0, w3[3].s1) << 32;
- #endif
-
+ st[ 0] = (u64) (w0[0]) | (u64) (w0[1]) << 32;
+ st[ 1] = (u64) (w0[2]) | (u64) (w0[3]) << 32;
+ st[ 2] = (u64) (w1[0]) | (u64) (w1[1]) << 32;
+ st[ 3] = (u64) (w1[2]) | (u64) (w1[3]) << 32;
+ st[ 4] = (u64) (w2[0]) | (u64) (w2[1]) << 32;
+ st[ 5] = (u64) (w2[2]) | (u64) (w2[3]) << 32;
+ st[ 6] = (u64) (w3[0]) | (u64) (w3[1]) << 32;
+ st[ 7] = (u64) (w3[2]) | (u64) (w3[3]) << 32;
st[ 8] = 0;
st[ 9] = 0;
st[10] = 0;
u64 st[25];
- #ifdef VECT_SIZE1
- st[ 0] = (u64x) (w0[0]) | (u64x) (w0[1]) << 32;
- st[ 1] = (u64x) (w0[2]) | (u64x) (w0[3]) << 32;
- st[ 2] = (u64x) (w1[0]) | (u64x) (w1[1]) << 32;
- st[ 3] = (u64x) (w1[2]) | (u64x) (w1[3]) << 32;
- st[ 4] = (u64x) (w2[0]) | (u64x) (w2[1]) << 32;
- st[ 5] = (u64x) (w2[2]) | (u64x) (w2[3]) << 32;
- st[ 6] = (u64x) (w3[0]) | (u64x) (w3[1]) << 32;
- st[ 7] = (u64x) (w3[2]) | (u64x) (w3[3]) << 32;
- #endif
-
- #ifdef VECT_SIZE2
- st[ 0] = (u64x) (w0[0].s0, w0[0].s1) | (u64x) (w0[1].s0, w0[1].s1) << 32;
- st[ 1] = (u64x) (w0[2].s0, w0[2].s1) | (u64x) (w0[3].s0, w0[3].s1) << 32;
- st[ 2] = (u64x) (w1[0].s0, w1[0].s1) | (u64x) (w1[1].s0, w1[1].s1) << 32;
- st[ 3] = (u64x) (w1[2].s0, w1[2].s1) | (u64x) (w1[3].s0, w1[3].s1) << 32;
- st[ 4] = (u64x) (w2[0].s0, w2[0].s1) | (u64x) (w2[1].s0, w2[1].s1) << 32;
- st[ 5] = (u64x) (w2[2].s0, w2[2].s1) | (u64x) (w2[3].s0, w2[3].s1) << 32;
- st[ 6] = (u64x) (w3[0].s0, w3[0].s1) | (u64x) (w3[1].s0, w3[1].s1) << 32;
- st[ 7] = (u64x) (w3[2].s0, w3[2].s1) | (u64x) (w3[3].s0, w3[3].s1) << 32;
- #endif
-
+ st[ 0] = (u64) (w0[0]) | (u64) (w0[1]) << 32;
+ st[ 1] = (u64) (w0[2]) | (u64) (w0[3]) << 32;
+ st[ 2] = (u64) (w1[0]) | (u64) (w1[1]) << 32;
+ st[ 3] = (u64) (w1[2]) | (u64) (w1[3]) << 32;
+ st[ 4] = (u64) (w2[0]) | (u64) (w2[1]) << 32;
+ st[ 5] = (u64) (w2[2]) | (u64) (w2[3]) << 32;
+ st[ 6] = (u64) (w3[0]) | (u64) (w3[1]) << 32;
+ st[ 7] = (u64) (w3[2]) | (u64) (w3[3]) << 32;
st[ 8] = 0;
st[ 9] = 0;
st[10] = 0;
__constant u32 shifts3s0[16] = { 1, 1, 2, 2, 2, 2, 2, 2, 1, 2, 2, 2, 2, 2, 2, 1 };
__constant u32 shifts3s1[16] = { 27, 27, 26, 26, 26, 26, 26, 26, 27, 26, 26, 26, 26, 26, 26, 27 };
-#ifdef VECT_SIZE1
-#define BOX(i,n,S) (u32x) ((S)[(n)][(i)])
-#endif
-
-#ifdef VECT_SIZE2
-#define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3])
-#endif
+#define BOX(i,n,S) (S)[(n)][(i)]
static void _des_crypt_encrypt (u32 iv[2], u32 data[2], u32 Kc[16], u32 Kd[16], __local u32 s_SPtrans[8][64])
{
static void transform_netntlmv1_key (const u32 w0, const u32 w1, u32 out[2])
{
- #ifdef VECT_SIZE1
const uchar4 t0 = as_uchar4 (w0);
const uchar4 t1 = as_uchar4 (w1);
out[0] = as_uint (k0);
out[1] = as_uint (k1);
- #endif
-
- #ifdef VECT_SIZE2
- const uchar8 t0 = as_uchar8 (w0);
- const uchar8 t1 = as_uchar8 (w1);
-
- uchar8 k0;
- uchar8 k1;
-
- k0.s0 = (t0.s0 >> 0);
- k0.s1 = (t0.s0 << 7) | (t0.s1 >> 1);
- k0.s2 = (t0.s1 << 6) | (t0.s2 >> 2);
- k0.s3 = (t0.s2 << 5) | (t0.s3 >> 3);
- k1.s0 = (t0.s3 << 4) | (t1.s0 >> 4);
- k1.s1 = (t1.s0 << 3) | (t1.s1 >> 5);
- k1.s2 = (t1.s1 << 2) | (t1.s2 >> 6);
- k1.s3 = (t1.s2 << 1);
-
- k0.s4 = (t0.s4 >> 0);
- k0.s5 = (t0.s4 << 7) | (t0.s5 >> 1);
- k0.s6 = (t0.s5 << 6) | (t0.s6 >> 2);
- k0.s7 = (t0.s6 << 5) | (t0.s7 >> 3);
- k1.s4 = (t0.s7 << 4) | (t1.s4 >> 4);
- k1.s5 = (t1.s4 << 3) | (t1.s5 >> 5);
- k1.s6 = (t1.s5 << 2) | (t1.s6 >> 6);
- k1.s7 = (t1.s6 << 1);
-
- out[0] = as_uint2 (k0);
- out[1] = as_uint2 (k1);
- #endif
-
- #ifdef VECT_SIZE4
- const uchar16 t0 = as_uchar16 (w0);
- const uchar16 t1 = as_uchar16 (w1);
-
- uchar16 k0;
- uchar16 k1;
-
- k0.s0 = (t0.s0 >> 0);
- k0.s1 = (t0.s0 << 7) | (t0.s1 >> 1);
- k0.s2 = (t0.s1 << 6) | (t0.s2 >> 2);
- k0.s3 = (t0.s2 << 5) | (t0.s3 >> 3);
- k1.s0 = (t0.s3 << 4) | (t1.s0 >> 4);
- k1.s1 = (t1.s0 << 3) | (t1.s1 >> 5);
- k1.s2 = (t1.s1 << 2) | (t1.s2 >> 6);
- k1.s3 = (t1.s2 << 1);
-
- k0.s4 = (t0.s4 >> 0);
- k0.s5 = (t0.s4 << 7) | (t0.s5 >> 1);
- k0.s6 = (t0.s5 << 6) | (t0.s6 >> 2);
- k0.s7 = (t0.s6 << 5) | (t0.s7 >> 3);
- k1.s4 = (t0.s7 << 4) | (t1.s4 >> 4);
- k1.s5 = (t1.s4 << 3) | (t1.s5 >> 5);
- k1.s6 = (t1.s5 << 2) | (t1.s6 >> 6);
- k1.s7 = (t1.s6 << 1);
-
- k0.s8 = (t0.s8 >> 0);
- k0.s9 = (t0.s8 << 7) | (t0.s9 >> 1);
- k0.sa = (t0.s9 << 6) | (t0.sa >> 2);
- k0.sb = (t0.sa << 5) | (t0.sb >> 3);
- k1.s8 = (t0.sb << 4) | (t1.s8 >> 4);
- k1.s9 = (t1.s8 << 3) | (t1.s9 >> 5);
- k1.sa = (t1.s9 << 2) | (t1.sa >> 6);
- k1.sb = (t1.sa << 1);
-
- k0.sc = (t0.sc >> 0);
- k0.sd = (t0.sc << 7) | (t0.sd >> 1);
- k0.se = (t0.sd << 6) | (t0.se >> 2);
- k0.sf = (t0.se << 5) | (t0.sf >> 3);
- k1.sc = (t0.sf << 4) | (t1.sc >> 4);
- k1.sd = (t1.sc << 3) | (t1.sd >> 5);
- k1.se = (t1.sd << 2) | (t1.se >> 6);
- k1.sf = (t1.se << 1);
-
- out[0] = as_uint4 (k0);
- out[1] = as_uint4 (k1);
- #endif
}
__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05500_m04 (__global pw_t *pws, __global gpu_rule_t * rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
__constant u32 shifts3s0[16] = { 1, 1, 2, 2, 2, 2, 2, 2, 1, 2, 2, 2, 2, 2, 2, 1 };
__constant u32 shifts3s1[16] = { 27, 27, 26, 26, 26, 26, 26, 26, 27, 26, 26, 26, 26, 26, 26, 27 };
-#ifdef VECT_SIZE1
-#define BOX(i,n,S) (u32x) ((S)[(n)][(i)])
-#endif
-
-#ifdef VECT_SIZE2
-#define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3])
-#endif
+#define BOX(i,n,S) (S)[(n)][(i)]
static void _des_crypt_encrypt (u32 iv[2], u32 data[2], u32 Kc[16], u32 Kd[16], __local u32 s_SPtrans[8][64])
{
static void transform_netntlmv1_key (const u32 w0, const u32 w1, u32 out[2])
{
- #ifdef VECT_SIZE1
const uchar4 t0 = as_uchar4 (w0);
const uchar4 t1 = as_uchar4 (w1);
out[0] = as_uint (k0);
out[1] = as_uint (k1);
- #endif
-
- #ifdef VECT_SIZE2
- const uchar8 t0 = as_uchar8 (w0);
- const uchar8 t1 = as_uchar8 (w1);
-
- uchar8 k0;
- uchar8 k1;
-
- k0.s0 = (t0.s0 >> 0);
- k0.s1 = (t0.s0 << 7) | (t0.s1 >> 1);
- k0.s2 = (t0.s1 << 6) | (t0.s2 >> 2);
- k0.s3 = (t0.s2 << 5) | (t0.s3 >> 3);
- k1.s0 = (t0.s3 << 4) | (t1.s0 >> 4);
- k1.s1 = (t1.s0 << 3) | (t1.s1 >> 5);
- k1.s2 = (t1.s1 << 2) | (t1.s2 >> 6);
- k1.s3 = (t1.s2 << 1);
-
- k0.s4 = (t0.s4 >> 0);
- k0.s5 = (t0.s4 << 7) | (t0.s5 >> 1);
- k0.s6 = (t0.s5 << 6) | (t0.s6 >> 2);
- k0.s7 = (t0.s6 << 5) | (t0.s7 >> 3);
- k1.s4 = (t0.s7 << 4) | (t1.s4 >> 4);
- k1.s5 = (t1.s4 << 3) | (t1.s5 >> 5);
- k1.s6 = (t1.s5 << 2) | (t1.s6 >> 6);
- k1.s7 = (t1.s6 << 1);
-
- out[0] = as_uint2 (k0);
- out[1] = as_uint2 (k1);
- #endif
-
- #ifdef VECT_SIZE4
- const uchar16 t0 = as_uchar16 (w0);
- const uchar16 t1 = as_uchar16 (w1);
-
- uchar16 k0;
- uchar16 k1;
-
- k0.s0 = (t0.s0 >> 0);
- k0.s1 = (t0.s0 << 7) | (t0.s1 >> 1);
- k0.s2 = (t0.s1 << 6) | (t0.s2 >> 2);
- k0.s3 = (t0.s2 << 5) | (t0.s3 >> 3);
- k1.s0 = (t0.s3 << 4) | (t1.s0 >> 4);
- k1.s1 = (t1.s0 << 3) | (t1.s1 >> 5);
- k1.s2 = (t1.s1 << 2) | (t1.s2 >> 6);
- k1.s3 = (t1.s2 << 1);
-
- k0.s4 = (t0.s4 >> 0);
- k0.s5 = (t0.s4 << 7) | (t0.s5 >> 1);
- k0.s6 = (t0.s5 << 6) | (t0.s6 >> 2);
- k0.s7 = (t0.s6 << 5) | (t0.s7 >> 3);
- k1.s4 = (t0.s7 << 4) | (t1.s4 >> 4);
- k1.s5 = (t1.s4 << 3) | (t1.s5 >> 5);
- k1.s6 = (t1.s5 << 2) | (t1.s6 >> 6);
- k1.s7 = (t1.s6 << 1);
-
- k0.s8 = (t0.s8 >> 0);
- k0.s9 = (t0.s8 << 7) | (t0.s9 >> 1);
- k0.sa = (t0.s9 << 6) | (t0.sa >> 2);
- k0.sb = (t0.sa << 5) | (t0.sb >> 3);
- k1.s8 = (t0.sb << 4) | (t1.s8 >> 4);
- k1.s9 = (t1.s8 << 3) | (t1.s9 >> 5);
- k1.sa = (t1.s9 << 2) | (t1.sa >> 6);
- k1.sb = (t1.sa << 1);
-
- k0.sc = (t0.sc >> 0);
- k0.sd = (t0.sc << 7) | (t0.sd >> 1);
- k0.se = (t0.sd << 6) | (t0.se >> 2);
- k0.sf = (t0.se << 5) | (t0.sf >> 3);
- k1.sc = (t0.sf << 4) | (t1.sc >> 4);
- k1.sd = (t1.sc << 3) | (t1.sd >> 5);
- k1.se = (t1.sd << 2) | (t1.se >> 6);
- k1.sf = (t1.se << 1);
-
- out[0] = as_uint4 (k0);
- out[1] = as_uint4 (k1);
- #endif
}
__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05500_m04 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
__constant u32 shifts3s0[16] = { 1, 1, 2, 2, 2, 2, 2, 2, 1, 2, 2, 2, 2, 2, 2, 1 };
__constant u32 shifts3s1[16] = { 27, 27, 26, 26, 26, 26, 26, 26, 27, 26, 26, 26, 26, 26, 26, 27 };
-#ifdef VECT_SIZE1
-#define BOX(i,n,S) (u32x) ((S)[(n)][(i)])
-#endif
-
-#ifdef VECT_SIZE2
-#define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3])
-#endif
+#define BOX(i,n,S) (S)[(n)][(i)]
static void _des_crypt_encrypt (u32 iv[2], u32 data[2], u32 Kc[16], u32 Kd[16], __local u32 s_SPtrans[8][64])
{
static void transform_netntlmv1_key (const u32 w0, const u32 w1, u32 out[2])
{
- #ifdef VECT_SIZE1
const uchar4 t0 = as_uchar4 (w0);
const uchar4 t1 = as_uchar4 (w1);
out[0] = as_uint (k0);
out[1] = as_uint (k1);
- #endif
-
- #ifdef VECT_SIZE2
- const uchar8 t0 = as_uchar8 (w0);
- const uchar8 t1 = as_uchar8 (w1);
-
- uchar8 k0;
- uchar8 k1;
-
- k0.s0 = (t0.s0 >> 0);
- k0.s1 = (t0.s0 << 7) | (t0.s1 >> 1);
- k0.s2 = (t0.s1 << 6) | (t0.s2 >> 2);
- k0.s3 = (t0.s2 << 5) | (t0.s3 >> 3);
- k1.s0 = (t0.s3 << 4) | (t1.s0 >> 4);
- k1.s1 = (t1.s0 << 3) | (t1.s1 >> 5);
- k1.s2 = (t1.s1 << 2) | (t1.s2 >> 6);
- k1.s3 = (t1.s2 << 1);
-
- k0.s4 = (t0.s4 >> 0);
- k0.s5 = (t0.s4 << 7) | (t0.s5 >> 1);
- k0.s6 = (t0.s5 << 6) | (t0.s6 >> 2);
- k0.s7 = (t0.s6 << 5) | (t0.s7 >> 3);
- k1.s4 = (t0.s7 << 4) | (t1.s4 >> 4);
- k1.s5 = (t1.s4 << 3) | (t1.s5 >> 5);
- k1.s6 = (t1.s5 << 2) | (t1.s6 >> 6);
- k1.s7 = (t1.s6 << 1);
-
- out[0] = as_uint2 (k0);
- out[1] = as_uint2 (k1);
- #endif
-
- #ifdef VECT_SIZE4
- const uchar16 t0 = as_uchar16 (w0);
- const uchar16 t1 = as_uchar16 (w1);
-
- uchar16 k0;
- uchar16 k1;
-
- k0.s0 = (t0.s0 >> 0);
- k0.s1 = (t0.s0 << 7) | (t0.s1 >> 1);
- k0.s2 = (t0.s1 << 6) | (t0.s2 >> 2);
- k0.s3 = (t0.s2 << 5) | (t0.s3 >> 3);
- k1.s0 = (t0.s3 << 4) | (t1.s0 >> 4);
- k1.s1 = (t1.s0 << 3) | (t1.s1 >> 5);
- k1.s2 = (t1.s1 << 2) | (t1.s2 >> 6);
- k1.s3 = (t1.s2 << 1);
-
- k0.s4 = (t0.s4 >> 0);
- k0.s5 = (t0.s4 << 7) | (t0.s5 >> 1);
- k0.s6 = (t0.s5 << 6) | (t0.s6 >> 2);
- k0.s7 = (t0.s6 << 5) | (t0.s7 >> 3);
- k1.s4 = (t0.s7 << 4) | (t1.s4 >> 4);
- k1.s5 = (t1.s4 << 3) | (t1.s5 >> 5);
- k1.s6 = (t1.s5 << 2) | (t1.s6 >> 6);
- k1.s7 = (t1.s6 << 1);
-
- k0.s8 = (t0.s8 >> 0);
- k0.s9 = (t0.s8 << 7) | (t0.s9 >> 1);
- k0.sa = (t0.s9 << 6) | (t0.sa >> 2);
- k0.sb = (t0.sa << 5) | (t0.sb >> 3);
- k1.s8 = (t0.sb << 4) | (t1.s8 >> 4);
- k1.s9 = (t1.s8 << 3) | (t1.s9 >> 5);
- k1.sa = (t1.s9 << 2) | (t1.sa >> 6);
- k1.sb = (t1.sa << 1);
-
- k0.sc = (t0.sc >> 0);
- k0.sd = (t0.sc << 7) | (t0.sd >> 1);
- k0.se = (t0.sd << 6) | (t0.se >> 2);
- k0.sf = (t0.se << 5) | (t0.sf >> 3);
- k1.sc = (t0.sf << 4) | (t1.sc >> 4);
- k1.sd = (t1.sc << 3) | (t1.sd >> 5);
- k1.se = (t1.sd << 2) | (t1.se >> 6);
- k1.sf = (t1.se << 1);
-
- out[0] = as_uint4 (k0);
- out[1] = as_uint4 (k1);
- #endif
}
static void m05500m (__local u32 s_SPtrans[8][64], __local u32 s_skb[8][64], u32 w[16], const u32 pw_len, __global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global u32 * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset)
#define R 10
-#ifdef VECT_SIZE1
-#define BOX(S,n,i) u32 ((S)[(n)][(i)])
-#endif
-
-#ifdef VECT_SIZE2
-#define BOX(S,n,i) u32 ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
-#endif
+#define BOX(S,n,i) (S)[(n)][(i)]
__constant u32 Ch[8][256] =
{
#define R 10
-#ifdef VECT_SIZE1
-#define BOX(S,n,i) u32 ((S)[(n)][(i)])
-#endif
-
-#ifdef VECT_SIZE2
-#define BOX(S,n,i) u32 ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
-#endif
+#define BOX(S,n,i) (S)[(n)][(i)]
__constant u32 Ch[8][256] =
{
#define R 10
-#ifdef VECT_SIZE1
-#define BOX(S,n,i) u32 ((S)[(n)][(i)])
-#endif
-
-#ifdef VECT_SIZE2
-#define BOX(S,n,i) u32 ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
-#endif
+#define BOX(S,n,i) (S)[(n)][(i)]
__constant u32 Ch[8][256] =
{
},
};
-#ifdef VECT_SIZE1
-#define BOX(S,n,i) (u32) ((S)[(n)][(i)])
-#endif
+#define BOX(S,n,i) (S)[(n)][(i)]
static void whirlpool_transform (const u32 w[16], u32 dgst[16], __local u32 s_Ch[8][256], __local u32 s_Cl[8][256])
{
},
};
-#ifdef VECT_SIZE1
-#define BOX(S,n,i) (u32) ((S)[(n)][(i)])
-#endif
+#define BOX(S,n,i) (S)[(n)][(i)]
static void whirlpool_transform (const u32 w[16], u32 dgst[16], __local u32 s_Ch[8][256], __local u32 s_Cl[8][256])
{
},
};
-#ifdef VECT_SIZE1
-#define BOX(S,n,i) (u32) ((S)[(n)][(i)])
-#endif
+#define BOX(S,n,i) (S)[(n)][(i)]
static void whirlpool_transform (const u32 w[16], u32 dgst[16], __local u32 s_Ch[8][256], __local u32 s_Cl[8][256])
{
u64 w2l[4];
u64 w3l[4];
- w0l[0] = (u64x) (w0[0]) << 32 | (u64x) (w0[1]);
- w0l[1] = (u64x) (w0[2]) << 32 | (u64x) (w0[3]);
- w0l[2] = (u64x) (w1[0]) << 32 | (u64x) (w1[1]);
- w0l[3] = (u64x) (w1[2]) << 32 | (u64x) (w1[3]);
- w1l[0] = (u64x) (w2[0]) << 32 | (u64x) (w2[1]);
- w1l[1] = (u64x) (w2[2]) << 32 | (u64x) (w2[3]);
- w1l[2] = (u64x) (w3[0]) << 32 | (u64x) (w3[1]);
- w1l[3] = (u64x) (w3[2]) << 32 | (u64x) (w3[3]);
+ w0l[0] = (u64) (w0[0]) << 32 | (u64) (w0[1]);
+ w0l[1] = (u64) (w0[2]) << 32 | (u64) (w0[3]);
+ w0l[2] = (u64) (w1[0]) << 32 | (u64) (w1[1]);
+ w0l[3] = (u64) (w1[2]) << 32 | (u64) (w1[3]);
+ w1l[0] = (u64) (w2[0]) << 32 | (u64) (w2[1]);
+ w1l[1] = (u64) (w2[2]) << 32 | (u64) (w2[3]);
+ w1l[2] = (u64) (w3[0]) << 32 | (u64) (w3[1]);
+ w1l[3] = (u64) (w3[2]) << 32 | (u64) (w3[3]);
w2l[0] = 0;
w2l[1] = 0;
w2l[2] = 0;
tmps[gid].opad[6] = opad[6];
tmps[gid].opad[7] = opad[7];
- w0l[0] = (u64x) salt_buf0[1] << 32 | (u64x) salt_buf0[0];
- w0l[1] = (u64x) salt_buf0[3] << 32 | (u64x) salt_buf0[2];
- w0l[2] = (u64x) salt_buf1[1] << 32 | (u64x) salt_buf1[0];
- w0l[3] = (u64x) salt_buf1[3] << 32 | (u64x) salt_buf1[2];
- w1l[0] = (u64x) salt_buf2[1] << 32 | (u64x) salt_buf2[0];
- w1l[1] = (u64x) salt_buf2[3] << 32 | (u64x) salt_buf2[2];
- w1l[2] = (u64x) salt_buf3[1] << 32 | (u64x) salt_buf3[0];
- w1l[3] = (u64x) salt_buf3[3] << 32 | (u64x) salt_buf3[2];
+ w0l[0] = (u64) salt_buf0[1] << 32 | (u64) salt_buf0[0];
+ w0l[1] = (u64) salt_buf0[3] << 32 | (u64) salt_buf0[2];
+ w0l[2] = (u64) salt_buf1[1] << 32 | (u64) salt_buf1[0];
+ w0l[3] = (u64) salt_buf1[3] << 32 | (u64) salt_buf1[2];
+ w1l[0] = (u64) salt_buf2[1] << 32 | (u64) salt_buf2[0];
+ w1l[1] = (u64) salt_buf2[3] << 32 | (u64) salt_buf2[2];
+ w1l[2] = (u64) salt_buf3[1] << 32 | (u64) salt_buf3[0];
+ w1l[3] = (u64) salt_buf3[3] << 32 | (u64) salt_buf3[2];
w2l[0] = 0;
w2l[1] = 0;
w2l[2] = 0;
#include "types_ocl.c"
#include "common.c"
-#ifdef VECT_SIZE1
-#define COMPARE_M "check_multi_vect1_comp4.c"
-#endif
-
-#ifdef VECT_SIZE2
-#define COMPARE_M "check_multi_vect2_comp4.c"
-#endif
+#define COMPARE_S "check_single_comp4.c"
+#define COMPARE_M "check_multi_comp4.c"
__constant u32 te0[256] =
{
#include "types_ocl.c"
#include "common.c"
-#ifdef VECT_SIZE1
-#define COMPARE_M "check_multi_vect1_comp4.c"
-#endif
-
-#ifdef VECT_SIZE2
-#define COMPARE_M "check_multi_vect2_comp4.c"
-#endif
+#define COMPARE_S "check_single_comp4.c"
+#define COMPARE_M "check_multi_comp4.c"
__constant u32 te0[256] =
{
}
};
-#ifdef VECT_SIZE1
-#define BOX(i,n,S) (u32x) ((S)[(n)][(i)])
-#endif
-
-#ifdef VECT_SIZE2
-#define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
-#endif
+#define BOX(i,n,S) (S)[(n)][(i)]
#define round(k1,k2,tbl) \
{ \
- u32 t; \
+ u32 t; \
t = (k1) + r; \
l ^= BOX (amd_bfe (t, 0, 8), 0, tbl) ^ \
BOX (amd_bfe (t, 8, 8), 1, tbl) ^ \
#define R(k,h,s,i,t) \
{ \
- u32 r; \
- u32 l; \
+ u32 r; \
+ u32 l; \
r = h[i + 0]; \
l = h[i + 1]; \
round (k[0], k[1], t); \
}
};
-#ifdef VECT_SIZE1
-#define BOX(i,n,S) (u32x) ((S)[(n)][(i)])
-#endif
-
-#ifdef VECT_SIZE2
-#define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
-#endif
+#define BOX(i,n,S) (S)[(n)][(i)]
#define round(k1,k2,tbl) \
{ \
- u32 t; \
+ u32 t; \
t = (k1) + r; \
l ^= BOX (amd_bfe (t, 0, 8), 0, tbl) ^ \
BOX (amd_bfe (t, 8, 8), 1, tbl) ^ \
#define R(k,h,s,i,t) \
{ \
- u32 r; \
- u32 l; \
+ u32 r; \
+ u32 l; \
r = h[i + 0]; \
l = h[i + 1]; \
round (k[0], k[1], t); \
}
};
-#ifdef VECT_SIZE1
-#define BOX(i,n,S) (u32x) ((S)[(n)][(i)])
-#endif
-
-#ifdef VECT_SIZE2
-#define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
-#endif
+#define BOX(i,n,S) (S)[(n)][(i)]
#define round(k1,k2,tbl) \
{ \
- u32 t; \
+ u32 t; \
t = (k1) + r; \
l ^= BOX (amd_bfe (t, 0, 8), 0, tbl) ^ \
BOX (amd_bfe (t, 8, 8), 1, tbl) ^ \
#define R(k,h,s,i,t) \
{ \
- u32 r; \
- u32 l; \
+ u32 r; \
+ u32 l; \
r = h[i + 0]; \
l = h[i + 1]; \
round (k[0], k[1], t); \
u32 tmp[4];
- #ifdef VECT_SIZE1
-
tmp[0] = digest[0];
tmp[1] = digest[1];
tmp[2] = digest[2];
d_return_buf[lid] = 1;
}
-
- #endif
}
}
u32 tmp[4];
- #ifdef VECT_SIZE1
-
tmp[0] = digest[0];
tmp[1] = digest[1];
tmp[2] = digest[2];
d_return_buf[lid] = 1;
}
-
- #endif
}
}
u32 tmp[4];
- #ifdef VECT_SIZE1
-
tmp[0] = digest[0];
tmp[1] = digest[1];
tmp[2] = digest[2];
d_return_buf[lid] = 1;
}
-
- #endif
}
}
u32 tmp[4];
- #ifdef VECT_SIZE1
-
tmp[0] = digest[0];
tmp[1] = digest[1];
tmp[2] = digest[2];
d_return_buf[lid] = 1;
}
-
- #endif
}
}
u32 tmp[4];
- #ifdef VECT_SIZE1
-
tmp[0] = digest[0];
tmp[1] = digest[1];
tmp[2] = digest[2];
d_return_buf[lid] = 1;
}
-
- #endif
}
}
#define COMPARE_S "check_single_comp4.c"
#define COMPARE_M "check_multi_comp4.c"
-#ifdef VECT_SIZE1
#define uint_to_hex_lower8(i) l_bin2asc[(i)]
-#endif
-
-#ifdef VECT_SIZE2
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3])
-#endif
__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m07600_m04 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
#define COMPARE_S "check_single_comp4.c"
#define COMPARE_M "check_multi_comp4.c"
-#ifdef VECT_SIZE1
#define uint_to_hex_lower8(i) l_bin2asc[(i)]
-#endif
-
-#ifdef VECT_SIZE2
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3])
-#endif
__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m07600_m04 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
#define COMPARE_S "check_single_comp4.c"
#define COMPARE_M "check_multi_comp4.c"
-#ifdef VECT_SIZE1
#define uint_to_hex_lower8(i) l_bin2asc[(i)]
-#endif
-
-#ifdef VECT_SIZE2
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3])
-#endif
static void m07600m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, __local u32 l_bin2asc[256])
{
{
u32 out = 0;
- #ifdef VECT_SIZE1
out |= (sapb_trans_tbl[(in >> 0) & 0xff]) << 0;
out |= (sapb_trans_tbl[(in >> 8) & 0xff]) << 8;
out |= (sapb_trans_tbl[(in >> 16) & 0xff]) << 16;
out |= (sapb_trans_tbl[(in >> 24) & 0xff]) << 24;
- #endif
return out;
}
{
u32 out = 0;
- #ifdef VECT_SIZE1
out |= (sapb_trans_tbl[(in >> 0) & 0xff]) << 0;
out |= (sapb_trans_tbl[(in >> 8) & 0xff]) << 8;
out |= (sapb_trans_tbl[(in >> 16) & 0xff]) << 16;
out |= (sapb_trans_tbl[(in >> 24) & 0xff]) << 24;
- #endif
return out;
}
{
u32 out = 0;
- #ifdef VECT_SIZE1
out |= (sapb_trans_tbl[(in >> 0) & 0xff]) << 0;
out |= (sapb_trans_tbl[(in >> 8) & 0xff]) << 8;
out |= (sapb_trans_tbl[(in >> 16) & 0xff]) << 16;
out |= (sapb_trans_tbl[(in >> 24) & 0xff]) << 24;
- #endif
return out;
}
#include "types_ocl.c"
#include "common.c"
-#ifdef VECT_SIZE1
-#define COMPARE_M "check_multi_vect1_comp4.c"
-#endif
+#define COMPARE_S "check_single_comp4.c"
+#define COMPARE_M "check_multi_comp4.c"
__constant u64 k_sha512[80] =
{
#include "types_ocl.c"
#include "common.c"
-#ifdef VECT_SIZE1
-#define COMPARE_M "check_multi_vect1_comp4.c"
-#endif
+#define COMPARE_S "check_single_comp4.c"
+#define COMPARE_M "check_multi_comp4.c"
__constant u32 k_sha256[64] =
{
#define COMPARE_S "check_single_comp4.c"
#define COMPARE_M "check_multi_comp4.c"
-#ifdef VECT_SIZE1
#define uint_to_hex_lower8_le(i) l_bin2asc[(i)]
-#endif
-
-#ifdef VECT_SIZE2
-#define uint_to_hex_lower8_le(i) u32 (l_bin2asc[(i).s0], l_bin2asc[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define uint_to_hex_lower8_le(i) u32 (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3])
-#endif
static void sha1_transform (const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], u32 digest[5])
{
#define COMPARE_S "check_single_comp4.c"
#define COMPARE_M "check_multi_comp4.c"
-#ifdef VECT_SIZE1
#define uint_to_hex_lower8_le(i) l_bin2asc[(i)]
-#endif
-
-#ifdef VECT_SIZE2
-#define uint_to_hex_lower8_le(i) u32 (l_bin2asc[(i).s0], l_bin2asc[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define uint_to_hex_lower8_le(i) u32 (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3])
-#endif
static void sha1_transform (const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], u32 digest[5])
{
#define COMPARE_S "check_single_comp4.c"
#define COMPARE_M "check_multi_comp4.c"
-#ifdef VECT_SIZE1
#define uint_to_hex_lower8_le(i) l_bin2asc[(i)]
-#endif
-
-#ifdef VECT_SIZE2
-#define uint_to_hex_lower8_le(i) u32 (l_bin2asc[(i).s0], l_bin2asc[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define uint_to_hex_lower8_le(i) u32 (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3])
-#endif
static void sha1_transform (const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], u32 digest[5])
{
u32 u = Kc[i] ^ r;
u32 t = Kd[i] ^ rotl32 (r, 28u);
- #ifdef VECT_SIZE1
- l ^= NBOX (((u >> 2) & 0x3f), 0, s_SPtrans)
- | NBOX (((u >> 10) & 0x3f), 2, s_SPtrans)
- | NBOX (((u >> 18) & 0x3f), 4, s_SPtrans)
- | NBOX (((u >> 26) & 0x3f), 6, s_SPtrans)
- | NBOX (((t >> 2) & 0x3f), 1, s_SPtrans)
- | NBOX (((t >> 10) & 0x3f), 3, s_SPtrans)
- | NBOX (((t >> 18) & 0x3f), 5, s_SPtrans)
- | NBOX (((t >> 26) & 0x3f), 7, s_SPtrans);
- #endif
-
- #ifdef VECT_SIZE2
- l.s0 ^= NBOX (((u.s0 >> 2) & 0x3f), 0, s_SPtrans)
- | NBOX (((u.s0 >> 10) & 0x3f), 2, s_SPtrans)
- | NBOX (((u.s0 >> 18) & 0x3f), 4, s_SPtrans)
- | NBOX (((u.s0 >> 26) & 0x3f), 6, s_SPtrans)
- | NBOX (((t.s0 >> 2) & 0x3f), 1, s_SPtrans)
- | NBOX (((t.s0 >> 10) & 0x3f), 3, s_SPtrans)
- | NBOX (((t.s0 >> 18) & 0x3f), 5, s_SPtrans)
- | NBOX (((t.s0 >> 26) & 0x3f), 7, s_SPtrans);
-
- l.s1 ^= NBOX (((u.s1 >> 2) & 0x3f), 0, s_SPtrans)
- | NBOX (((u.s1 >> 10) & 0x3f), 2, s_SPtrans)
- | NBOX (((u.s1 >> 18) & 0x3f), 4, s_SPtrans)
- | NBOX (((u.s1 >> 26) & 0x3f), 6, s_SPtrans)
- | NBOX (((t.s1 >> 2) & 0x3f), 1, s_SPtrans)
- | NBOX (((t.s1 >> 10) & 0x3f), 3, s_SPtrans)
- | NBOX (((t.s1 >> 18) & 0x3f), 5, s_SPtrans)
- | NBOX (((t.s1 >> 26) & 0x3f), 7, s_SPtrans);
- #endif
-
- #ifdef VECT_SIZE4
- l.s0 ^= NBOX (((u.s0 >> 2) & 0x3f), 0, s_SPtrans)
- | NBOX (((u.s0 >> 10) & 0x3f), 2, s_SPtrans)
- | NBOX (((u.s0 >> 18) & 0x3f), 4, s_SPtrans)
- | NBOX (((u.s0 >> 26) & 0x3f), 6, s_SPtrans)
- | NBOX (((t.s0 >> 2) & 0x3f), 1, s_SPtrans)
- | NBOX (((t.s0 >> 10) & 0x3f), 3, s_SPtrans)
- | NBOX (((t.s0 >> 18) & 0x3f), 5, s_SPtrans)
- | NBOX (((t.s0 >> 26) & 0x3f), 7, s_SPtrans);
-
- l.s1 ^= NBOX (((u.s1 >> 2) & 0x3f), 0, s_SPtrans)
- | NBOX (((u.s1 >> 10) & 0x3f), 2, s_SPtrans)
- | NBOX (((u.s1 >> 18) & 0x3f), 4, s_SPtrans)
- | NBOX (((u.s1 >> 26) & 0x3f), 6, s_SPtrans)
- | NBOX (((t.s1 >> 2) & 0x3f), 1, s_SPtrans)
- | NBOX (((t.s1 >> 10) & 0x3f), 3, s_SPtrans)
- | NBOX (((t.s1 >> 18) & 0x3f), 5, s_SPtrans)
- | NBOX (((t.s1 >> 26) & 0x3f), 7, s_SPtrans);
-
- l.s2 ^= NBOX (((u.s2 >> 2) & 0x3f), 0, s_SPtrans)
- | NBOX (((u.s2 >> 10) & 0x3f), 2, s_SPtrans)
- | NBOX (((u.s2 >> 18) & 0x3f), 4, s_SPtrans)
- | NBOX (((u.s2 >> 26) & 0x3f), 6, s_SPtrans)
- | NBOX (((t.s2 >> 2) & 0x3f), 1, s_SPtrans)
- | NBOX (((t.s2 >> 10) & 0x3f), 3, s_SPtrans)
- | NBOX (((t.s2 >> 18) & 0x3f), 5, s_SPtrans)
- | NBOX (((t.s2 >> 26) & 0x3f), 7, s_SPtrans);
-
- l.s3 ^= NBOX (((u.s3 >> 2) & 0x3f), 0, s_SPtrans)
- | NBOX (((u.s3 >> 10) & 0x3f), 2, s_SPtrans)
- | NBOX (((u.s3 >> 18) & 0x3f), 4, s_SPtrans)
- | NBOX (((u.s3 >> 26) & 0x3f), 6, s_SPtrans)
- | NBOX (((t.s3 >> 2) & 0x3f), 1, s_SPtrans)
- | NBOX (((t.s3 >> 10) & 0x3f), 3, s_SPtrans)
- | NBOX (((t.s3 >> 18) & 0x3f), 5, s_SPtrans)
- | NBOX (((t.s3 >> 26) & 0x3f), 7, s_SPtrans);
- #endif
+ l ^= NBOX (((u >> 2) & 0x3f), 0, s_SPtrans)
+ | NBOX (((u >> 10) & 0x3f), 2, s_SPtrans)
+ | NBOX (((u >> 18) & 0x3f), 4, s_SPtrans)
+ | NBOX (((u >> 26) & 0x3f), 6, s_SPtrans)
+ | NBOX (((t >> 2) & 0x3f), 1, s_SPtrans)
+ | NBOX (((t >> 10) & 0x3f), 3, s_SPtrans)
+ | NBOX (((t >> 18) & 0x3f), 5, s_SPtrans)
+ | NBOX (((t >> 26) & 0x3f), 7, s_SPtrans);
tt = l;
l = r;
u32 s;
u32 t;
- #ifdef VECT_SIZE1
- s = NBOX ((( c >> 0) & 0x3f), 0, s_skb)
- | NBOX ((((c >> 6) & 0x03)
- | ((c >> 7) & 0x3c)), 1, s_skb)
- | NBOX ((((c >> 13) & 0x0f)
- | ((c >> 14) & 0x30)), 2, s_skb)
- | NBOX ((((c >> 20) & 0x01)
- | ((c >> 21) & 0x06)
- | ((c >> 22) & 0x38)), 3, s_skb);
-
- t = NBOX ((( d >> 0) & 0x3f), 4, s_skb)
- | NBOX ((((d >> 7) & 0x03)
- | ((d >> 8) & 0x3c)), 5, s_skb)
- | NBOX ((((d >> 15) & 0x3f)), 6, s_skb)
- | NBOX ((((d >> 21) & 0x0f)
- | ((d >> 22) & 0x30)), 7, s_skb);
- #endif
-
- #ifdef VECT_SIZE2
- s.s0 = NBOX ((( c.s0 >> 0) & 0x3f), 0, s_skb)
- | NBOX ((((c.s0 >> 6) & 0x03)
- | ((c.s0 >> 7) & 0x3c)), 1, s_skb)
- | NBOX ((((c.s0 >> 13) & 0x0f)
- | ((c.s0 >> 14) & 0x30)), 2, s_skb)
- | NBOX ((((c.s0 >> 20) & 0x01)
- | ((c.s0 >> 21) & 0x06)
- | ((c.s0 >> 22) & 0x38)), 3, s_skb);
-
- t.s0 = NBOX ((( d.s0 >> 0) & 0x3f), 4, s_skb)
- | NBOX ((((d.s0 >> 7) & 0x03)
- | ((d.s0 >> 8) & 0x3c)), 5, s_skb)
- | NBOX ((((d.s0 >> 15) & 0x3f)), 6, s_skb)
- | NBOX ((((d.s0 >> 21) & 0x0f)
- | ((d.s0 >> 22) & 0x30)), 7, s_skb);
-
- s.s1 = NBOX ((( c.s1 >> 0) & 0x3f), 0, s_skb)
- | NBOX ((((c.s1 >> 6) & 0x03)
- | ((c.s1 >> 7) & 0x3c)), 1, s_skb)
- | NBOX ((((c.s1 >> 13) & 0x0f)
- | ((c.s1 >> 14) & 0x30)), 2, s_skb)
- | NBOX ((((c.s1 >> 20) & 0x01)
- | ((c.s1 >> 21) & 0x06)
- | ((c.s1 >> 22) & 0x38)), 3, s_skb);
-
- t.s1 = NBOX ((( d.s1 >> 0) & 0x3f), 4, s_skb)
- | NBOX ((((d.s1 >> 7) & 0x03)
- | ((d.s1 >> 8) & 0x3c)), 5, s_skb)
- | NBOX ((((d.s1 >> 15) & 0x3f)), 6, s_skb)
- | NBOX ((((d.s1 >> 21) & 0x0f)
- | ((d.s1 >> 22) & 0x30)), 7, s_skb);
- #endif
-
- #ifdef VECT_SIZE4
- s.s0 = NBOX ((( c.s0 >> 0) & 0x3f), 0, s_skb)
- | NBOX ((((c.s0 >> 6) & 0x03)
- | ((c.s0 >> 7) & 0x3c)), 1, s_skb)
- | NBOX ((((c.s0 >> 13) & 0x0f)
- | ((c.s0 >> 14) & 0x30)), 2, s_skb)
- | NBOX ((((c.s0 >> 20) & 0x01)
- | ((c.s0 >> 21) & 0x06)
- | ((c.s0 >> 22) & 0x38)), 3, s_skb);
-
- t.s0 = NBOX ((( d.s0 >> 0) & 0x3f), 4, s_skb)
- | NBOX ((((d.s0 >> 7) & 0x03)
- | ((d.s0 >> 8) & 0x3c)), 5, s_skb)
- | NBOX ((((d.s0 >> 15) & 0x3f)), 6, s_skb)
- | NBOX ((((d.s0 >> 21) & 0x0f)
- | ((d.s0 >> 22) & 0x30)), 7, s_skb);
-
- s.s1 = NBOX ((( c.s1 >> 0) & 0x3f), 0, s_skb)
- | NBOX ((((c.s1 >> 6) & 0x03)
- | ((c.s1 >> 7) & 0x3c)), 1, s_skb)
- | NBOX ((((c.s1 >> 13) & 0x0f)
- | ((c.s1 >> 14) & 0x30)), 2, s_skb)
- | NBOX ((((c.s1 >> 20) & 0x01)
- | ((c.s1 >> 21) & 0x06)
- | ((c.s1 >> 22) & 0x38)), 3, s_skb);
-
- t.s1 = NBOX ((( d.s1 >> 0) & 0x3f), 4, s_skb)
- | NBOX ((((d.s1 >> 7) & 0x03)
- | ((d.s1 >> 8) & 0x3c)), 5, s_skb)
- | NBOX ((((d.s1 >> 15) & 0x3f)), 6, s_skb)
- | NBOX ((((d.s1 >> 21) & 0x0f)
- | ((d.s1 >> 22) & 0x30)), 7, s_skb);
-
- s.s2 = NBOX ((( c.s2 >> 0) & 0x3f), 0, s_skb)
- | NBOX ((((c.s2 >> 6) & 0x03)
- | ((c.s2 >> 7) & 0x3c)), 1, s_skb)
- | NBOX ((((c.s2 >> 13) & 0x0f)
- | ((c.s2 >> 14) & 0x30)), 2, s_skb)
- | NBOX ((((c.s2 >> 20) & 0x01)
- | ((c.s2 >> 21) & 0x06)
- | ((c.s2 >> 22) & 0x38)), 3, s_skb);
-
- t.s2 = NBOX ((( d.s2 >> 0) & 0x3f), 4, s_skb)
- | NBOX ((((d.s2 >> 7) & 0x03)
- | ((d.s2 >> 8) & 0x3c)), 5, s_skb)
- | NBOX ((((d.s2 >> 15) & 0x3f)), 6, s_skb)
- | NBOX ((((d.s2 >> 21) & 0x0f)
- | ((d.s2 >> 22) & 0x30)), 7, s_skb);
-
- s.s3 = NBOX ((( c.s3 >> 0) & 0x3f), 0, s_skb)
- | NBOX ((((c.s3 >> 6) & 0x03)
- | ((c.s3 >> 7) & 0x3c)), 1, s_skb)
- | NBOX ((((c.s3 >> 13) & 0x0f)
- | ((c.s3 >> 14) & 0x30)), 2, s_skb)
- | NBOX ((((c.s3 >> 20) & 0x01)
- | ((c.s3 >> 21) & 0x06)
- | ((c.s3 >> 22) & 0x38)), 3, s_skb);
-
- t.s3 = NBOX ((( d.s3 >> 0) & 0x3f), 4, s_skb)
- | NBOX ((((d.s3 >> 7) & 0x03)
- | ((d.s3 >> 8) & 0x3c)), 5, s_skb)
- | NBOX ((((d.s3 >> 15) & 0x3f)), 6, s_skb)
- | NBOX ((((d.s3 >> 21) & 0x0f)
- | ((d.s3 >> 22) & 0x30)), 7, s_skb);
- #endif
+ s = NBOX ((( c >> 0) & 0x3f), 0, s_skb)
+ | NBOX ((((c >> 6) & 0x03)
+ | ((c >> 7) & 0x3c)), 1, s_skb)
+ | NBOX ((((c >> 13) & 0x0f)
+ | ((c >> 14) & 0x30)), 2, s_skb)
+ | NBOX ((((c >> 20) & 0x01)
+ | ((c >> 21) & 0x06)
+ | ((c >> 22) & 0x38)), 3, s_skb);
+
+ t = NBOX ((( d >> 0) & 0x3f), 4, s_skb)
+ | NBOX ((((d >> 7) & 0x03)
+ | ((d >> 8) & 0x3c)), 5, s_skb)
+ | NBOX ((((d >> 15) & 0x3f)), 6, s_skb)
+ | NBOX ((((d >> 21) & 0x0f)
+ | ((d >> 22) & 0x30)), 7, s_skb);
#if defined cl_amd_media_ops
Kc[i] = amd_bytealign (t, s << 16, 2);
static void transform_racf_key (const u32 w0, const u32 w1, u32 key[2])
{
- #ifdef VECT_SIZE1
-
- key[0] = (ascii_to_ebcdic_pc[(w0 >> 0) & 0xff]) << 0
- | (ascii_to_ebcdic_pc[(w0 >> 8) & 0xff]) << 8
- | (ascii_to_ebcdic_pc[(w0 >> 16) & 0xff]) << 16
- | (ascii_to_ebcdic_pc[(w0 >> 24) & 0xff]) << 24;
-
- key[1] = (ascii_to_ebcdic_pc[(w1 >> 0) & 0xff]) << 0
- | (ascii_to_ebcdic_pc[(w1 >> 8) & 0xff]) << 8
- | (ascii_to_ebcdic_pc[(w1 >> 16) & 0xff]) << 16
- | (ascii_to_ebcdic_pc[(w1 >> 24) & 0xff]) << 24;
- #endif
-
- #ifdef VECT_SIZE2
-
- key[0].s0 = (ascii_to_ebcdic_pc[(w0.s0 >> 0) & 0xff]) << 0
- | (ascii_to_ebcdic_pc[(w0.s0 >> 8) & 0xff]) << 8
- | (ascii_to_ebcdic_pc[(w0.s0 >> 16) & 0xff]) << 16
- | (ascii_to_ebcdic_pc[(w0.s0 >> 24) & 0xff]) << 24;
-
- key[0].s1 = (ascii_to_ebcdic_pc[(w0.s1 >> 0) & 0xff]) << 0
- | (ascii_to_ebcdic_pc[(w0.s1 >> 8) & 0xff]) << 8
- | (ascii_to_ebcdic_pc[(w0.s1 >> 16) & 0xff]) << 16
- | (ascii_to_ebcdic_pc[(w0.s1 >> 24) & 0xff]) << 24;
-
- key[1].s0 = (ascii_to_ebcdic_pc[(w1.s0 >> 0) & 0xff]) << 0
- | (ascii_to_ebcdic_pc[(w1.s0 >> 8) & 0xff]) << 8
- | (ascii_to_ebcdic_pc[(w1.s0 >> 16) & 0xff]) << 16
- | (ascii_to_ebcdic_pc[(w1.s0 >> 24) & 0xff]) << 24;
-
- key[1].s1 = (ascii_to_ebcdic_pc[(w1.s1 >> 0) & 0xff]) << 0
- | (ascii_to_ebcdic_pc[(w1.s1 >> 8) & 0xff]) << 8
- | (ascii_to_ebcdic_pc[(w1.s1 >> 16) & 0xff]) << 16
- | (ascii_to_ebcdic_pc[(w1.s1 >> 24) & 0xff]) << 24;
- #endif
-
- #ifdef VECT_SIZE4
- key[0].s0 = (ascii_to_ebcdic_pc[(w0.s0 >> 0) & 0xff]) << 0
- | (ascii_to_ebcdic_pc[(w0.s0 >> 8) & 0xff]) << 8
- | (ascii_to_ebcdic_pc[(w0.s0 >> 16) & 0xff]) << 16
- | (ascii_to_ebcdic_pc[(w0.s0 >> 24) & 0xff]) << 24;
-
- key[0].s1 = (ascii_to_ebcdic_pc[(w0.s1 >> 0) & 0xff]) << 0
- | (ascii_to_ebcdic_pc[(w0.s1 >> 8) & 0xff]) << 8
- | (ascii_to_ebcdic_pc[(w0.s1 >> 16) & 0xff]) << 16
- | (ascii_to_ebcdic_pc[(w0.s1 >> 24) & 0xff]) << 24;
-
- key[0].s2 = (ascii_to_ebcdic_pc[(w0.s2 >> 0) & 0xff]) << 0
- | (ascii_to_ebcdic_pc[(w0.s2 >> 8) & 0xff]) << 8
- | (ascii_to_ebcdic_pc[(w0.s2 >> 16) & 0xff]) << 16
- | (ascii_to_ebcdic_pc[(w0.s2 >> 24) & 0xff]) << 24;
-
- key[0].s3 = (ascii_to_ebcdic_pc[(w0.s3 >> 0) & 0xff]) << 0
- | (ascii_to_ebcdic_pc[(w0.s3 >> 8) & 0xff]) << 8
- | (ascii_to_ebcdic_pc[(w0.s3 >> 16) & 0xff]) << 16
- | (ascii_to_ebcdic_pc[(w0.s3 >> 24) & 0xff]) << 24;
-
- key[1].s0 = (ascii_to_ebcdic_pc[(w1.s0 >> 0) & 0xff]) << 0
- | (ascii_to_ebcdic_pc[(w1.s0 >> 8) & 0xff]) << 8
- | (ascii_to_ebcdic_pc[(w1.s0 >> 16) & 0xff]) << 16
- | (ascii_to_ebcdic_pc[(w1.s0 >> 24) & 0xff]) << 24;
-
- key[1].s1 = (ascii_to_ebcdic_pc[(w1.s1 >> 0) & 0xff]) << 0
- | (ascii_to_ebcdic_pc[(w1.s1 >> 8) & 0xff]) << 8
- | (ascii_to_ebcdic_pc[(w1.s1 >> 16) & 0xff]) << 16
- | (ascii_to_ebcdic_pc[(w1.s1 >> 24) & 0xff]) << 24;
-
- key[1].s2 = (ascii_to_ebcdic_pc[(w1.s2 >> 0) & 0xff]) << 0
- | (ascii_to_ebcdic_pc[(w1.s2 >> 8) & 0xff]) << 8
- | (ascii_to_ebcdic_pc[(w1.s2 >> 16) & 0xff]) << 16
- | (ascii_to_ebcdic_pc[(w1.s2 >> 24) & 0xff]) << 24;
-
- key[1].s3 = (ascii_to_ebcdic_pc[(w1.s3 >> 0) & 0xff]) << 0
- | (ascii_to_ebcdic_pc[(w1.s3 >> 8) & 0xff]) << 8
- | (ascii_to_ebcdic_pc[(w1.s3 >> 16) & 0xff]) << 16
- | (ascii_to_ebcdic_pc[(w1.s3 >> 24) & 0xff]) << 24;
- #endif
+ key[0] = (ascii_to_ebcdic_pc[(w0 >> 0) & 0xff]) << 0
+ | (ascii_to_ebcdic_pc[(w0 >> 8) & 0xff]) << 8
+ | (ascii_to_ebcdic_pc[(w0 >> 16) & 0xff]) << 16
+ | (ascii_to_ebcdic_pc[(w0 >> 24) & 0xff]) << 24;
+
+ key[1] = (ascii_to_ebcdic_pc[(w1 >> 0) & 0xff]) << 0
+ | (ascii_to_ebcdic_pc[(w1 >> 8) & 0xff]) << 8
+ | (ascii_to_ebcdic_pc[(w1 >> 16) & 0xff]) << 16
+ | (ascii_to_ebcdic_pc[(w1 >> 24) & 0xff]) << 24;
}
__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m08500_m04 (__global pw_t *pws, __global gpu_rule_t * rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
u32 u = Kc[i] ^ r;
u32 t = Kd[i] ^ rotl32 (r, 28u);
- #ifdef VECT_SIZE1
- l ^= NBOX (((u >> 2) & 0x3f), 0, s_SPtrans)
- | NBOX (((u >> 10) & 0x3f), 2, s_SPtrans)
- | NBOX (((u >> 18) & 0x3f), 4, s_SPtrans)
- | NBOX (((u >> 26) & 0x3f), 6, s_SPtrans)
- | NBOX (((t >> 2) & 0x3f), 1, s_SPtrans)
- | NBOX (((t >> 10) & 0x3f), 3, s_SPtrans)
- | NBOX (((t >> 18) & 0x3f), 5, s_SPtrans)
- | NBOX (((t >> 26) & 0x3f), 7, s_SPtrans);
- #endif
-
- #ifdef VECT_SIZE2
- l.s0 ^= NBOX (((u.s0 >> 2) & 0x3f), 0, s_SPtrans)
- | NBOX (((u.s0 >> 10) & 0x3f), 2, s_SPtrans)
- | NBOX (((u.s0 >> 18) & 0x3f), 4, s_SPtrans)
- | NBOX (((u.s0 >> 26) & 0x3f), 6, s_SPtrans)
- | NBOX (((t.s0 >> 2) & 0x3f), 1, s_SPtrans)
- | NBOX (((t.s0 >> 10) & 0x3f), 3, s_SPtrans)
- | NBOX (((t.s0 >> 18) & 0x3f), 5, s_SPtrans)
- | NBOX (((t.s0 >> 26) & 0x3f), 7, s_SPtrans);
-
- l.s1 ^= NBOX (((u.s1 >> 2) & 0x3f), 0, s_SPtrans)
- | NBOX (((u.s1 >> 10) & 0x3f), 2, s_SPtrans)
- | NBOX (((u.s1 >> 18) & 0x3f), 4, s_SPtrans)
- | NBOX (((u.s1 >> 26) & 0x3f), 6, s_SPtrans)
- | NBOX (((t.s1 >> 2) & 0x3f), 1, s_SPtrans)
- | NBOX (((t.s1 >> 10) & 0x3f), 3, s_SPtrans)
- | NBOX (((t.s1 >> 18) & 0x3f), 5, s_SPtrans)
- | NBOX (((t.s1 >> 26) & 0x3f), 7, s_SPtrans);
- #endif
-
- #ifdef VECT_SIZE4
- l.s0 ^= NBOX (((u.s0 >> 2) & 0x3f), 0, s_SPtrans)
- | NBOX (((u.s0 >> 10) & 0x3f), 2, s_SPtrans)
- | NBOX (((u.s0 >> 18) & 0x3f), 4, s_SPtrans)
- | NBOX (((u.s0 >> 26) & 0x3f), 6, s_SPtrans)
- | NBOX (((t.s0 >> 2) & 0x3f), 1, s_SPtrans)
- | NBOX (((t.s0 >> 10) & 0x3f), 3, s_SPtrans)
- | NBOX (((t.s0 >> 18) & 0x3f), 5, s_SPtrans)
- | NBOX (((t.s0 >> 26) & 0x3f), 7, s_SPtrans);
-
- l.s1 ^= NBOX (((u.s1 >> 2) & 0x3f), 0, s_SPtrans)
- | NBOX (((u.s1 >> 10) & 0x3f), 2, s_SPtrans)
- | NBOX (((u.s1 >> 18) & 0x3f), 4, s_SPtrans)
- | NBOX (((u.s1 >> 26) & 0x3f), 6, s_SPtrans)
- | NBOX (((t.s1 >> 2) & 0x3f), 1, s_SPtrans)
- | NBOX (((t.s1 >> 10) & 0x3f), 3, s_SPtrans)
- | NBOX (((t.s1 >> 18) & 0x3f), 5, s_SPtrans)
- | NBOX (((t.s1 >> 26) & 0x3f), 7, s_SPtrans);
-
- l.s2 ^= NBOX (((u.s2 >> 2) & 0x3f), 0, s_SPtrans)
- | NBOX (((u.s2 >> 10) & 0x3f), 2, s_SPtrans)
- | NBOX (((u.s2 >> 18) & 0x3f), 4, s_SPtrans)
- | NBOX (((u.s2 >> 26) & 0x3f), 6, s_SPtrans)
- | NBOX (((t.s2 >> 2) & 0x3f), 1, s_SPtrans)
- | NBOX (((t.s2 >> 10) & 0x3f), 3, s_SPtrans)
- | NBOX (((t.s2 >> 18) & 0x3f), 5, s_SPtrans)
- | NBOX (((t.s2 >> 26) & 0x3f), 7, s_SPtrans);
-
- l.s3 ^= NBOX (((u.s3 >> 2) & 0x3f), 0, s_SPtrans)
- | NBOX (((u.s3 >> 10) & 0x3f), 2, s_SPtrans)
- | NBOX (((u.s3 >> 18) & 0x3f), 4, s_SPtrans)
- | NBOX (((u.s3 >> 26) & 0x3f), 6, s_SPtrans)
- | NBOX (((t.s3 >> 2) & 0x3f), 1, s_SPtrans)
- | NBOX (((t.s3 >> 10) & 0x3f), 3, s_SPtrans)
- | NBOX (((t.s3 >> 18) & 0x3f), 5, s_SPtrans)
- | NBOX (((t.s3 >> 26) & 0x3f), 7, s_SPtrans);
- #endif
+ l ^= NBOX (((u >> 2) & 0x3f), 0, s_SPtrans)
+ | NBOX (((u >> 10) & 0x3f), 2, s_SPtrans)
+ | NBOX (((u >> 18) & 0x3f), 4, s_SPtrans)
+ | NBOX (((u >> 26) & 0x3f), 6, s_SPtrans)
+ | NBOX (((t >> 2) & 0x3f), 1, s_SPtrans)
+ | NBOX (((t >> 10) & 0x3f), 3, s_SPtrans)
+ | NBOX (((t >> 18) & 0x3f), 5, s_SPtrans)
+ | NBOX (((t >> 26) & 0x3f), 7, s_SPtrans);
tt = l;
l = r;
u32 s;
u32 t;
- #ifdef VECT_SIZE1
- s = NBOX ((( c >> 0) & 0x3f), 0, s_skb)
- | NBOX ((((c >> 6) & 0x03)
- | ((c >> 7) & 0x3c)), 1, s_skb)
- | NBOX ((((c >> 13) & 0x0f)
- | ((c >> 14) & 0x30)), 2, s_skb)
- | NBOX ((((c >> 20) & 0x01)
- | ((c >> 21) & 0x06)
- | ((c >> 22) & 0x38)), 3, s_skb);
-
- t = NBOX ((( d >> 0) & 0x3f), 4, s_skb)
- | NBOX ((((d >> 7) & 0x03)
- | ((d >> 8) & 0x3c)), 5, s_skb)
- | NBOX ((((d >> 15) & 0x3f)), 6, s_skb)
- | NBOX ((((d >> 21) & 0x0f)
- | ((d >> 22) & 0x30)), 7, s_skb);
- #endif
-
- #ifdef VECT_SIZE2
- s.s0 = NBOX ((( c.s0 >> 0) & 0x3f), 0, s_skb)
- | NBOX ((((c.s0 >> 6) & 0x03)
- | ((c.s0 >> 7) & 0x3c)), 1, s_skb)
- | NBOX ((((c.s0 >> 13) & 0x0f)
- | ((c.s0 >> 14) & 0x30)), 2, s_skb)
- | NBOX ((((c.s0 >> 20) & 0x01)
- | ((c.s0 >> 21) & 0x06)
- | ((c.s0 >> 22) & 0x38)), 3, s_skb);
-
- t.s0 = NBOX ((( d.s0 >> 0) & 0x3f), 4, s_skb)
- | NBOX ((((d.s0 >> 7) & 0x03)
- | ((d.s0 >> 8) & 0x3c)), 5, s_skb)
- | NBOX ((((d.s0 >> 15) & 0x3f)), 6, s_skb)
- | NBOX ((((d.s0 >> 21) & 0x0f)
- | ((d.s0 >> 22) & 0x30)), 7, s_skb);
-
- s.s1 = NBOX ((( c.s1 >> 0) & 0x3f), 0, s_skb)
- | NBOX ((((c.s1 >> 6) & 0x03)
- | ((c.s1 >> 7) & 0x3c)), 1, s_skb)
- | NBOX ((((c.s1 >> 13) & 0x0f)
- | ((c.s1 >> 14) & 0x30)), 2, s_skb)
- | NBOX ((((c.s1 >> 20) & 0x01)
- | ((c.s1 >> 21) & 0x06)
- | ((c.s1 >> 22) & 0x38)), 3, s_skb);
-
- t.s1 = NBOX ((( d.s1 >> 0) & 0x3f), 4, s_skb)
- | NBOX ((((d.s1 >> 7) & 0x03)
- | ((d.s1 >> 8) & 0x3c)), 5, s_skb)
- | NBOX ((((d.s1 >> 15) & 0x3f)), 6, s_skb)
- | NBOX ((((d.s1 >> 21) & 0x0f)
- | ((d.s1 >> 22) & 0x30)), 7, s_skb);
- #endif
-
- #ifdef VECT_SIZE4
- s.s0 = NBOX ((( c.s0 >> 0) & 0x3f), 0, s_skb)
- | NBOX ((((c.s0 >> 6) & 0x03)
- | ((c.s0 >> 7) & 0x3c)), 1, s_skb)
- | NBOX ((((c.s0 >> 13) & 0x0f)
- | ((c.s0 >> 14) & 0x30)), 2, s_skb)
- | NBOX ((((c.s0 >> 20) & 0x01)
- | ((c.s0 >> 21) & 0x06)
- | ((c.s0 >> 22) & 0x38)), 3, s_skb);
-
- t.s0 = NBOX ((( d.s0 >> 0) & 0x3f), 4, s_skb)
- | NBOX ((((d.s0 >> 7) & 0x03)
- | ((d.s0 >> 8) & 0x3c)), 5, s_skb)
- | NBOX ((((d.s0 >> 15) & 0x3f)), 6, s_skb)
- | NBOX ((((d.s0 >> 21) & 0x0f)
- | ((d.s0 >> 22) & 0x30)), 7, s_skb);
-
- s.s1 = NBOX ((( c.s1 >> 0) & 0x3f), 0, s_skb)
- | NBOX ((((c.s1 >> 6) & 0x03)
- | ((c.s1 >> 7) & 0x3c)), 1, s_skb)
- | NBOX ((((c.s1 >> 13) & 0x0f)
- | ((c.s1 >> 14) & 0x30)), 2, s_skb)
- | NBOX ((((c.s1 >> 20) & 0x01)
- | ((c.s1 >> 21) & 0x06)
- | ((c.s1 >> 22) & 0x38)), 3, s_skb);
-
- t.s1 = NBOX ((( d.s1 >> 0) & 0x3f), 4, s_skb)
- | NBOX ((((d.s1 >> 7) & 0x03)
- | ((d.s1 >> 8) & 0x3c)), 5, s_skb)
- | NBOX ((((d.s1 >> 15) & 0x3f)), 6, s_skb)
- | NBOX ((((d.s1 >> 21) & 0x0f)
- | ((d.s1 >> 22) & 0x30)), 7, s_skb);
-
- s.s2 = NBOX ((( c.s2 >> 0) & 0x3f), 0, s_skb)
- | NBOX ((((c.s2 >> 6) & 0x03)
- | ((c.s2 >> 7) & 0x3c)), 1, s_skb)
- | NBOX ((((c.s2 >> 13) & 0x0f)
- | ((c.s2 >> 14) & 0x30)), 2, s_skb)
- | NBOX ((((c.s2 >> 20) & 0x01)
- | ((c.s2 >> 21) & 0x06)
- | ((c.s2 >> 22) & 0x38)), 3, s_skb);
-
- t.s2 = NBOX ((( d.s2 >> 0) & 0x3f), 4, s_skb)
- | NBOX ((((d.s2 >> 7) & 0x03)
- | ((d.s2 >> 8) & 0x3c)), 5, s_skb)
- | NBOX ((((d.s2 >> 15) & 0x3f)), 6, s_skb)
- | NBOX ((((d.s2 >> 21) & 0x0f)
- | ((d.s2 >> 22) & 0x30)), 7, s_skb);
-
- s.s3 = NBOX ((( c.s3 >> 0) & 0x3f), 0, s_skb)
- | NBOX ((((c.s3 >> 6) & 0x03)
- | ((c.s3 >> 7) & 0x3c)), 1, s_skb)
- | NBOX ((((c.s3 >> 13) & 0x0f)
- | ((c.s3 >> 14) & 0x30)), 2, s_skb)
- | NBOX ((((c.s3 >> 20) & 0x01)
- | ((c.s3 >> 21) & 0x06)
- | ((c.s3 >> 22) & 0x38)), 3, s_skb);
-
- t.s3 = NBOX ((( d.s3 >> 0) & 0x3f), 4, s_skb)
- | NBOX ((((d.s3 >> 7) & 0x03)
- | ((d.s3 >> 8) & 0x3c)), 5, s_skb)
- | NBOX ((((d.s3 >> 15) & 0x3f)), 6, s_skb)
- | NBOX ((((d.s3 >> 21) & 0x0f)
- | ((d.s3 >> 22) & 0x30)), 7, s_skb);
- #endif
+ s = NBOX ((( c >> 0) & 0x3f), 0, s_skb)
+ | NBOX ((((c >> 6) & 0x03)
+ | ((c >> 7) & 0x3c)), 1, s_skb)
+ | NBOX ((((c >> 13) & 0x0f)
+ | ((c >> 14) & 0x30)), 2, s_skb)
+ | NBOX ((((c >> 20) & 0x01)
+ | ((c >> 21) & 0x06)
+ | ((c >> 22) & 0x38)), 3, s_skb);
+
+ t = NBOX ((( d >> 0) & 0x3f), 4, s_skb)
+ | NBOX ((((d >> 7) & 0x03)
+ | ((d >> 8) & 0x3c)), 5, s_skb)
+ | NBOX ((((d >> 15) & 0x3f)), 6, s_skb)
+ | NBOX ((((d >> 21) & 0x0f)
+ | ((d >> 22) & 0x30)), 7, s_skb);
#if defined cl_amd_media_ops
Kc[i] = amd_bytealign (t, s << 16, 2);
static void transform_racf_key (const u32 w0, const u32 w1, u32 key[2])
{
- #ifdef VECT_SIZE1
-
- key[0] = (ascii_to_ebcdic_pc[(w0 >> 0) & 0xff]) << 0
- | (ascii_to_ebcdic_pc[(w0 >> 8) & 0xff]) << 8
- | (ascii_to_ebcdic_pc[(w0 >> 16) & 0xff]) << 16
- | (ascii_to_ebcdic_pc[(w0 >> 24) & 0xff]) << 24;
-
- key[1] = (ascii_to_ebcdic_pc[(w1 >> 0) & 0xff]) << 0
- | (ascii_to_ebcdic_pc[(w1 >> 8) & 0xff]) << 8
- | (ascii_to_ebcdic_pc[(w1 >> 16) & 0xff]) << 16
- | (ascii_to_ebcdic_pc[(w1 >> 24) & 0xff]) << 24;
- #endif
-
- #ifdef VECT_SIZE2
-
- key[0].s0 = (ascii_to_ebcdic_pc[(w0.s0 >> 0) & 0xff]) << 0
- | (ascii_to_ebcdic_pc[(w0.s0 >> 8) & 0xff]) << 8
- | (ascii_to_ebcdic_pc[(w0.s0 >> 16) & 0xff]) << 16
- | (ascii_to_ebcdic_pc[(w0.s0 >> 24) & 0xff]) << 24;
-
- key[0].s1 = (ascii_to_ebcdic_pc[(w0.s1 >> 0) & 0xff]) << 0
- | (ascii_to_ebcdic_pc[(w0.s1 >> 8) & 0xff]) << 8
- | (ascii_to_ebcdic_pc[(w0.s1 >> 16) & 0xff]) << 16
- | (ascii_to_ebcdic_pc[(w0.s1 >> 24) & 0xff]) << 24;
-
- key[1].s0 = (ascii_to_ebcdic_pc[(w1.s0 >> 0) & 0xff]) << 0
- | (ascii_to_ebcdic_pc[(w1.s0 >> 8) & 0xff]) << 8
- | (ascii_to_ebcdic_pc[(w1.s0 >> 16) & 0xff]) << 16
- | (ascii_to_ebcdic_pc[(w1.s0 >> 24) & 0xff]) << 24;
-
- key[1].s1 = (ascii_to_ebcdic_pc[(w1.s1 >> 0) & 0xff]) << 0
- | (ascii_to_ebcdic_pc[(w1.s1 >> 8) & 0xff]) << 8
- | (ascii_to_ebcdic_pc[(w1.s1 >> 16) & 0xff]) << 16
- | (ascii_to_ebcdic_pc[(w1.s1 >> 24) & 0xff]) << 24;
- #endif
-
- #ifdef VECT_SIZE4
- key[0].s0 = (ascii_to_ebcdic_pc[(w0.s0 >> 0) & 0xff]) << 0
- | (ascii_to_ebcdic_pc[(w0.s0 >> 8) & 0xff]) << 8
- | (ascii_to_ebcdic_pc[(w0.s0 >> 16) & 0xff]) << 16
- | (ascii_to_ebcdic_pc[(w0.s0 >> 24) & 0xff]) << 24;
-
- key[0].s1 = (ascii_to_ebcdic_pc[(w0.s1 >> 0) & 0xff]) << 0
- | (ascii_to_ebcdic_pc[(w0.s1 >> 8) & 0xff]) << 8
- | (ascii_to_ebcdic_pc[(w0.s1 >> 16) & 0xff]) << 16
- | (ascii_to_ebcdic_pc[(w0.s1 >> 24) & 0xff]) << 24;
-
- key[0].s2 = (ascii_to_ebcdic_pc[(w0.s2 >> 0) & 0xff]) << 0
- | (ascii_to_ebcdic_pc[(w0.s2 >> 8) & 0xff]) << 8
- | (ascii_to_ebcdic_pc[(w0.s2 >> 16) & 0xff]) << 16
- | (ascii_to_ebcdic_pc[(w0.s2 >> 24) & 0xff]) << 24;
-
- key[0].s3 = (ascii_to_ebcdic_pc[(w0.s3 >> 0) & 0xff]) << 0
- | (ascii_to_ebcdic_pc[(w0.s3 >> 8) & 0xff]) << 8
- | (ascii_to_ebcdic_pc[(w0.s3 >> 16) & 0xff]) << 16
- | (ascii_to_ebcdic_pc[(w0.s3 >> 24) & 0xff]) << 24;
-
- key[1].s0 = (ascii_to_ebcdic_pc[(w1.s0 >> 0) & 0xff]) << 0
- | (ascii_to_ebcdic_pc[(w1.s0 >> 8) & 0xff]) << 8
- | (ascii_to_ebcdic_pc[(w1.s0 >> 16) & 0xff]) << 16
- | (ascii_to_ebcdic_pc[(w1.s0 >> 24) & 0xff]) << 24;
-
- key[1].s1 = (ascii_to_ebcdic_pc[(w1.s1 >> 0) & 0xff]) << 0
- | (ascii_to_ebcdic_pc[(w1.s1 >> 8) & 0xff]) << 8
- | (ascii_to_ebcdic_pc[(w1.s1 >> 16) & 0xff]) << 16
- | (ascii_to_ebcdic_pc[(w1.s1 >> 24) & 0xff]) << 24;
-
- key[1].s2 = (ascii_to_ebcdic_pc[(w1.s2 >> 0) & 0xff]) << 0
- | (ascii_to_ebcdic_pc[(w1.s2 >> 8) & 0xff]) << 8
- | (ascii_to_ebcdic_pc[(w1.s2 >> 16) & 0xff]) << 16
- | (ascii_to_ebcdic_pc[(w1.s2 >> 24) & 0xff]) << 24;
-
- key[1].s3 = (ascii_to_ebcdic_pc[(w1.s3 >> 0) & 0xff]) << 0
- | (ascii_to_ebcdic_pc[(w1.s3 >> 8) & 0xff]) << 8
- | (ascii_to_ebcdic_pc[(w1.s3 >> 16) & 0xff]) << 16
- | (ascii_to_ebcdic_pc[(w1.s3 >> 24) & 0xff]) << 24;
- #endif
+ key[0] = (ascii_to_ebcdic_pc[(w0 >> 0) & 0xff]) << 0
+ | (ascii_to_ebcdic_pc[(w0 >> 8) & 0xff]) << 8
+ | (ascii_to_ebcdic_pc[(w0 >> 16) & 0xff]) << 16
+ | (ascii_to_ebcdic_pc[(w0 >> 24) & 0xff]) << 24;
+
+ key[1] = (ascii_to_ebcdic_pc[(w1 >> 0) & 0xff]) << 0
+ | (ascii_to_ebcdic_pc[(w1 >> 8) & 0xff]) << 8
+ | (ascii_to_ebcdic_pc[(w1 >> 16) & 0xff]) << 16
+ | (ascii_to_ebcdic_pc[(w1 >> 24) & 0xff]) << 24;
}
__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m08500_m04 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
u32 u = Kc[i] ^ r;
u32 t = Kd[i] ^ rotl32 (r, 28u);
- #ifdef VECT_SIZE1
- l ^= NBOX (((u >> 2) & 0x3f), 0, s_SPtrans)
- | NBOX (((u >> 10) & 0x3f), 2, s_SPtrans)
- | NBOX (((u >> 18) & 0x3f), 4, s_SPtrans)
- | NBOX (((u >> 26) & 0x3f), 6, s_SPtrans)
- | NBOX (((t >> 2) & 0x3f), 1, s_SPtrans)
- | NBOX (((t >> 10) & 0x3f), 3, s_SPtrans)
- | NBOX (((t >> 18) & 0x3f), 5, s_SPtrans)
- | NBOX (((t >> 26) & 0x3f), 7, s_SPtrans);
- #endif
-
- #ifdef VECT_SIZE2
- l.s0 ^= NBOX (((u.s0 >> 2) & 0x3f), 0, s_SPtrans)
- | NBOX (((u.s0 >> 10) & 0x3f), 2, s_SPtrans)
- | NBOX (((u.s0 >> 18) & 0x3f), 4, s_SPtrans)
- | NBOX (((u.s0 >> 26) & 0x3f), 6, s_SPtrans)
- | NBOX (((t.s0 >> 2) & 0x3f), 1, s_SPtrans)
- | NBOX (((t.s0 >> 10) & 0x3f), 3, s_SPtrans)
- | NBOX (((t.s0 >> 18) & 0x3f), 5, s_SPtrans)
- | NBOX (((t.s0 >> 26) & 0x3f), 7, s_SPtrans);
-
- l.s1 ^= NBOX (((u.s1 >> 2) & 0x3f), 0, s_SPtrans)
- | NBOX (((u.s1 >> 10) & 0x3f), 2, s_SPtrans)
- | NBOX (((u.s1 >> 18) & 0x3f), 4, s_SPtrans)
- | NBOX (((u.s1 >> 26) & 0x3f), 6, s_SPtrans)
- | NBOX (((t.s1 >> 2) & 0x3f), 1, s_SPtrans)
- | NBOX (((t.s1 >> 10) & 0x3f), 3, s_SPtrans)
- | NBOX (((t.s1 >> 18) & 0x3f), 5, s_SPtrans)
- | NBOX (((t.s1 >> 26) & 0x3f), 7, s_SPtrans);
- #endif
-
- #ifdef VECT_SIZE4
- l.s0 ^= NBOX (((u.s0 >> 2) & 0x3f), 0, s_SPtrans)
- | NBOX (((u.s0 >> 10) & 0x3f), 2, s_SPtrans)
- | NBOX (((u.s0 >> 18) & 0x3f), 4, s_SPtrans)
- | NBOX (((u.s0 >> 26) & 0x3f), 6, s_SPtrans)
- | NBOX (((t.s0 >> 2) & 0x3f), 1, s_SPtrans)
- | NBOX (((t.s0 >> 10) & 0x3f), 3, s_SPtrans)
- | NBOX (((t.s0 >> 18) & 0x3f), 5, s_SPtrans)
- | NBOX (((t.s0 >> 26) & 0x3f), 7, s_SPtrans);
-
- l.s1 ^= NBOX (((u.s1 >> 2) & 0x3f), 0, s_SPtrans)
- | NBOX (((u.s1 >> 10) & 0x3f), 2, s_SPtrans)
- | NBOX (((u.s1 >> 18) & 0x3f), 4, s_SPtrans)
- | NBOX (((u.s1 >> 26) & 0x3f), 6, s_SPtrans)
- | NBOX (((t.s1 >> 2) & 0x3f), 1, s_SPtrans)
- | NBOX (((t.s1 >> 10) & 0x3f), 3, s_SPtrans)
- | NBOX (((t.s1 >> 18) & 0x3f), 5, s_SPtrans)
- | NBOX (((t.s1 >> 26) & 0x3f), 7, s_SPtrans);
-
- l.s2 ^= NBOX (((u.s2 >> 2) & 0x3f), 0, s_SPtrans)
- | NBOX (((u.s2 >> 10) & 0x3f), 2, s_SPtrans)
- | NBOX (((u.s2 >> 18) & 0x3f), 4, s_SPtrans)
- | NBOX (((u.s2 >> 26) & 0x3f), 6, s_SPtrans)
- | NBOX (((t.s2 >> 2) & 0x3f), 1, s_SPtrans)
- | NBOX (((t.s2 >> 10) & 0x3f), 3, s_SPtrans)
- | NBOX (((t.s2 >> 18) & 0x3f), 5, s_SPtrans)
- | NBOX (((t.s2 >> 26) & 0x3f), 7, s_SPtrans);
-
- l.s3 ^= NBOX (((u.s3 >> 2) & 0x3f), 0, s_SPtrans)
- | NBOX (((u.s3 >> 10) & 0x3f), 2, s_SPtrans)
- | NBOX (((u.s3 >> 18) & 0x3f), 4, s_SPtrans)
- | NBOX (((u.s3 >> 26) & 0x3f), 6, s_SPtrans)
- | NBOX (((t.s3 >> 2) & 0x3f), 1, s_SPtrans)
- | NBOX (((t.s3 >> 10) & 0x3f), 3, s_SPtrans)
- | NBOX (((t.s3 >> 18) & 0x3f), 5, s_SPtrans)
- | NBOX (((t.s3 >> 26) & 0x3f), 7, s_SPtrans);
- #endif
+ l ^= NBOX (((u >> 2) & 0x3f), 0, s_SPtrans)
+ | NBOX (((u >> 10) & 0x3f), 2, s_SPtrans)
+ | NBOX (((u >> 18) & 0x3f), 4, s_SPtrans)
+ | NBOX (((u >> 26) & 0x3f), 6, s_SPtrans)
+ | NBOX (((t >> 2) & 0x3f), 1, s_SPtrans)
+ | NBOX (((t >> 10) & 0x3f), 3, s_SPtrans)
+ | NBOX (((t >> 18) & 0x3f), 5, s_SPtrans)
+ | NBOX (((t >> 26) & 0x3f), 7, s_SPtrans);
tt = l;
l = r;
u32 s;
u32 t;
- #ifdef VECT_SIZE1
- s = NBOX ((( c >> 0) & 0x3f), 0, s_skb)
- | NBOX ((((c >> 6) & 0x03)
- | ((c >> 7) & 0x3c)), 1, s_skb)
- | NBOX ((((c >> 13) & 0x0f)
- | ((c >> 14) & 0x30)), 2, s_skb)
- | NBOX ((((c >> 20) & 0x01)
- | ((c >> 21) & 0x06)
- | ((c >> 22) & 0x38)), 3, s_skb);
-
- t = NBOX ((( d >> 0) & 0x3f), 4, s_skb)
- | NBOX ((((d >> 7) & 0x03)
- | ((d >> 8) & 0x3c)), 5, s_skb)
- | NBOX ((((d >> 15) & 0x3f)), 6, s_skb)
- | NBOX ((((d >> 21) & 0x0f)
- | ((d >> 22) & 0x30)), 7, s_skb);
- #endif
-
- #ifdef VECT_SIZE2
- s.s0 = NBOX ((( c.s0 >> 0) & 0x3f), 0, s_skb)
- | NBOX ((((c.s0 >> 6) & 0x03)
- | ((c.s0 >> 7) & 0x3c)), 1, s_skb)
- | NBOX ((((c.s0 >> 13) & 0x0f)
- | ((c.s0 >> 14) & 0x30)), 2, s_skb)
- | NBOX ((((c.s0 >> 20) & 0x01)
- | ((c.s0 >> 21) & 0x06)
- | ((c.s0 >> 22) & 0x38)), 3, s_skb);
-
- t.s0 = NBOX ((( d.s0 >> 0) & 0x3f), 4, s_skb)
- | NBOX ((((d.s0 >> 7) & 0x03)
- | ((d.s0 >> 8) & 0x3c)), 5, s_skb)
- | NBOX ((((d.s0 >> 15) & 0x3f)), 6, s_skb)
- | NBOX ((((d.s0 >> 21) & 0x0f)
- | ((d.s0 >> 22) & 0x30)), 7, s_skb);
-
- s.s1 = NBOX ((( c.s1 >> 0) & 0x3f), 0, s_skb)
- | NBOX ((((c.s1 >> 6) & 0x03)
- | ((c.s1 >> 7) & 0x3c)), 1, s_skb)
- | NBOX ((((c.s1 >> 13) & 0x0f)
- | ((c.s1 >> 14) & 0x30)), 2, s_skb)
- | NBOX ((((c.s1 >> 20) & 0x01)
- | ((c.s1 >> 21) & 0x06)
- | ((c.s1 >> 22) & 0x38)), 3, s_skb);
-
- t.s1 = NBOX ((( d.s1 >> 0) & 0x3f), 4, s_skb)
- | NBOX ((((d.s1 >> 7) & 0x03)
- | ((d.s1 >> 8) & 0x3c)), 5, s_skb)
- | NBOX ((((d.s1 >> 15) & 0x3f)), 6, s_skb)
- | NBOX ((((d.s1 >> 21) & 0x0f)
- | ((d.s1 >> 22) & 0x30)), 7, s_skb);
- #endif
-
- #ifdef VECT_SIZE4
- s.s0 = NBOX ((( c.s0 >> 0) & 0x3f), 0, s_skb)
- | NBOX ((((c.s0 >> 6) & 0x03)
- | ((c.s0 >> 7) & 0x3c)), 1, s_skb)
- | NBOX ((((c.s0 >> 13) & 0x0f)
- | ((c.s0 >> 14) & 0x30)), 2, s_skb)
- | NBOX ((((c.s0 >> 20) & 0x01)
- | ((c.s0 >> 21) & 0x06)
- | ((c.s0 >> 22) & 0x38)), 3, s_skb);
-
- t.s0 = NBOX ((( d.s0 >> 0) & 0x3f), 4, s_skb)
- | NBOX ((((d.s0 >> 7) & 0x03)
- | ((d.s0 >> 8) & 0x3c)), 5, s_skb)
- | NBOX ((((d.s0 >> 15) & 0x3f)), 6, s_skb)
- | NBOX ((((d.s0 >> 21) & 0x0f)
- | ((d.s0 >> 22) & 0x30)), 7, s_skb);
-
- s.s1 = NBOX ((( c.s1 >> 0) & 0x3f), 0, s_skb)
- | NBOX ((((c.s1 >> 6) & 0x03)
- | ((c.s1 >> 7) & 0x3c)), 1, s_skb)
- | NBOX ((((c.s1 >> 13) & 0x0f)
- | ((c.s1 >> 14) & 0x30)), 2, s_skb)
- | NBOX ((((c.s1 >> 20) & 0x01)
- | ((c.s1 >> 21) & 0x06)
- | ((c.s1 >> 22) & 0x38)), 3, s_skb);
-
- t.s1 = NBOX ((( d.s1 >> 0) & 0x3f), 4, s_skb)
- | NBOX ((((d.s1 >> 7) & 0x03)
- | ((d.s1 >> 8) & 0x3c)), 5, s_skb)
- | NBOX ((((d.s1 >> 15) & 0x3f)), 6, s_skb)
- | NBOX ((((d.s1 >> 21) & 0x0f)
- | ((d.s1 >> 22) & 0x30)), 7, s_skb);
-
- s.s2 = NBOX ((( c.s2 >> 0) & 0x3f), 0, s_skb)
- | NBOX ((((c.s2 >> 6) & 0x03)
- | ((c.s2 >> 7) & 0x3c)), 1, s_skb)
- | NBOX ((((c.s2 >> 13) & 0x0f)
- | ((c.s2 >> 14) & 0x30)), 2, s_skb)
- | NBOX ((((c.s2 >> 20) & 0x01)
- | ((c.s2 >> 21) & 0x06)
- | ((c.s2 >> 22) & 0x38)), 3, s_skb);
-
- t.s2 = NBOX ((( d.s2 >> 0) & 0x3f), 4, s_skb)
- | NBOX ((((d.s2 >> 7) & 0x03)
- | ((d.s2 >> 8) & 0x3c)), 5, s_skb)
- | NBOX ((((d.s2 >> 15) & 0x3f)), 6, s_skb)
- | NBOX ((((d.s2 >> 21) & 0x0f)
- | ((d.s2 >> 22) & 0x30)), 7, s_skb);
-
- s.s3 = NBOX ((( c.s3 >> 0) & 0x3f), 0, s_skb)
- | NBOX ((((c.s3 >> 6) & 0x03)
- | ((c.s3 >> 7) & 0x3c)), 1, s_skb)
- | NBOX ((((c.s3 >> 13) & 0x0f)
- | ((c.s3 >> 14) & 0x30)), 2, s_skb)
- | NBOX ((((c.s3 >> 20) & 0x01)
- | ((c.s3 >> 21) & 0x06)
- | ((c.s3 >> 22) & 0x38)), 3, s_skb);
-
- t.s3 = NBOX ((( d.s3 >> 0) & 0x3f), 4, s_skb)
- | NBOX ((((d.s3 >> 7) & 0x03)
- | ((d.s3 >> 8) & 0x3c)), 5, s_skb)
- | NBOX ((((d.s3 >> 15) & 0x3f)), 6, s_skb)
- | NBOX ((((d.s3 >> 21) & 0x0f)
- | ((d.s3 >> 22) & 0x30)), 7, s_skb);
- #endif
+ s = NBOX ((( c >> 0) & 0x3f), 0, s_skb)
+ | NBOX ((((c >> 6) & 0x03)
+ | ((c >> 7) & 0x3c)), 1, s_skb)
+ | NBOX ((((c >> 13) & 0x0f)
+ | ((c >> 14) & 0x30)), 2, s_skb)
+ | NBOX ((((c >> 20) & 0x01)
+ | ((c >> 21) & 0x06)
+ | ((c >> 22) & 0x38)), 3, s_skb);
+
+ t = NBOX ((( d >> 0) & 0x3f), 4, s_skb)
+ | NBOX ((((d >> 7) & 0x03)
+ | ((d >> 8) & 0x3c)), 5, s_skb)
+ | NBOX ((((d >> 15) & 0x3f)), 6, s_skb)
+ | NBOX ((((d >> 21) & 0x0f)
+ | ((d >> 22) & 0x30)), 7, s_skb);
#if defined cl_amd_media_ops
Kc[i] = amd_bytealign (t, s << 16, 2);
static void transform_racf_key (const u32 w0, const u32 w1, u32 key[2])
{
- #ifdef VECT_SIZE1
-
- key[0] = (ascii_to_ebcdic_pc[(w0 >> 0) & 0xff]) << 0
- | (ascii_to_ebcdic_pc[(w0 >> 8) & 0xff]) << 8
- | (ascii_to_ebcdic_pc[(w0 >> 16) & 0xff]) << 16
- | (ascii_to_ebcdic_pc[(w0 >> 24) & 0xff]) << 24;
-
- key[1] = (ascii_to_ebcdic_pc[(w1 >> 0) & 0xff]) << 0
- | (ascii_to_ebcdic_pc[(w1 >> 8) & 0xff]) << 8
- | (ascii_to_ebcdic_pc[(w1 >> 16) & 0xff]) << 16
- | (ascii_to_ebcdic_pc[(w1 >> 24) & 0xff]) << 24;
- #endif
-
- #ifdef VECT_SIZE2
-
- key[0].s0 = (ascii_to_ebcdic_pc[(w0.s0 >> 0) & 0xff]) << 0
- | (ascii_to_ebcdic_pc[(w0.s0 >> 8) & 0xff]) << 8
- | (ascii_to_ebcdic_pc[(w0.s0 >> 16) & 0xff]) << 16
- | (ascii_to_ebcdic_pc[(w0.s0 >> 24) & 0xff]) << 24;
-
- key[0].s1 = (ascii_to_ebcdic_pc[(w0.s1 >> 0) & 0xff]) << 0
- | (ascii_to_ebcdic_pc[(w0.s1 >> 8) & 0xff]) << 8
- | (ascii_to_ebcdic_pc[(w0.s1 >> 16) & 0xff]) << 16
- | (ascii_to_ebcdic_pc[(w0.s1 >> 24) & 0xff]) << 24;
-
- key[1].s0 = (ascii_to_ebcdic_pc[(w1.s0 >> 0) & 0xff]) << 0
- | (ascii_to_ebcdic_pc[(w1.s0 >> 8) & 0xff]) << 8
- | (ascii_to_ebcdic_pc[(w1.s0 >> 16) & 0xff]) << 16
- | (ascii_to_ebcdic_pc[(w1.s0 >> 24) & 0xff]) << 24;
-
- key[1].s1 = (ascii_to_ebcdic_pc[(w1.s1 >> 0) & 0xff]) << 0
- | (ascii_to_ebcdic_pc[(w1.s1 >> 8) & 0xff]) << 8
- | (ascii_to_ebcdic_pc[(w1.s1 >> 16) & 0xff]) << 16
- | (ascii_to_ebcdic_pc[(w1.s1 >> 24) & 0xff]) << 24;
- #endif
-
- #ifdef VECT_SIZE4
- key[0].s0 = (ascii_to_ebcdic_pc[(w0.s0 >> 0) & 0xff]) << 0
- | (ascii_to_ebcdic_pc[(w0.s0 >> 8) & 0xff]) << 8
- | (ascii_to_ebcdic_pc[(w0.s0 >> 16) & 0xff]) << 16
- | (ascii_to_ebcdic_pc[(w0.s0 >> 24) & 0xff]) << 24;
-
- key[0].s1 = (ascii_to_ebcdic_pc[(w0.s1 >> 0) & 0xff]) << 0
- | (ascii_to_ebcdic_pc[(w0.s1 >> 8) & 0xff]) << 8
- | (ascii_to_ebcdic_pc[(w0.s1 >> 16) & 0xff]) << 16
- | (ascii_to_ebcdic_pc[(w0.s1 >> 24) & 0xff]) << 24;
-
- key[0].s2 = (ascii_to_ebcdic_pc[(w0.s2 >> 0) & 0xff]) << 0
- | (ascii_to_ebcdic_pc[(w0.s2 >> 8) & 0xff]) << 8
- | (ascii_to_ebcdic_pc[(w0.s2 >> 16) & 0xff]) << 16
- | (ascii_to_ebcdic_pc[(w0.s2 >> 24) & 0xff]) << 24;
-
- key[0].s3 = (ascii_to_ebcdic_pc[(w0.s3 >> 0) & 0xff]) << 0
- | (ascii_to_ebcdic_pc[(w0.s3 >> 8) & 0xff]) << 8
- | (ascii_to_ebcdic_pc[(w0.s3 >> 16) & 0xff]) << 16
- | (ascii_to_ebcdic_pc[(w0.s3 >> 24) & 0xff]) << 24;
-
- key[1].s0 = (ascii_to_ebcdic_pc[(w1.s0 >> 0) & 0xff]) << 0
- | (ascii_to_ebcdic_pc[(w1.s0 >> 8) & 0xff]) << 8
- | (ascii_to_ebcdic_pc[(w1.s0 >> 16) & 0xff]) << 16
- | (ascii_to_ebcdic_pc[(w1.s0 >> 24) & 0xff]) << 24;
-
- key[1].s1 = (ascii_to_ebcdic_pc[(w1.s1 >> 0) & 0xff]) << 0
- | (ascii_to_ebcdic_pc[(w1.s1 >> 8) & 0xff]) << 8
- | (ascii_to_ebcdic_pc[(w1.s1 >> 16) & 0xff]) << 16
- | (ascii_to_ebcdic_pc[(w1.s1 >> 24) & 0xff]) << 24;
-
- key[1].s2 = (ascii_to_ebcdic_pc[(w1.s2 >> 0) & 0xff]) << 0
- | (ascii_to_ebcdic_pc[(w1.s2 >> 8) & 0xff]) << 8
- | (ascii_to_ebcdic_pc[(w1.s2 >> 16) & 0xff]) << 16
- | (ascii_to_ebcdic_pc[(w1.s2 >> 24) & 0xff]) << 24;
-
- key[1].s3 = (ascii_to_ebcdic_pc[(w1.s3 >> 0) & 0xff]) << 0
- | (ascii_to_ebcdic_pc[(w1.s3 >> 8) & 0xff]) << 8
- | (ascii_to_ebcdic_pc[(w1.s3 >> 16) & 0xff]) << 16
- | (ascii_to_ebcdic_pc[(w1.s3 >> 24) & 0xff]) << 24;
- #endif
+ key[0] = (ascii_to_ebcdic_pc[(w0 >> 0) & 0xff]) << 0
+ | (ascii_to_ebcdic_pc[(w0 >> 8) & 0xff]) << 8
+ | (ascii_to_ebcdic_pc[(w0 >> 16) & 0xff]) << 16
+ | (ascii_to_ebcdic_pc[(w0 >> 24) & 0xff]) << 24;
+
+ key[1] = (ascii_to_ebcdic_pc[(w1 >> 0) & 0xff]) << 0
+ | (ascii_to_ebcdic_pc[(w1 >> 8) & 0xff]) << 8
+ | (ascii_to_ebcdic_pc[(w1 >> 16) & 0xff]) << 16
+ | (ascii_to_ebcdic_pc[(w1 >> 24) & 0xff]) << 24;
}
static void m08500m (__local u32 s_SPtrans[8][64], __local u32 s_skb[8][64], u32 w[16], const u32 pw_len, __global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global u32 * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset)
0x29, 0x39, 0xb9, 0xe9, 0x4c, 0xff, 0x43, 0xab,
};
-#ifdef VECT_SIZE1
-#define BOX(S,i) u32 ((S)[(i)])
-#endif
-
-#ifdef VECT_SIZE2
-#define BOX(S,i) u32 ((S)[(i).s0], (S)[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define BOX(S,i) u32 ((S)[(i).s0], (S)[(i).s1], (S)[(i).s2], (S)[(i).s3])
-#endif
+#define BOX(S,i) (S)[(i)]
static void lotus_mix (u32 *in, __local u32 s_lotus_magic_table[256])
{
0x29, 0x39, 0xb9, 0xe9, 0x4c, 0xff, 0x43, 0xab,
};
-#ifdef VECT_SIZE1
-#define BOX(S,i) u32 ((S)[(i)])
-#endif
-
-#ifdef VECT_SIZE2
-#define BOX(S,i) u32 ((S)[(i).s0], (S)[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define BOX(S,i) u32 ((S)[(i).s0], (S)[(i).s1], (S)[(i).s2], (S)[(i).s3])
-#endif
+#define BOX(S,i) (S)[(i)]
static void lotus_mix (u32 *in, __local u32 s_lotus_magic_table[256])
{
0x29, 0x39, 0xb9, 0xe9, 0x4c, 0xff, 0x43, 0xab,
};
-#ifdef VECT_SIZE1
-#define BOX(S,i) u32 ((S)[(i)])
-#endif
-
-#ifdef VECT_SIZE2
-#define BOX(S,i) u32 ((S)[(i).s0], (S)[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define BOX(S,i) u32 ((S)[(i).s0], (S)[(i).s1], (S)[(i).s2], (S)[(i).s3])
-#endif
+#define BOX(S,i) (S)[(i)]
static void lotus_mix (u32 *in, __local u32 s_lotus_magic_table[256])
{
0x29, 0x39, 0xb9, 0xe9, 0x4c, 0xff, 0x43, 0xab,
};
-#ifdef VECT_SIZE1
-#define BOX(S,i) (u32x) ((S)[(i)])
-#endif
+#define BOX(S,i) (S)[(i)]
-#ifdef VECT_SIZE2
-#define BOX(S,i) (u32x) ((S)[(i).s0], (S)[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define BOX(S,i) (u32x) ((S)[(i).s0], (S)[(i).s1], (S)[(i).s2], (S)[(i).s3])
-#endif
-
-#ifdef VECT_SIZE1
-#define uint_to_hex_upper8(i) (u32x) (l_bin2asc[(i)])
-#endif
-
-#ifdef VECT_SIZE2
-#define uint_to_hex_upper8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define uint_to_hex_upper8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3])
-#endif
+#define uint_to_hex_upper8(i) l_bin2asc[(i)]
static void lotus_mix (u32 *in, __local u32 s_lotus_magic_table[256])
{
0x29, 0x39, 0xb9, 0xe9, 0x4c, 0xff, 0x43, 0xab,
};
-#ifdef VECT_SIZE1
-#define BOX(S,i) (u32x) ((S)[(i)])
-#endif
+#define BOX(S,i) (S)[(i)]
-#ifdef VECT_SIZE2
-#define BOX(S,i) (u32x) ((S)[(i).s0], (S)[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define BOX(S,i) (u32x) ((S)[(i).s0], (S)[(i).s1], (S)[(i).s2], (S)[(i).s3])
-#endif
-
-#ifdef VECT_SIZE1
-#define uint_to_hex_upper8(i) (u32x) (l_bin2asc[(i)])
-#endif
-
-#ifdef VECT_SIZE2
-#define uint_to_hex_upper8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define uint_to_hex_upper8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3])
-#endif
+#define uint_to_hex_upper8(i) l_bin2asc[(i)]
static void lotus_mix (u32 *in, __local u32 s_lotus_magic_table[256])
{
0x29, 0x39, 0xb9, 0xe9, 0x4c, 0xff, 0x43, 0xab,
};
-#ifdef VECT_SIZE1
-#define BOX(S,i) (u32x) ((S)[(i)])
-#endif
+#define BOX(S,i) (S)[(i)]
-#ifdef VECT_SIZE2
-#define BOX(S,i) (u32x) ((S)[(i).s0], (S)[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define BOX(S,i) (u32x) ((S)[(i).s0], (S)[(i).s1], (S)[(i).s2], (S)[(i).s3])
-#endif
-
-#ifdef VECT_SIZE1
-#define uint_to_hex_upper8(i) (u32x) (l_bin2asc[(i)])
-#endif
-
-#ifdef VECT_SIZE2
-#define uint_to_hex_upper8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define uint_to_hex_upper8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3])
-#endif
+#define uint_to_hex_upper8(i) l_bin2asc[(i)]
static void lotus_mix (u32 *in, __local u32 s_lotus_magic_table[256])
{
#include "types_ocl.c"
#include "common.c"
-#ifdef VECT_SIZE1
-#define COMPARE_M "check_multi_vect1_comp4.c"
-#endif
+#define COMPARE_S "check_single_comp4.c"
+#define COMPARE_M "check_multi_comp4.c"
// http://www.schneier.com/code/constants.txt
0x29, 0x39, 0xb9, 0xe9, 0x4c, 0xff, 0x43, 0xab,
};
-#ifdef VECT_SIZE1
-#define BOX(S,i) u32 ((S)[(i)])
-#endif
+#define BOX(S,i) (S)[(i)]
-#ifdef VECT_SIZE2
-#define BOX(S,i) u32 ((S)[(i).s0], (S)[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define BOX(S,i) u32 ((S)[(i).s0], (S)[(i).s1], (S)[(i).s2], (S)[(i).s3])
-#endif
-
-#ifdef VECT_SIZE1
-#define uint_to_hex_upper8(i) u32 (l_bin2asc[(i)])
-#endif
-
-#ifdef VECT_SIZE2
-#define uint_to_hex_upper8(i) u32 (l_bin2asc[(i).s0], l_bin2asc[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define uint_to_hex_upper8(i) u32 (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3])
-#endif
+#define uint_to_hex_upper8(i) l_bin2asc[(i)]
static void lotus_mix (u32 *in, __local u32 s_lotus_magic_table[256])
{
uchar4 salt0c = as_uchar4 (salt0);
uchar4 salt1c = as_uchar4 (salt1);
- #ifdef VECT_SIZE1
uchar4 ac;
uchar4 bc;
uchar4 cc;
ac = as_uchar4 (a);
bc = as_uchar4 (b);
cc = as_uchar4 (c);
- #endif
-
- #ifdef VECT_SIZE2
- uchar4 ac[2];
- uchar4 bc[2];
- uchar4 cc[2];
-
- ac[0] = as_uchar4 (a.s0);
- bc[0] = as_uchar4 (b.s0);
- cc[0] = as_uchar4 (c.s0);
-
- ac[1] = as_uchar4 (a.s1);
- bc[1] = as_uchar4 (b.s1);
- cc[1] = as_uchar4 (c.s1);
- #endif
u8 tmp[24]; // size 22 (=pw_len) is needed but base64 needs size divisible by 4
base64_plain[ 3] = salt0c.s3;
base64_plain[ 3] -= -4; // dont ask!
base64_plain[ 4] = salt1c.s0;
-
- #ifdef VECT_SIZE1
base64_plain[ 5] = ac.s0;
base64_plain[ 6] = ac.s1;
base64_plain[ 7] = ac.s2;
base64_hash[19] = tmp[19];
base64_hash[20] = tmp[20];
base64_hash[21] = ')';
- #endif
-
- #ifdef VECT_SIZE2
- base64_plain[ 5] = ac[0].s0;
- base64_plain[ 6] = ac[0].s1;
- base64_plain[ 7] = ac[0].s2;
- base64_plain[ 8] = ac[0].s3;
- base64_plain[ 9] = bc[0].s0;
- base64_plain[10] = bc[0].s1;
- base64_plain[11] = bc[0].s2;
- base64_plain[12] = bc[0].s3;
- base64_plain[13] = cc[0].s0;
- base64_plain[14] = cc[0].s1;
- base64_plain[15] = cc[0].s2;
-
- /*
- * base64 encode the $salt.$digest string
- */
-
- base64_encode (tmp + 2, 14, base64_plain);
-
- base64_hash[ 0].s0 = '(';
- base64_hash[ 1].s0 = 'G';
- base64_hash[ 2].s0 = tmp[ 2];
- base64_hash[ 3].s0 = tmp[ 3];
- base64_hash[ 4].s0 = tmp[ 4];
- base64_hash[ 5].s0 = tmp[ 5];
- base64_hash[ 6].s0 = tmp[ 6];
- base64_hash[ 7].s0 = tmp[ 7];
- base64_hash[ 8].s0 = tmp[ 8];
- base64_hash[ 9].s0 = tmp[ 9];
- base64_hash[10].s0 = tmp[10];
- base64_hash[11].s0 = tmp[11];
- base64_hash[12].s0 = tmp[12];
- base64_hash[13].s0 = tmp[13];
- base64_hash[14].s0 = tmp[14];
- base64_hash[15].s0 = tmp[15];
- base64_hash[16].s0 = tmp[16];
- base64_hash[17].s0 = tmp[17];
- base64_hash[18].s0 = tmp[18];
- base64_hash[19].s0 = tmp[19];
- base64_hash[20].s0 = tmp[20];
- base64_hash[21].s0 = ')';
-
- base64_plain[ 5] = ac[1].s0;
- base64_plain[ 6] = ac[1].s1;
- base64_plain[ 7] = ac[1].s2;
- base64_plain[ 8] = ac[1].s3;
- base64_plain[ 9] = bc[1].s0;
- base64_plain[10] = bc[1].s1;
- base64_plain[11] = bc[1].s2;
- base64_plain[12] = bc[1].s3;
- base64_plain[13] = cc[1].s0;
- base64_plain[14] = cc[1].s1;
- base64_plain[15] = cc[1].s2;
-
- /*
- * base64 encode the $salt.$digest string
- */
-
- base64_encode (tmp + 2, 14, base64_plain);
-
- base64_hash[ 0].s1 = '(';
- base64_hash[ 1].s1 = 'G';
- base64_hash[ 2].s1 = tmp[ 2];
- base64_hash[ 3].s1 = tmp[ 3];
- base64_hash[ 4].s1 = tmp[ 4];
- base64_hash[ 5].s1 = tmp[ 5];
- base64_hash[ 6].s1 = tmp[ 6];
- base64_hash[ 7].s1 = tmp[ 7];
- base64_hash[ 8].s1 = tmp[ 8];
- base64_hash[ 9].s1 = tmp[ 9];
- base64_hash[10].s1 = tmp[10];
- base64_hash[11].s1 = tmp[11];
- base64_hash[12].s1 = tmp[12];
- base64_hash[13].s1 = tmp[13];
- base64_hash[14].s1 = tmp[14];
- base64_hash[15].s1 = tmp[15];
- base64_hash[16].s1 = tmp[16];
- base64_hash[17].s1 = tmp[17];
- base64_hash[18].s1 = tmp[18];
- base64_hash[19].s1 = tmp[19];
- base64_hash[20].s1 = tmp[20];
- base64_hash[21].s1 = ')';
- #endif
-
}
-
__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m09100_init (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global lotus8_tmp_t *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global wpa_t *wpa_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
/**
lotus6_base64_encode (base64_hash, salt_buf0[0], salt_buf0[1], a, b, c);
-
/**
* PBKDF2 - HMACSHA1 - 1st iteration
*/
- #ifdef VECT_SIZE1
u32 w0[4];
+ u32 w1[4];
+ u32 w2[4];
+ u32 w3[4];
w0[0] = (base64_hash[ 0] << 24) | (base64_hash[ 1] << 16) | (base64_hash[ 2] << 8) | base64_hash[ 3];
w0[1] = (base64_hash[ 4] << 24) | (base64_hash[ 5] << 16) | (base64_hash[ 6] << 8) | base64_hash[ 7];
w0[2] = (base64_hash[ 8] << 24) | (base64_hash[ 9] << 16) | (base64_hash[10] << 8) | base64_hash[11];
w0[3] = (base64_hash[12] << 24) | (base64_hash[13] << 16) | (base64_hash[14] << 8) | base64_hash[15];
-
- u32 w1[4];
-
w1[0] = (base64_hash[16] << 24) | (base64_hash[17] << 16) | (base64_hash[18] << 8) | base64_hash[19];
w1[1] = (base64_hash[20] << 24) | (base64_hash[21] << 16);
w1[2] = 0;
w1[3] = 0;
- #endif
-
- #ifdef VECT_SIZE2
- u32 w0[4];
-
- w0[0].s0 = (base64_hash[ 0].s0 << 24) | (base64_hash[ 1].s0 << 16) | (base64_hash[ 2].s0 << 8) | base64_hash[ 3].s0;
- w0[1].s0 = (base64_hash[ 4].s0 << 24) | (base64_hash[ 5].s0 << 16) | (base64_hash[ 6].s0 << 8) | base64_hash[ 7].s0;
- w0[2].s0 = (base64_hash[ 8].s0 << 24) | (base64_hash[ 9].s0 << 16) | (base64_hash[10].s0 << 8) | base64_hash[11].s0;
- w0[3].s0 = (base64_hash[12].s0 << 24) | (base64_hash[13].s0 << 16) | (base64_hash[14].s0 << 8) | base64_hash[15].s0;
-
- w0[0].s1 = (base64_hash[ 0].s1 << 24) | (base64_hash[ 1].s1 << 16) | (base64_hash[ 2].s1 << 8) | base64_hash[ 3].s1;
- w0[1].s1 = (base64_hash[ 4].s1 << 24) | (base64_hash[ 5].s1 << 16) | (base64_hash[ 6].s1 << 8) | base64_hash[ 7].s1;
- w0[2].s1 = (base64_hash[ 8].s1 << 24) | (base64_hash[ 9].s1 << 16) | (base64_hash[10].s1 << 8) | base64_hash[11].s1;
- w0[3].s1 = (base64_hash[12].s1 << 24) | (base64_hash[13].s1 << 16) | (base64_hash[14].s1 << 8) | base64_hash[15].s1;
-
- u32 w1[4];
-
- w1[0].s0 = (base64_hash[16].s0 << 24) | (base64_hash[17].s0 << 16) | (base64_hash[18].s0 << 8) | base64_hash[19].s0;
- w1[1].s0 = (base64_hash[20].s0 << 24) | (base64_hash[21].s0 << 16);
- w1[2].s0 = 0;
- w1[3].s0 = 0;
-
- w1[0].s1 = (base64_hash[16].s1 << 24) | (base64_hash[17].s1 << 16) | (base64_hash[18].s1 << 8) | base64_hash[19].s1;
- w1[1].s1 = (base64_hash[20].s1 << 24) | (base64_hash[21].s1 << 16);
- w1[2].s1 = 0;
- w1[3].s1 = 0;
- #endif
-
- u32 w2[4];
-
w2[0] = 0;
w2[1] = 0;
w2[2] = 0;
w2[3] = 0;
-
- u32 w3[4];
-
w3[0] = 0;
w3[1] = 0;
w3[2] = 0;
#define COMPARE_S "check_single_comp4.c"
#define COMPARE_M "check_multi_comp4.c"
-#ifdef VECT_SIZE1
#define SIPROUND(v0,v1,v2,v3) \
(v0) += (v1); \
(v1) = rotl64 ((v1), 13); \
(v1) = rotl64 ((v1), 17); \
(v1) ^= (v2); \
(v2) = as_ulong (as_uint2 ((v2)).s10);
-#else
-#define SIPROUND(v0,v1,v2,v3) \
- (v0) += (v1); \
- (v1) = rotl64 ((v1), 13); \
- (v1) ^= (v0); \
- (v0) = rotl64 ((v0), 32); \
- (v2) += (v3); \
- (v3) = rotl64 ((v3), 16); \
- (v3) ^= (v2); \
- (v0) += (v3); \
- (v3) = rotl64 ((v3), 21); \
- (v3) ^= (v0); \
- (v2) += (v1); \
- (v1) = rotl64 ((v1), 17); \
- (v1) ^= (v2); \
- (v2) = rotl64 ((v2), 32);
-#endif
__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m10100_m04 (__global pw_t *pws, __global gpu_rule_t * rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
#define COMPARE_S "check_single_comp4.c"
#define COMPARE_M "check_multi_comp4.c"
-#ifdef VECT_SIZE1
#define SIPROUND(v0,v1,v2,v3) \
(v0) += (v1); \
(v1) = rotl64 ((v1), 13); \
(v1) = rotl64 ((v1), 17); \
(v1) ^= (v2); \
(v2) = as_ulong (as_uint2 ((v2)).s10);
-#else
-#define SIPROUND(v0,v1,v2,v3) \
- (v0) += (v1); \
- (v1) = rotl64 ((v1), 13); \
- (v1) ^= (v0); \
- (v0) = rotl64 ((v0), 32); \
- (v2) += (v3); \
- (v3) = rotl64 ((v3), 16); \
- (v3) ^= (v2); \
- (v0) += (v3); \
- (v3) = rotl64 ((v3), 21); \
- (v3) ^= (v0); \
- (v2) += (v1); \
- (v1) = rotl64 ((v1), 17); \
- (v1) ^= (v2); \
- (v2) = rotl64 ((v2), 32);
-#endif
__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m10100_m04 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
#define COMPARE_S "check_single_comp4.c"
#define COMPARE_M "check_multi_comp4.c"
-#ifdef VECT_SIZE1
#define SIPROUND(v0,v1,v2,v3) \
(v0) += (v1); \
(v1) = rotl64 ((v1), 13); \
(v1) = rotl64 ((v1), 17); \
(v1) ^= (v2); \
(v2) = as_ulong (as_uint2 ((v2)).s10);
-#else
-#define SIPROUND(v0,v1,v2,v3) \
- (v0) += (v1); \
- (v1) = rotl64 ((v1), 13); \
- (v1) ^= (v0); \
- (v0) = rotl64 ((v0), 32); \
- (v2) += (v3); \
- (v3) = rotl64 ((v3), 16); \
- (v3) ^= (v2); \
- (v0) += (v3); \
- (v3) = rotl64 ((v3), 21); \
- (v3) ^= (v0); \
- (v2) += (v1); \
- (v1) = rotl64 ((v1), 17); \
- (v1) ^= (v2); \
- (v2) = rotl64 ((v2), 32);
-#endif
static void m10100m (u32 w[16], const u32 pw_len, __global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global u32 * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset)
{
#include "types_ocl.c"
#include "common.c"
-#ifdef VECT_SIZE1
-#define COMPARE_M "check_multi_vect1_comp4.c"
-#endif
+#define COMPARE_S "check_single_comp4.c"
+#define COMPARE_M "check_multi_comp4.c"
__constant u32 k_sha256[64] =
{
#define COMPARE_S "check_single_comp4.c"
#define COMPARE_M "check_multi_comp4.c"
-#ifdef VECT_SIZE1
#define uint_to_hex_lower8(i) l_bin2asc[(i)]
-#endif
-
-#ifdef VECT_SIZE2
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3])
-#endif
__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11100_m04 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
#define COMPARE_S "check_single_comp4.c"
#define COMPARE_M "check_multi_comp4.c"
-#ifdef VECT_SIZE1
#define uint_to_hex_lower8(i) l_bin2asc[(i)]
-#endif
-
-#ifdef VECT_SIZE2
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3])
-#endif
__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11100_m04 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
#define COMPARE_S "check_single_comp4.c"
#define COMPARE_M "check_multi_comp4.c"
-#ifdef VECT_SIZE1
#define uint_to_hex_lower8(i) l_bin2asc[(i)]
-#endif
-
-#ifdef VECT_SIZE2
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3])
-#endif
static void m11100m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, __local u32 l_bin2asc[256])
{
const u32 s = a >> 8;
- #ifdef VECT_SIZE1
- a = crc32tab[k];
- #endif
-
- #ifdef VECT_SIZE2
- a.s0 = crc32tab[k.s0];
- a.s1 = crc32tab[k.s1];
- #endif
-
- #ifdef VECT_SIZE4
- a.s0 = crc32tab[k.s0];
- a.s1 = crc32tab[k.s1];
- a.s2 = crc32tab[k.s2];
- a.s3 = crc32tab[k.s3];
- #endif
+ a = crc32tab[k];
a ^= s;
const u32 s = a >> 8;
- #ifdef VECT_SIZE1
- a = crc32tab[k];
- #endif
-
- #ifdef VECT_SIZE2
- a.s0 = crc32tab[k.s0];
- a.s1 = crc32tab[k.s1];
- #endif
-
- #ifdef VECT_SIZE4
- a.s0 = crc32tab[k.s0];
- a.s1 = crc32tab[k.s1];
- a.s2 = crc32tab[k.s2];
- a.s3 = crc32tab[k.s3];
- #endif
+ a = crc32tab[k];
a ^= s;
const u32 s = a >> 8;
- #ifdef VECT_SIZE1
- a = crc32tab[k];
- #endif
-
- #ifdef VECT_SIZE2
- a.s0 = crc32tab[k.s0];
- a.s1 = crc32tab[k.s1];
- #endif
-
- #ifdef VECT_SIZE4
- a.s0 = crc32tab[k.s0];
- a.s1 = crc32tab[k.s1];
- a.s2 = crc32tab[k.s2];
- a.s3 = crc32tab[k.s3];
- #endif
+ a = crc32tab[k];
a ^= s;
const u32 s = a >> 8;
- #ifdef VECT_SIZE1
- a = crc32tab[k];
- #endif
-
- #ifdef VECT_SIZE2
- a.s0 = crc32tab[k.s0];
- a.s1 = crc32tab[k.s1];
- #endif
-
- #ifdef VECT_SIZE4
- a.s0 = crc32tab[k.s0];
- a.s1 = crc32tab[k.s1];
- a.s2 = crc32tab[k.s2];
- a.s3 = crc32tab[k.s3];
- #endif
+ a = crc32tab[k];
a ^= s;
#include "types_ocl.c"
#include "common.c"
-#ifdef VECT_SIZE1
-#define COMPARE_M "check_multi_vect1_comp4.c"
-#endif
+#define COMPARE_S "check_single_comp4.c"
+#define COMPARE_M "check_multi_comp4.c"
static void md5_transform (const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], u32 digest[4])
{
#include "types_ocl.c"
#include "common.c"
-#ifdef VECT_SIZE1
-#define COMPARE_M "check_multi_vect1_comp4.c"
-#endif
+#define COMPARE_S "check_single_comp4.c"
+#define COMPARE_M "check_multi_comp4.c"
static void sha1_transform (const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], u32 digest[5])
{
#include "types_ocl.c"
#include "common.c"
-#ifdef VECT_SIZE1
-#define COMPARE_M "check_multi_vect1_comp4.c"
-#endif
+#define COMPARE_S "check_single_comp4.c"
+#define COMPARE_M "check_multi_comp4.c"
__constant u64 k_sha512[80] =
{
#include "types_ocl.c"
#include "common.c"
-#ifdef VECT_SIZE1
-#define COMPARE_M "check_multi_vect1_comp4.c"
-#endif
+#define COMPARE_S "check_single_comp4.c"
+#define COMPARE_M "check_multi_comp4.c"
__constant u64 k_sha512[80] =
{
0x00002822, 0x04002822, 0x00042822, 0x04042822
};
-#ifdef VECT_SIZE1
-#define BOX(i,n,S) (u32x) ((S)[(n)][(i)])
-#endif
-
-#ifdef VECT_SIZE2
-#define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3])
-#endif
+#define BOX(i,n,S) (S)[(n)][(i)]
static void _des_crypt_keysetup (u32 c, u32 d, u32 Kc[16], u32 Kd[16], __local u32 s_skb[8][64])
{
#include "types_ocl.c"
#include "common.c"
-#ifdef VECT_SIZE1
-#define COMPARE_M "check_multi_vect1_comp4.c"
-#endif
+#define COMPARE_S "check_single_comp4.c"
+#define COMPARE_M "check_multi_comp4.c"
#define ROUNDS 0x40000
#define COMPARE_S "check_single_comp4.c"
#define COMPARE_M "check_multi_comp4.c"
-#ifdef VECT_SIZE1
#define uint_to_hex_upper8(i) l_bin2asc[(i)]
-#endif
-
-#ifdef VECT_SIZE2
-#define uint_to_hex_upper8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define uint_to_hex_upper8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3])
-#endif
__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m12600_m04 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
#define COMPARE_S "check_single_comp4.c"
#define COMPARE_M "check_multi_comp4.c"
-#ifdef VECT_SIZE1
#define uint_to_hex_upper8(i) l_bin2asc[(i)]
-#endif
-
-#ifdef VECT_SIZE2
-#define uint_to_hex_upper8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define uint_to_hex_upper8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3])
-#endif
__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m12600_m04 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
#define COMPARE_S "check_single_comp4.c"
#define COMPARE_M "check_multi_comp4.c"
-#ifdef VECT_SIZE1
#define uint_to_hex_upper8(i) l_bin2asc[(i)]
-#endif
-
-#ifdef VECT_SIZE2
-#define uint_to_hex_upper8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1])
-#endif
-
-#ifdef VECT_SIZE4
-#define uint_to_hex_upper8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3])
-#endif
static void m12600m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, __local u32 l_bin2asc[256])
{
#include "types_ocl.c"
#include "common.c"
-#ifdef VECT_SIZE1
-#define COMPARE_M "check_multi_vect1_comp4.c"
-#endif
-
-#ifdef VECT_SIZE2
-#define COMPARE_M "check_multi_vect2_comp4.c"
-#endif
+#define COMPARE_S "check_single_comp4.c"
+#define COMPARE_M "check_multi_comp4.c"
__constant u32 te0[256] =
{
#include "types_ocl.c"
#include "common.c"
-#ifdef VECT_SIZE1
-#define COMPARE_M "check_multi_vect1_comp4.c"
-#endif
+#define COMPARE_S "check_single_comp4.c"
+#define COMPARE_M "check_multi_comp4.c"
-#ifdef VECT_SIZE1
#define uint_to_hex_lower8(i) l_bin2asc[(i)]
-#endif
__constant u32 k_sha256[64] =
{
projects / hashcat.git / commitdiff