#define _DES_
-#include "include/constants.h"
-#include "include/kernel_vendor.h"
+#include "inc_vendor.cl"
+#include "inc_hash_constants.h"
+#include "inc_hash_functions.cl"
+#include "inc_types.cl"
+#include "inc_common.cl"
-#define DGST_R0 0
-#define DGST_R1 1
-#define DGST_R2 2
-#define DGST_R3 3
-
-#include "include/kernel_functions.c"
-#include "OpenCL/types_ocl.c"
-#include "OpenCL/common.c"
-
-#define COMPARE_S "OpenCL/check_single_comp4_bs.c"
-#define COMPARE_M "OpenCL/check_multi_comp4_bs.c"
+#define COMPARE_S "inc_comp_single_bs.cl"
+#define COMPARE_M "inc_comp_multi_bs.cl"
#ifdef IS_NV
#define KXX_DECL
KXX_DECL u32 k36, k37, k38, k39, k40, k41;
KXX_DECL u32 k42, k43, k44, k45, k46, k47;
- #ifdef IS_NV
- #if CUDA_ARCH >= 500
- #else
+ #ifdef _unroll
#pragma unroll
#endif
- #endif
-
- #ifdef IS_AMD
- #pragma unroll
- #endif
-
- #ifdef IS_GENERIC
- #pragma unroll 1
- #endif
-
for (u32 i = 0; i < 2; i++)
{
if (i) KEYSET10 else KEYSET00
swap (data[30], data[31], 1, 0x55555555);
}
-void m03000m (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * 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 il_cnt, const u32 digests_cnt, const u32 digests_offset)
+void m03000m (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * 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_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset)
{
/**
* base
const u32 w0s = pws[gid].i[0];
const u32 w1s = pws[gid].i[1];
- #define K00 -((w0s >> ( 0 + 7)) & 1)
- #define K01 -((w0s >> ( 0 + 6)) & 1)
- #define K02 -((w0s >> ( 0 + 5)) & 1)
- #define K03 -((w0s >> ( 0 + 4)) & 1)
- #define K04 -((w0s >> ( 0 + 3)) & 1)
- #define K05 -((w0s >> ( 0 + 2)) & 1)
- #define K06 -((w0s >> ( 0 + 1)) & 1)
- #define K07 -((w0s >> ( 0 + 0)) & 1)
- #define K08 -((w0s >> ( 8 + 7)) & 1)
- #define K09 -((w0s >> ( 8 + 6)) & 1)
- #define K10 -((w0s >> ( 8 + 5)) & 1)
- #define K11 -((w0s >> ( 8 + 4)) & 1)
- #define K12 -((w0s >> ( 8 + 3)) & 1)
- #define K13 -((w0s >> ( 8 + 2)) & 1)
- #define K14 -((w0s >> ( 8 + 1)) & 1)
- #define K15 -((w0s >> ( 8 + 0)) & 1)
- #define K16 -((w0s >> (16 + 7)) & 1)
- #define K17 -((w0s >> (16 + 6)) & 1)
- #define K18 -((w0s >> (16 + 5)) & 1)
- #define K19 -((w0s >> (16 + 4)) & 1)
- #define K20 -((w0s >> (16 + 3)) & 1)
- #define K21 -((w0s >> (16 + 2)) & 1)
- #define K22 -((w0s >> (16 + 1)) & 1)
- #define K23 -((w0s >> (16 + 0)) & 1)
- #define K24 -((w0s >> (24 + 7)) & 1)
- #define K25 -((w0s >> (24 + 6)) & 1)
- #define K26 -((w0s >> (24 + 5)) & 1)
- #define K27 -((w0s >> (24 + 4)) & 1)
- #define K28 -((w0s >> (24 + 3)) & 1)
- #define K29 -((w0s >> (24 + 2)) & 1)
- #define K30 -((w0s >> (24 + 1)) & 1)
- #define K31 -((w0s >> (24 + 0)) & 1)
- #define K32 -((w1s >> ( 0 + 7)) & 1)
- #define K33 -((w1s >> ( 0 + 6)) & 1)
- #define K34 -((w1s >> ( 0 + 5)) & 1)
- #define K35 -((w1s >> ( 0 + 4)) & 1)
- #define K36 -((w1s >> ( 0 + 3)) & 1)
- #define K37 -((w1s >> ( 0 + 2)) & 1)
- #define K38 -((w1s >> ( 0 + 1)) & 1)
- #define K39 -((w1s >> ( 0 + 0)) & 1)
- #define K40 -((w1s >> ( 8 + 7)) & 1)
- #define K41 -((w1s >> ( 8 + 6)) & 1)
- #define K42 -((w1s >> ( 8 + 5)) & 1)
- #define K43 -((w1s >> ( 8 + 4)) & 1)
- #define K44 -((w1s >> ( 8 + 3)) & 1)
- #define K45 -((w1s >> ( 8 + 2)) & 1)
- #define K46 -((w1s >> ( 8 + 1)) & 1)
- #define K47 -((w1s >> ( 8 + 0)) & 1)
- #define K48 -((w1s >> (16 + 7)) & 1)
- #define K49 -((w1s >> (16 + 6)) & 1)
- #define K50 -((w1s >> (16 + 5)) & 1)
- #define K51 -((w1s >> (16 + 4)) & 1)
- #define K52 -((w1s >> (16 + 3)) & 1)
- #define K53 -((w1s >> (16 + 2)) & 1)
- #define K54 -((w1s >> (16 + 1)) & 1)
- #define K55 -((w1s >> (16 + 0)) & 1)
+ #define K00 (((w0s >> ( 0 + 7)) & 1) ? -1 : 0)
+ #define K01 (((w0s >> ( 0 + 6)) & 1) ? -1 : 0)
+ #define K02 (((w0s >> ( 0 + 5)) & 1) ? -1 : 0)
+ #define K03 (((w0s >> ( 0 + 4)) & 1) ? -1 : 0)
+ #define K04 (((w0s >> ( 0 + 3)) & 1) ? -1 : 0)
+ #define K05 (((w0s >> ( 0 + 2)) & 1) ? -1 : 0)
+ #define K06 (((w0s >> ( 0 + 1)) & 1) ? -1 : 0)
+ #define K07 (((w0s >> ( 0 + 0)) & 1) ? -1 : 0)
+ #define K08 (((w0s >> ( 8 + 7)) & 1) ? -1 : 0)
+ #define K09 (((w0s >> ( 8 + 6)) & 1) ? -1 : 0)
+ #define K10 (((w0s >> ( 8 + 5)) & 1) ? -1 : 0)
+ #define K11 (((w0s >> ( 8 + 4)) & 1) ? -1 : 0)
+ #define K12 (((w0s >> ( 8 + 3)) & 1) ? -1 : 0)
+ #define K13 (((w0s >> ( 8 + 2)) & 1) ? -1 : 0)
+ #define K14 (((w0s >> ( 8 + 1)) & 1) ? -1 : 0)
+ #define K15 (((w0s >> ( 8 + 0)) & 1) ? -1 : 0)
+ #define K16 (((w0s >> (16 + 7)) & 1) ? -1 : 0)
+ #define K17 (((w0s >> (16 + 6)) & 1) ? -1 : 0)
+ #define K18 (((w0s >> (16 + 5)) & 1) ? -1 : 0)
+ #define K19 (((w0s >> (16 + 4)) & 1) ? -1 : 0)
+ #define K20 (((w0s >> (16 + 3)) & 1) ? -1 : 0)
+ #define K21 (((w0s >> (16 + 2)) & 1) ? -1 : 0)
+ #define K22 (((w0s >> (16 + 1)) & 1) ? -1 : 0)
+ #define K23 (((w0s >> (16 + 0)) & 1) ? -1 : 0)
+ #define K24 (((w0s >> (24 + 7)) & 1) ? -1 : 0)
+ #define K25 (((w0s >> (24 + 6)) & 1) ? -1 : 0)
+ #define K26 (((w0s >> (24 + 5)) & 1) ? -1 : 0)
+ #define K27 (((w0s >> (24 + 4)) & 1) ? -1 : 0)
+ #define K28 (((w0s >> (24 + 3)) & 1) ? -1 : 0)
+ #define K29 (((w0s >> (24 + 2)) & 1) ? -1 : 0)
+ #define K30 (((w0s >> (24 + 1)) & 1) ? -1 : 0)
+ #define K31 (((w0s >> (24 + 0)) & 1) ? -1 : 0)
+ #define K32 (((w1s >> ( 0 + 7)) & 1) ? -1 : 0)
+ #define K33 (((w1s >> ( 0 + 6)) & 1) ? -1 : 0)
+ #define K34 (((w1s >> ( 0 + 5)) & 1) ? -1 : 0)
+ #define K35 (((w1s >> ( 0 + 4)) & 1) ? -1 : 0)
+ #define K36 (((w1s >> ( 0 + 3)) & 1) ? -1 : 0)
+ #define K37 (((w1s >> ( 0 + 2)) & 1) ? -1 : 0)
+ #define K38 (((w1s >> ( 0 + 1)) & 1) ? -1 : 0)
+ #define K39 (((w1s >> ( 0 + 0)) & 1) ? -1 : 0)
+ #define K40 (((w1s >> ( 8 + 7)) & 1) ? -1 : 0)
+ #define K41 (((w1s >> ( 8 + 6)) & 1) ? -1 : 0)
+ #define K42 (((w1s >> ( 8 + 5)) & 1) ? -1 : 0)
+ #define K43 (((w1s >> ( 8 + 4)) & 1) ? -1 : 0)
+ #define K44 (((w1s >> ( 8 + 3)) & 1) ? -1 : 0)
+ #define K45 (((w1s >> ( 8 + 2)) & 1) ? -1 : 0)
+ #define K46 (((w1s >> ( 8 + 1)) & 1) ? -1 : 0)
+ #define K47 (((w1s >> ( 8 + 0)) & 1) ? -1 : 0)
+ #define K48 (((w1s >> (16 + 7)) & 1) ? -1 : 0)
+ #define K49 (((w1s >> (16 + 6)) & 1) ? -1 : 0)
+ #define K50 (((w1s >> (16 + 5)) & 1) ? -1 : 0)
+ #define K51 (((w1s >> (16 + 4)) & 1) ? -1 : 0)
+ #define K52 (((w1s >> (16 + 3)) & 1) ? -1 : 0)
+ #define K53 (((w1s >> (16 + 2)) & 1) ? -1 : 0)
+ #define K54 (((w1s >> (16 + 1)) & 1) ? -1 : 0)
+ #define K55 (((w1s >> (16 + 0)) & 1) ? -1 : 0)
/**
* inner loop
u32 tmpResult = 0;
+ #ifdef _unroll
#pragma unroll
+ #endif
for (int i = 0; i < 32; i++)
{
const u32 b0 = -((search[0] >> i) & 1);
u32 out0[32];
u32 out1[32];
+ #ifdef _unroll
#pragma unroll
+ #endif
for (int i = 0; i < 32; i++)
{
out0[i] = out[ 0 + 31 - i];
transpose32c (out0);
transpose32c (out1);
+ #ifdef _unroll
#pragma unroll
+ #endif
for (int slice = 0; slice < 32; slice++)
{
const u32 r0 = out0[31 - slice];
}
}
-void m03000s (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * 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 il_cnt, const u32 digests_cnt, const u32 digests_offset)
+void m03000s (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * 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_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset)
{
/**
* base
const u32 s0 = digests_buf[0].digest_buf[0];
const u32 s1 = digests_buf[0].digest_buf[1];
- #define S00 -((s0 >> 0) & 1)
- #define S01 -((s0 >> 1) & 1)
- #define S02 -((s0 >> 2) & 1)
- #define S03 -((s0 >> 3) & 1)
- #define S04 -((s0 >> 4) & 1)
- #define S05 -((s0 >> 5) & 1)
- #define S06 -((s0 >> 6) & 1)
- #define S07 -((s0 >> 7) & 1)
- #define S08 -((s0 >> 8) & 1)
- #define S09 -((s0 >> 9) & 1)
- #define S10 -((s0 >> 10) & 1)
- #define S11 -((s0 >> 11) & 1)
- #define S12 -((s0 >> 12) & 1)
- #define S13 -((s0 >> 13) & 1)
- #define S14 -((s0 >> 14) & 1)
- #define S15 -((s0 >> 15) & 1)
- #define S16 -((s0 >> 16) & 1)
- #define S17 -((s0 >> 17) & 1)
- #define S18 -((s0 >> 18) & 1)
- #define S19 -((s0 >> 19) & 1)
- #define S20 -((s0 >> 20) & 1)
- #define S21 -((s0 >> 21) & 1)
- #define S22 -((s0 >> 22) & 1)
- #define S23 -((s0 >> 23) & 1)
- #define S24 -((s0 >> 24) & 1)
- #define S25 -((s0 >> 25) & 1)
- #define S26 -((s0 >> 26) & 1)
- #define S27 -((s0 >> 27) & 1)
- #define S28 -((s0 >> 28) & 1)
- #define S29 -((s0 >> 29) & 1)
- #define S30 -((s0 >> 30) & 1)
- #define S31 -((s0 >> 31) & 1)
- #define S32 -((s1 >> 0) & 1)
- #define S33 -((s1 >> 1) & 1)
- #define S34 -((s1 >> 2) & 1)
- #define S35 -((s1 >> 3) & 1)
- #define S36 -((s1 >> 4) & 1)
- #define S37 -((s1 >> 5) & 1)
- #define S38 -((s1 >> 6) & 1)
- #define S39 -((s1 >> 7) & 1)
- #define S40 -((s1 >> 8) & 1)
- #define S41 -((s1 >> 9) & 1)
- #define S42 -((s1 >> 10) & 1)
- #define S43 -((s1 >> 11) & 1)
- #define S44 -((s1 >> 12) & 1)
- #define S45 -((s1 >> 13) & 1)
- #define S46 -((s1 >> 14) & 1)
- #define S47 -((s1 >> 15) & 1)
- #define S48 -((s1 >> 16) & 1)
- #define S49 -((s1 >> 17) & 1)
- #define S50 -((s1 >> 18) & 1)
- #define S51 -((s1 >> 19) & 1)
- #define S52 -((s1 >> 20) & 1)
- #define S53 -((s1 >> 21) & 1)
- #define S54 -((s1 >> 22) & 1)
- #define S55 -((s1 >> 23) & 1)
- #define S56 -((s1 >> 24) & 1)
- #define S57 -((s1 >> 25) & 1)
- #define S58 -((s1 >> 26) & 1)
- #define S59 -((s1 >> 27) & 1)
- #define S60 -((s1 >> 28) & 1)
- #define S61 -((s1 >> 29) & 1)
- #define S62 -((s1 >> 30) & 1)
- #define S63 -((s1 >> 31) & 1)
+ #define S00 (((s0 >> 0) & 1) ? -1 : 0)
+ #define S01 (((s0 >> 1) & 1) ? -1 : 0)
+ #define S02 (((s0 >> 2) & 1) ? -1 : 0)
+ #define S03 (((s0 >> 3) & 1) ? -1 : 0)
+ #define S04 (((s0 >> 4) & 1) ? -1 : 0)
+ #define S05 (((s0 >> 5) & 1) ? -1 : 0)
+ #define S06 (((s0 >> 6) & 1) ? -1 : 0)
+ #define S07 (((s0 >> 7) & 1) ? -1 : 0)
+ #define S08 (((s0 >> 8) & 1) ? -1 : 0)
+ #define S09 (((s0 >> 9) & 1) ? -1 : 0)
+ #define S10 (((s0 >> 10) & 1) ? -1 : 0)
+ #define S11 (((s0 >> 11) & 1) ? -1 : 0)
+ #define S12 (((s0 >> 12) & 1) ? -1 : 0)
+ #define S13 (((s0 >> 13) & 1) ? -1 : 0)
+ #define S14 (((s0 >> 14) & 1) ? -1 : 0)
+ #define S15 (((s0 >> 15) & 1) ? -1 : 0)
+ #define S16 (((s0 >> 16) & 1) ? -1 : 0)
+ #define S17 (((s0 >> 17) & 1) ? -1 : 0)
+ #define S18 (((s0 >> 18) & 1) ? -1 : 0)
+ #define S19 (((s0 >> 19) & 1) ? -1 : 0)
+ #define S20 (((s0 >> 20) & 1) ? -1 : 0)
+ #define S21 (((s0 >> 21) & 1) ? -1 : 0)
+ #define S22 (((s0 >> 22) & 1) ? -1 : 0)
+ #define S23 (((s0 >> 23) & 1) ? -1 : 0)
+ #define S24 (((s0 >> 24) & 1) ? -1 : 0)
+ #define S25 (((s0 >> 25) & 1) ? -1 : 0)
+ #define S26 (((s0 >> 26) & 1) ? -1 : 0)
+ #define S27 (((s0 >> 27) & 1) ? -1 : 0)
+ #define S28 (((s0 >> 28) & 1) ? -1 : 0)
+ #define S29 (((s0 >> 29) & 1) ? -1 : 0)
+ #define S30 (((s0 >> 30) & 1) ? -1 : 0)
+ #define S31 (((s0 >> 31) & 1) ? -1 : 0)
+ #define S32 (((s1 >> 0) & 1) ? -1 : 0)
+ #define S33 (((s1 >> 1) & 1) ? -1 : 0)
+ #define S34 (((s1 >> 2) & 1) ? -1 : 0)
+ #define S35 (((s1 >> 3) & 1) ? -1 : 0)
+ #define S36 (((s1 >> 4) & 1) ? -1 : 0)
+ #define S37 (((s1 >> 5) & 1) ? -1 : 0)
+ #define S38 (((s1 >> 6) & 1) ? -1 : 0)
+ #define S39 (((s1 >> 7) & 1) ? -1 : 0)
+ #define S40 (((s1 >> 8) & 1) ? -1 : 0)
+ #define S41 (((s1 >> 9) & 1) ? -1 : 0)
+ #define S42 (((s1 >> 10) & 1) ? -1 : 0)
+ #define S43 (((s1 >> 11) & 1) ? -1 : 0)
+ #define S44 (((s1 >> 12) & 1) ? -1 : 0)
+ #define S45 (((s1 >> 13) & 1) ? -1 : 0)
+ #define S46 (((s1 >> 14) & 1) ? -1 : 0)
+ #define S47 (((s1 >> 15) & 1) ? -1 : 0)
+ #define S48 (((s1 >> 16) & 1) ? -1 : 0)
+ #define S49 (((s1 >> 17) & 1) ? -1 : 0)
+ #define S50 (((s1 >> 18) & 1) ? -1 : 0)
+ #define S51 (((s1 >> 19) & 1) ? -1 : 0)
+ #define S52 (((s1 >> 20) & 1) ? -1 : 0)
+ #define S53 (((s1 >> 21) & 1) ? -1 : 0)
+ #define S54 (((s1 >> 22) & 1) ? -1 : 0)
+ #define S55 (((s1 >> 23) & 1) ? -1 : 0)
+ #define S56 (((s1 >> 24) & 1) ? -1 : 0)
+ #define S57 (((s1 >> 25) & 1) ? -1 : 0)
+ #define S58 (((s1 >> 26) & 1) ? -1 : 0)
+ #define S59 (((s1 >> 27) & 1) ? -1 : 0)
+ #define S60 (((s1 >> 28) & 1) ? -1 : 0)
+ #define S61 (((s1 >> 29) & 1) ? -1 : 0)
+ #define S62 (((s1 >> 30) & 1) ? -1 : 0)
+ #define S63 (((s1 >> 31) & 1) ? -1 : 0)
/**
* base
const u32 w0s = pws[gid].i[0];
const u32 w1s = pws[gid].i[1];
- #define K00 -((w0s >> ( 0 + 7)) & 1)
- #define K01 -((w0s >> ( 0 + 6)) & 1)
- #define K02 -((w0s >> ( 0 + 5)) & 1)
- #define K03 -((w0s >> ( 0 + 4)) & 1)
- #define K04 -((w0s >> ( 0 + 3)) & 1)
- #define K05 -((w0s >> ( 0 + 2)) & 1)
- #define K06 -((w0s >> ( 0 + 1)) & 1)
- #define K07 -((w0s >> ( 0 + 0)) & 1)
- #define K08 -((w0s >> ( 8 + 7)) & 1)
- #define K09 -((w0s >> ( 8 + 6)) & 1)
- #define K10 -((w0s >> ( 8 + 5)) & 1)
- #define K11 -((w0s >> ( 8 + 4)) & 1)
- #define K12 -((w0s >> ( 8 + 3)) & 1)
- #define K13 -((w0s >> ( 8 + 2)) & 1)
- #define K14 -((w0s >> ( 8 + 1)) & 1)
- #define K15 -((w0s >> ( 8 + 0)) & 1)
- #define K16 -((w0s >> (16 + 7)) & 1)
- #define K17 -((w0s >> (16 + 6)) & 1)
- #define K18 -((w0s >> (16 + 5)) & 1)
- #define K19 -((w0s >> (16 + 4)) & 1)
- #define K20 -((w0s >> (16 + 3)) & 1)
- #define K21 -((w0s >> (16 + 2)) & 1)
- #define K22 -((w0s >> (16 + 1)) & 1)
- #define K23 -((w0s >> (16 + 0)) & 1)
- #define K24 -((w0s >> (24 + 7)) & 1)
- #define K25 -((w0s >> (24 + 6)) & 1)
- #define K26 -((w0s >> (24 + 5)) & 1)
- #define K27 -((w0s >> (24 + 4)) & 1)
- #define K28 -((w0s >> (24 + 3)) & 1)
- #define K29 -((w0s >> (24 + 2)) & 1)
- #define K30 -((w0s >> (24 + 1)) & 1)
- #define K31 -((w0s >> (24 + 0)) & 1)
- #define K32 -((w1s >> ( 0 + 7)) & 1)
- #define K33 -((w1s >> ( 0 + 6)) & 1)
- #define K34 -((w1s >> ( 0 + 5)) & 1)
- #define K35 -((w1s >> ( 0 + 4)) & 1)
- #define K36 -((w1s >> ( 0 + 3)) & 1)
- #define K37 -((w1s >> ( 0 + 2)) & 1)
- #define K38 -((w1s >> ( 0 + 1)) & 1)
- #define K39 -((w1s >> ( 0 + 0)) & 1)
- #define K40 -((w1s >> ( 8 + 7)) & 1)
- #define K41 -((w1s >> ( 8 + 6)) & 1)
- #define K42 -((w1s >> ( 8 + 5)) & 1)
- #define K43 -((w1s >> ( 8 + 4)) & 1)
- #define K44 -((w1s >> ( 8 + 3)) & 1)
- #define K45 -((w1s >> ( 8 + 2)) & 1)
- #define K46 -((w1s >> ( 8 + 1)) & 1)
- #define K47 -((w1s >> ( 8 + 0)) & 1)
- #define K48 -((w1s >> (16 + 7)) & 1)
- #define K49 -((w1s >> (16 + 6)) & 1)
- #define K50 -((w1s >> (16 + 5)) & 1)
- #define K51 -((w1s >> (16 + 4)) & 1)
- #define K52 -((w1s >> (16 + 3)) & 1)
- #define K53 -((w1s >> (16 + 2)) & 1)
- #define K54 -((w1s >> (16 + 1)) & 1)
- #define K55 -((w1s >> (16 + 0)) & 1)
+ #define K00 (((w0s >> ( 0 + 7)) & 1) ? -1 : 0)
+ #define K01 (((w0s >> ( 0 + 6)) & 1) ? -1 : 0)
+ #define K02 (((w0s >> ( 0 + 5)) & 1) ? -1 : 0)
+ #define K03 (((w0s >> ( 0 + 4)) & 1) ? -1 : 0)
+ #define K04 (((w0s >> ( 0 + 3)) & 1) ? -1 : 0)
+ #define K05 (((w0s >> ( 0 + 2)) & 1) ? -1 : 0)
+ #define K06 (((w0s >> ( 0 + 1)) & 1) ? -1 : 0)
+ #define K07 (((w0s >> ( 0 + 0)) & 1) ? -1 : 0)
+ #define K08 (((w0s >> ( 8 + 7)) & 1) ? -1 : 0)
+ #define K09 (((w0s >> ( 8 + 6)) & 1) ? -1 : 0)
+ #define K10 (((w0s >> ( 8 + 5)) & 1) ? -1 : 0)
+ #define K11 (((w0s >> ( 8 + 4)) & 1) ? -1 : 0)
+ #define K12 (((w0s >> ( 8 + 3)) & 1) ? -1 : 0)
+ #define K13 (((w0s >> ( 8 + 2)) & 1) ? -1 : 0)
+ #define K14 (((w0s >> ( 8 + 1)) & 1) ? -1 : 0)
+ #define K15 (((w0s >> ( 8 + 0)) & 1) ? -1 : 0)
+ #define K16 (((w0s >> (16 + 7)) & 1) ? -1 : 0)
+ #define K17 (((w0s >> (16 + 6)) & 1) ? -1 : 0)
+ #define K18 (((w0s >> (16 + 5)) & 1) ? -1 : 0)
+ #define K19 (((w0s >> (16 + 4)) & 1) ? -1 : 0)
+ #define K20 (((w0s >> (16 + 3)) & 1) ? -1 : 0)
+ #define K21 (((w0s >> (16 + 2)) & 1) ? -1 : 0)
+ #define K22 (((w0s >> (16 + 1)) & 1) ? -1 : 0)
+ #define K23 (((w0s >> (16 + 0)) & 1) ? -1 : 0)
+ #define K24 (((w0s >> (24 + 7)) & 1) ? -1 : 0)
+ #define K25 (((w0s >> (24 + 6)) & 1) ? -1 : 0)
+ #define K26 (((w0s >> (24 + 5)) & 1) ? -1 : 0)
+ #define K27 (((w0s >> (24 + 4)) & 1) ? -1 : 0)
+ #define K28 (((w0s >> (24 + 3)) & 1) ? -1 : 0)
+ #define K29 (((w0s >> (24 + 2)) & 1) ? -1 : 0)
+ #define K30 (((w0s >> (24 + 1)) & 1) ? -1 : 0)
+ #define K31 (((w0s >> (24 + 0)) & 1) ? -1 : 0)
+ #define K32 (((w1s >> ( 0 + 7)) & 1) ? -1 : 0)
+ #define K33 (((w1s >> ( 0 + 6)) & 1) ? -1 : 0)
+ #define K34 (((w1s >> ( 0 + 5)) & 1) ? -1 : 0)
+ #define K35 (((w1s >> ( 0 + 4)) & 1) ? -1 : 0)
+ #define K36 (((w1s >> ( 0 + 3)) & 1) ? -1 : 0)
+ #define K37 (((w1s >> ( 0 + 2)) & 1) ? -1 : 0)
+ #define K38 (((w1s >> ( 0 + 1)) & 1) ? -1 : 0)
+ #define K39 (((w1s >> ( 0 + 0)) & 1) ? -1 : 0)
+ #define K40 (((w1s >> ( 8 + 7)) & 1) ? -1 : 0)
+ #define K41 (((w1s >> ( 8 + 6)) & 1) ? -1 : 0)
+ #define K42 (((w1s >> ( 8 + 5)) & 1) ? -1 : 0)
+ #define K43 (((w1s >> ( 8 + 4)) & 1) ? -1 : 0)
+ #define K44 (((w1s >> ( 8 + 3)) & 1) ? -1 : 0)
+ #define K45 (((w1s >> ( 8 + 2)) & 1) ? -1 : 0)
+ #define K46 (((w1s >> ( 8 + 1)) & 1) ? -1 : 0)
+ #define K47 (((w1s >> ( 8 + 0)) & 1) ? -1 : 0)
+ #define K48 (((w1s >> (16 + 7)) & 1) ? -1 : 0)
+ #define K49 (((w1s >> (16 + 6)) & 1) ? -1 : 0)
+ #define K50 (((w1s >> (16 + 5)) & 1) ? -1 : 0)
+ #define K51 (((w1s >> (16 + 4)) & 1) ? -1 : 0)
+ #define K52 (((w1s >> (16 + 3)) & 1) ? -1 : 0)
+ #define K53 (((w1s >> (16 + 2)) & 1) ? -1 : 0)
+ #define K54 (((w1s >> (16 + 1)) & 1) ? -1 : 0)
+ #define K55 (((w1s >> (16 + 0)) & 1) ? -1 : 0)
/**
* inner loop
}
}
-__kernel void m03000_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * 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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+__kernel void m03000_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * 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_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
/**
* base
* main
*/
- m03000m (pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset);
+ m03000m (pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV0_buf, d_scryptV1_buf, d_scryptV2_buf, d_scryptV3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset);
}
-__kernel void m03000_m08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * 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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+__kernel void m03000_m08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * 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_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
}
-__kernel void m03000_m16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * 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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+__kernel void m03000_m16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * 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_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
}
-__kernel void m03000_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * 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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+__kernel void m03000_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * 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_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
/**
* base
* main
*/
- m03000s (pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset);
+ m03000s (pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV0_buf, d_scryptV1_buf, d_scryptV2_buf, d_scryptV3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset);
}
-__kernel void m03000_s08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * 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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+__kernel void m03000_s08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * 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_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
}
-__kernel void m03000_s16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * 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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+__kernel void m03000_s16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bs_word_t * 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_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
}
projects / hashcat.git / blobdiff