diff --git a/OpenCL/m00200_a0.cl b/OpenCL/m00200_a0-optimized.cl similarity index 100% rename from OpenCL/m00200_a0.cl rename to OpenCL/m00200_a0-optimized.cl diff --git a/OpenCL/m00200_a1.cl b/OpenCL/m00200_a1-optimized.cl similarity index 100% rename from OpenCL/m00200_a1.cl rename to OpenCL/m00200_a1-optimized.cl diff --git a/OpenCL/m00200_a3.cl b/OpenCL/m00200_a3-optimized.cl similarity index 100% rename from OpenCL/m00200_a3.cl rename to OpenCL/m00200_a3-optimized.cl diff --git a/OpenCL/m00600_a0.cl b/OpenCL/m00600_a0-optimized.cl similarity index 100% rename from OpenCL/m00600_a0.cl rename to OpenCL/m00600_a0-optimized.cl diff --git a/OpenCL/m00600_a1.cl b/OpenCL/m00600_a1-optimized.cl similarity index 100% rename from OpenCL/m00600_a1.cl rename to OpenCL/m00600_a1-optimized.cl diff --git a/OpenCL/m00600_a3.cl b/OpenCL/m00600_a3-optimized.cl similarity index 100% rename from OpenCL/m00600_a3.cl rename to OpenCL/m00600_a3-optimized.cl diff --git a/OpenCL/m01500_a0.cl b/OpenCL/m01500_a0.cl index 960476517..aca153b0e 100644 --- a/OpenCL/m01500_a0.cl +++ b/OpenCL/m01500_a0.cl @@ -3,7 +3,7 @@ * License.....: MIT */ -#define NEW_SIMD_CODE +//#define NEW_SIMD_CODE #include "inc_vendor.cl" #include "inc_hash_constants.h" @@ -481,7 +481,7 @@ void _des_crypt_encrypt (u32x iv[2], u32 mask, u32x Kc[16], u32x Kd[16], __local iv[1] = rotl32 (l, 31); } -__kernel void m01500_m04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 m01500_mxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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) { /** * modifier @@ -580,15 +580,7 @@ __kernel void m01500_m04 (__global pw_t *pws, __global const kernel_rule_t *rule } } -__kernel void m01500_m08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 m01500_m16 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 m01500_s04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 m01500_sxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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) { /** * modifier @@ -698,11 +690,3 @@ __kernel void m01500_s04 (__global pw_t *pws, __global const kernel_rule_t *rule COMPARE_S_SIMD (iv[0], iv[1], z, z); } } - -__kernel void m01500_s08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 m01500_s16 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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) -{ -} diff --git a/OpenCL/m01500_a1.cl b/OpenCL/m01500_a1.cl index 3cd9c339f..fc85a9b4b 100644 --- a/OpenCL/m01500_a1.cl +++ b/OpenCL/m01500_a1.cl @@ -3,7 +3,7 @@ * License.....: MIT */ -#define NEW_SIMD_CODE +//#define NEW_SIMD_CODE #include "inc_vendor.cl" #include "inc_hash_constants.h" @@ -479,7 +479,7 @@ void _des_crypt_encrypt (u32x iv[2], u32 mask, u32x Kc[16], u32x Kd[16], __local iv[1] = rotl32 (l, 31); } -__kernel void m01500_m04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 m01500_mxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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) { /** * modifier @@ -642,15 +642,7 @@ __kernel void m01500_m04 (__global pw_t *pws, __global const kernel_rule_t *rule } } -__kernel void m01500_m08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 m01500_m16 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 m01500_s04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 m01500_sxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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) { /** * modifier @@ -824,11 +816,3 @@ __kernel void m01500_s04 (__global pw_t *pws, __global const kernel_rule_t *rule COMPARE_S_SIMD (iv[0], iv[1], z, z); } } - -__kernel void m01500_s08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 m01500_s16 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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) -{ -} diff --git a/OpenCL/m01500_a3.cl b/OpenCL/m01500_a3.cl index a6c640069..f54007b40 100644 --- a/OpenCL/m01500_a3.cl +++ b/OpenCL/m01500_a3.cl @@ -2642,7 +2642,7 @@ __kernel void m01500_tm (__global u32 *mod, __global bs_word_t *words_buf_r) } } -__kernel void m01500_m04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 m01500_mxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 @@ -2660,15 +2660,7 @@ __kernel void m01500_m04 (__global pw_t *pws, __global const kernel_rule_t *rule m01500m (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 m01500_m08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 m01500_m16 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 m01500_s04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 m01500_sxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 @@ -2685,11 +2677,3 @@ __kernel void m01500_s04 (__global pw_t *pws, __global const kernel_rule_t *rule m01500s (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 m01500_s08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 m01500_s16 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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) -{ -} diff --git a/OpenCL/m02000_a0.cl b/OpenCL/m02000_a0.cl index 0505812cd..b722fca79 100644 --- a/OpenCL/m02000_a0.cl +++ b/OpenCL/m02000_a0.cl @@ -12,26 +12,10 @@ __kernel void gpu_memset (__global uint4 *buf, const uint value, const uint gid_ buf[gid] = (uint4) (value); } -__kernel void m02000_m04 (__global void *pws, __global void *rules_buf, __global void *combs_buf, __global void * words_buf_r, __global void *tmps, __global void *hooks, __global void *bitmaps_buf_s1_a, __global void *bitmaps_buf_s1_b, __global void *bitmaps_buf_s1_c, __global void *bitmaps_buf_s1_d, __global void *bitmaps_buf_s2_a, __global void *bitmaps_buf_s2_b, __global void *bitmaps_buf_s2_c, __global void *bitmaps_buf_s2_d, __global void *plains_buf, __global void *digests_buf, __global void *hashes_shown, __global void *salt_bufs, __global const void *esalt_bufs, __global void *d_return_buf, __global void *d_scryptV0_buf, __global void *d_scryptV1_buf, __global void *d_scryptV2_buf, __global void *d_scryptV3_buf, const uint bitmap_mask, const uint bitmap_shift1, const uint bitmap_shift2, const uint salt_pos, const uint loop_pos, const uint loop_cnt, const uint il_cnt, const uint digests_cnt, const uint digests_offset, const uint combs_mode, const uint gid_max) +__kernel void m02000_mxx (__global void *pws, __global void *rules_buf, __global void *combs_buf, __global void * words_buf_r, __global void *tmps, __global void *hooks, __global void *bitmaps_buf_s1_a, __global void *bitmaps_buf_s1_b, __global void *bitmaps_buf_s1_c, __global void *bitmaps_buf_s1_d, __global void *bitmaps_buf_s2_a, __global void *bitmaps_buf_s2_b, __global void *bitmaps_buf_s2_c, __global void *bitmaps_buf_s2_d, __global void *plains_buf, __global void *digests_buf, __global void *hashes_shown, __global void *salt_bufs, __global const void *esalt_bufs, __global void *d_return_buf, __global void *d_scryptV0_buf, __global void *d_scryptV1_buf, __global void *d_scryptV2_buf, __global void *d_scryptV3_buf, const uint bitmap_mask, const uint bitmap_shift1, const uint bitmap_shift2, const uint salt_pos, const uint loop_pos, const uint loop_cnt, const uint il_cnt, const uint digests_cnt, const uint digests_offset, const uint combs_mode, const uint gid_max) { } -__kernel void m02000_m08 (__global void *pws, __global void *rules_buf, __global void *combs_buf, __global void * words_buf_r, __global void *tmps, __global void *hooks, __global void *bitmaps_buf_s1_a, __global void *bitmaps_buf_s1_b, __global void *bitmaps_buf_s1_c, __global void *bitmaps_buf_s1_d, __global void *bitmaps_buf_s2_a, __global void *bitmaps_buf_s2_b, __global void *bitmaps_buf_s2_c, __global void *bitmaps_buf_s2_d, __global void *plains_buf, __global void *digests_buf, __global void *hashes_shown, __global void *salt_bufs, __global const void *esalt_bufs, __global void *d_return_buf, __global void *d_scryptV0_buf, __global void *d_scryptV1_buf, __global void *d_scryptV2_buf, __global void *d_scryptV3_buf, const uint bitmap_mask, const uint bitmap_shift1, const uint bitmap_shift2, const uint salt_pos, const uint loop_pos, const uint loop_cnt, const uint il_cnt, const uint digests_cnt, const uint digests_offset, const uint combs_mode, const uint gid_max) -{ -} - -__kernel void m02000_m16 (__global void *pws, __global void *rules_buf, __global void *combs_buf, __global void * words_buf_r, __global void *tmps, __global void *hooks, __global void *bitmaps_buf_s1_a, __global void *bitmaps_buf_s1_b, __global void *bitmaps_buf_s1_c, __global void *bitmaps_buf_s1_d, __global void *bitmaps_buf_s2_a, __global void *bitmaps_buf_s2_b, __global void *bitmaps_buf_s2_c, __global void *bitmaps_buf_s2_d, __global void *plains_buf, __global void *digests_buf, __global void *hashes_shown, __global void *salt_bufs, __global const void *esalt_bufs, __global void *d_return_buf, __global void *d_scryptV0_buf, __global void *d_scryptV1_buf, __global void *d_scryptV2_buf, __global void *d_scryptV3_buf, const uint bitmap_mask, const uint bitmap_shift1, const uint bitmap_shift2, const uint salt_pos, const uint loop_pos, const uint loop_cnt, const uint il_cnt, const uint digests_cnt, const uint digests_offset, const uint combs_mode, const uint gid_max) -{ -} - -__kernel void m02000_s04 (__global void *pws, __global void *rules_buf, __global void *combs_buf, __global void * words_buf_r, __global void *tmps, __global void *hooks, __global void *bitmaps_buf_s1_a, __global void *bitmaps_buf_s1_b, __global void *bitmaps_buf_s1_c, __global void *bitmaps_buf_s1_d, __global void *bitmaps_buf_s2_a, __global void *bitmaps_buf_s2_b, __global void *bitmaps_buf_s2_c, __global void *bitmaps_buf_s2_d, __global void *plains_buf, __global void *digests_buf, __global void *hashes_shown, __global void *salt_bufs, __global const void *esalt_bufs, __global void *d_return_buf, __global void *d_scryptV0_buf, __global void *d_scryptV1_buf, __global void *d_scryptV2_buf, __global void *d_scryptV3_buf, const uint bitmap_mask, const uint bitmap_shift1, const uint bitmap_shift2, const uint salt_pos, const uint loop_pos, const uint loop_cnt, const uint il_cnt, const uint digests_cnt, const uint digests_offset, const uint combs_mode, const uint gid_max) -{ -} - -__kernel void m02000_s08 (__global void *pws, __global void *rules_buf, __global void *combs_buf, __global void * words_buf_r, __global void *tmps, __global void *hooks, __global void *bitmaps_buf_s1_a, __global void *bitmaps_buf_s1_b, __global void *bitmaps_buf_s1_c, __global void *bitmaps_buf_s1_d, __global void *bitmaps_buf_s2_a, __global void *bitmaps_buf_s2_b, __global void *bitmaps_buf_s2_c, __global void *bitmaps_buf_s2_d, __global void *plains_buf, __global void *digests_buf, __global void *hashes_shown, __global void *salt_bufs, __global const void *esalt_bufs, __global void *d_return_buf, __global void *d_scryptV0_buf, __global void *d_scryptV1_buf, __global void *d_scryptV2_buf, __global void *d_scryptV3_buf, const uint bitmap_mask, const uint bitmap_shift1, const uint bitmap_shift2, const uint salt_pos, const uint loop_pos, const uint loop_cnt, const uint il_cnt, const uint digests_cnt, const uint digests_offset, const uint combs_mode, const uint gid_max) -{ -} - -__kernel void m02000_s16 (__global void *pws, __global void *rules_buf, __global void *combs_buf, __global void * words_buf_r, __global void *tmps, __global void *hooks, __global void *bitmaps_buf_s1_a, __global void *bitmaps_buf_s1_b, __global void *bitmaps_buf_s1_c, __global void *bitmaps_buf_s1_d, __global void *bitmaps_buf_s2_a, __global void *bitmaps_buf_s2_b, __global void *bitmaps_buf_s2_c, __global void *bitmaps_buf_s2_d, __global void *plains_buf, __global void *digests_buf, __global void *hashes_shown, __global void *salt_bufs, __global const void *esalt_bufs, __global void *d_return_buf, __global void *d_scryptV0_buf, __global void *d_scryptV1_buf, __global void *d_scryptV2_buf, __global void *d_scryptV3_buf, const uint bitmap_mask, const uint bitmap_shift1, const uint bitmap_shift2, const uint salt_pos, const uint loop_pos, const uint loop_cnt, const uint il_cnt, const uint digests_cnt, const uint digests_offset, const uint combs_mode, const uint gid_max) +__kernel void m02000_sxx (__global void *pws, __global void *rules_buf, __global void *combs_buf, __global void * words_buf_r, __global void *tmps, __global void *hooks, __global void *bitmaps_buf_s1_a, __global void *bitmaps_buf_s1_b, __global void *bitmaps_buf_s1_c, __global void *bitmaps_buf_s1_d, __global void *bitmaps_buf_s2_a, __global void *bitmaps_buf_s2_b, __global void *bitmaps_buf_s2_c, __global void *bitmaps_buf_s2_d, __global void *plains_buf, __global void *digests_buf, __global void *hashes_shown, __global void *salt_bufs, __global const void *esalt_bufs, __global void *d_return_buf, __global void *d_scryptV0_buf, __global void *d_scryptV1_buf, __global void *d_scryptV2_buf, __global void *d_scryptV3_buf, const uint bitmap_mask, const uint bitmap_shift1, const uint bitmap_shift2, const uint salt_pos, const uint loop_pos, const uint loop_cnt, const uint il_cnt, const uint digests_cnt, const uint digests_offset, const uint combs_mode, const uint gid_max) { } diff --git a/OpenCL/m02000_a1.cl b/OpenCL/m02000_a1.cl index 0505812cd..b722fca79 100644 --- a/OpenCL/m02000_a1.cl +++ b/OpenCL/m02000_a1.cl @@ -12,26 +12,10 @@ __kernel void gpu_memset (__global uint4 *buf, const uint value, const uint gid_ buf[gid] = (uint4) (value); } -__kernel void m02000_m04 (__global void *pws, __global void *rules_buf, __global void *combs_buf, __global void * words_buf_r, __global void *tmps, __global void *hooks, __global void *bitmaps_buf_s1_a, __global void *bitmaps_buf_s1_b, __global void *bitmaps_buf_s1_c, __global void *bitmaps_buf_s1_d, __global void *bitmaps_buf_s2_a, __global void *bitmaps_buf_s2_b, __global void *bitmaps_buf_s2_c, __global void *bitmaps_buf_s2_d, __global void *plains_buf, __global void *digests_buf, __global void *hashes_shown, __global void *salt_bufs, __global const void *esalt_bufs, __global void *d_return_buf, __global void *d_scryptV0_buf, __global void *d_scryptV1_buf, __global void *d_scryptV2_buf, __global void *d_scryptV3_buf, const uint bitmap_mask, const uint bitmap_shift1, const uint bitmap_shift2, const uint salt_pos, const uint loop_pos, const uint loop_cnt, const uint il_cnt, const uint digests_cnt, const uint digests_offset, const uint combs_mode, const uint gid_max) +__kernel void m02000_mxx (__global void *pws, __global void *rules_buf, __global void *combs_buf, __global void * words_buf_r, __global void *tmps, __global void *hooks, __global void *bitmaps_buf_s1_a, __global void *bitmaps_buf_s1_b, __global void *bitmaps_buf_s1_c, __global void *bitmaps_buf_s1_d, __global void *bitmaps_buf_s2_a, __global void *bitmaps_buf_s2_b, __global void *bitmaps_buf_s2_c, __global void *bitmaps_buf_s2_d, __global void *plains_buf, __global void *digests_buf, __global void *hashes_shown, __global void *salt_bufs, __global const void *esalt_bufs, __global void *d_return_buf, __global void *d_scryptV0_buf, __global void *d_scryptV1_buf, __global void *d_scryptV2_buf, __global void *d_scryptV3_buf, const uint bitmap_mask, const uint bitmap_shift1, const uint bitmap_shift2, const uint salt_pos, const uint loop_pos, const uint loop_cnt, const uint il_cnt, const uint digests_cnt, const uint digests_offset, const uint combs_mode, const uint gid_max) { } -__kernel void m02000_m08 (__global void *pws, __global void *rules_buf, __global void *combs_buf, __global void * words_buf_r, __global void *tmps, __global void *hooks, __global void *bitmaps_buf_s1_a, __global void *bitmaps_buf_s1_b, __global void *bitmaps_buf_s1_c, __global void *bitmaps_buf_s1_d, __global void *bitmaps_buf_s2_a, __global void *bitmaps_buf_s2_b, __global void *bitmaps_buf_s2_c, __global void *bitmaps_buf_s2_d, __global void *plains_buf, __global void *digests_buf, __global void *hashes_shown, __global void *salt_bufs, __global const void *esalt_bufs, __global void *d_return_buf, __global void *d_scryptV0_buf, __global void *d_scryptV1_buf, __global void *d_scryptV2_buf, __global void *d_scryptV3_buf, const uint bitmap_mask, const uint bitmap_shift1, const uint bitmap_shift2, const uint salt_pos, const uint loop_pos, const uint loop_cnt, const uint il_cnt, const uint digests_cnt, const uint digests_offset, const uint combs_mode, const uint gid_max) -{ -} - -__kernel void m02000_m16 (__global void *pws, __global void *rules_buf, __global void *combs_buf, __global void * words_buf_r, __global void *tmps, __global void *hooks, __global void *bitmaps_buf_s1_a, __global void *bitmaps_buf_s1_b, __global void *bitmaps_buf_s1_c, __global void *bitmaps_buf_s1_d, __global void *bitmaps_buf_s2_a, __global void *bitmaps_buf_s2_b, __global void *bitmaps_buf_s2_c, __global void *bitmaps_buf_s2_d, __global void *plains_buf, __global void *digests_buf, __global void *hashes_shown, __global void *salt_bufs, __global const void *esalt_bufs, __global void *d_return_buf, __global void *d_scryptV0_buf, __global void *d_scryptV1_buf, __global void *d_scryptV2_buf, __global void *d_scryptV3_buf, const uint bitmap_mask, const uint bitmap_shift1, const uint bitmap_shift2, const uint salt_pos, const uint loop_pos, const uint loop_cnt, const uint il_cnt, const uint digests_cnt, const uint digests_offset, const uint combs_mode, const uint gid_max) -{ -} - -__kernel void m02000_s04 (__global void *pws, __global void *rules_buf, __global void *combs_buf, __global void * words_buf_r, __global void *tmps, __global void *hooks, __global void *bitmaps_buf_s1_a, __global void *bitmaps_buf_s1_b, __global void *bitmaps_buf_s1_c, __global void *bitmaps_buf_s1_d, __global void *bitmaps_buf_s2_a, __global void *bitmaps_buf_s2_b, __global void *bitmaps_buf_s2_c, __global void *bitmaps_buf_s2_d, __global void *plains_buf, __global void *digests_buf, __global void *hashes_shown, __global void *salt_bufs, __global const void *esalt_bufs, __global void *d_return_buf, __global void *d_scryptV0_buf, __global void *d_scryptV1_buf, __global void *d_scryptV2_buf, __global void *d_scryptV3_buf, const uint bitmap_mask, const uint bitmap_shift1, const uint bitmap_shift2, const uint salt_pos, const uint loop_pos, const uint loop_cnt, const uint il_cnt, const uint digests_cnt, const uint digests_offset, const uint combs_mode, const uint gid_max) -{ -} - -__kernel void m02000_s08 (__global void *pws, __global void *rules_buf, __global void *combs_buf, __global void * words_buf_r, __global void *tmps, __global void *hooks, __global void *bitmaps_buf_s1_a, __global void *bitmaps_buf_s1_b, __global void *bitmaps_buf_s1_c, __global void *bitmaps_buf_s1_d, __global void *bitmaps_buf_s2_a, __global void *bitmaps_buf_s2_b, __global void *bitmaps_buf_s2_c, __global void *bitmaps_buf_s2_d, __global void *plains_buf, __global void *digests_buf, __global void *hashes_shown, __global void *salt_bufs, __global const void *esalt_bufs, __global void *d_return_buf, __global void *d_scryptV0_buf, __global void *d_scryptV1_buf, __global void *d_scryptV2_buf, __global void *d_scryptV3_buf, const uint bitmap_mask, const uint bitmap_shift1, const uint bitmap_shift2, const uint salt_pos, const uint loop_pos, const uint loop_cnt, const uint il_cnt, const uint digests_cnt, const uint digests_offset, const uint combs_mode, const uint gid_max) -{ -} - -__kernel void m02000_s16 (__global void *pws, __global void *rules_buf, __global void *combs_buf, __global void * words_buf_r, __global void *tmps, __global void *hooks, __global void *bitmaps_buf_s1_a, __global void *bitmaps_buf_s1_b, __global void *bitmaps_buf_s1_c, __global void *bitmaps_buf_s1_d, __global void *bitmaps_buf_s2_a, __global void *bitmaps_buf_s2_b, __global void *bitmaps_buf_s2_c, __global void *bitmaps_buf_s2_d, __global void *plains_buf, __global void *digests_buf, __global void *hashes_shown, __global void *salt_bufs, __global const void *esalt_bufs, __global void *d_return_buf, __global void *d_scryptV0_buf, __global void *d_scryptV1_buf, __global void *d_scryptV2_buf, __global void *d_scryptV3_buf, const uint bitmap_mask, const uint bitmap_shift1, const uint bitmap_shift2, const uint salt_pos, const uint loop_pos, const uint loop_cnt, const uint il_cnt, const uint digests_cnt, const uint digests_offset, const uint combs_mode, const uint gid_max) +__kernel void m02000_sxx (__global void *pws, __global void *rules_buf, __global void *combs_buf, __global void * words_buf_r, __global void *tmps, __global void *hooks, __global void *bitmaps_buf_s1_a, __global void *bitmaps_buf_s1_b, __global void *bitmaps_buf_s1_c, __global void *bitmaps_buf_s1_d, __global void *bitmaps_buf_s2_a, __global void *bitmaps_buf_s2_b, __global void *bitmaps_buf_s2_c, __global void *bitmaps_buf_s2_d, __global void *plains_buf, __global void *digests_buf, __global void *hashes_shown, __global void *salt_bufs, __global const void *esalt_bufs, __global void *d_return_buf, __global void *d_scryptV0_buf, __global void *d_scryptV1_buf, __global void *d_scryptV2_buf, __global void *d_scryptV3_buf, const uint bitmap_mask, const uint bitmap_shift1, const uint bitmap_shift2, const uint salt_pos, const uint loop_pos, const uint loop_cnt, const uint il_cnt, const uint digests_cnt, const uint digests_offset, const uint combs_mode, const uint gid_max) { } diff --git a/OpenCL/m02000_a3.cl b/OpenCL/m02000_a3.cl index 0505812cd..b722fca79 100644 --- a/OpenCL/m02000_a3.cl +++ b/OpenCL/m02000_a3.cl @@ -12,26 +12,10 @@ __kernel void gpu_memset (__global uint4 *buf, const uint value, const uint gid_ buf[gid] = (uint4) (value); } -__kernel void m02000_m04 (__global void *pws, __global void *rules_buf, __global void *combs_buf, __global void * words_buf_r, __global void *tmps, __global void *hooks, __global void *bitmaps_buf_s1_a, __global void *bitmaps_buf_s1_b, __global void *bitmaps_buf_s1_c, __global void *bitmaps_buf_s1_d, __global void *bitmaps_buf_s2_a, __global void *bitmaps_buf_s2_b, __global void *bitmaps_buf_s2_c, __global void *bitmaps_buf_s2_d, __global void *plains_buf, __global void *digests_buf, __global void *hashes_shown, __global void *salt_bufs, __global const void *esalt_bufs, __global void *d_return_buf, __global void *d_scryptV0_buf, __global void *d_scryptV1_buf, __global void *d_scryptV2_buf, __global void *d_scryptV3_buf, const uint bitmap_mask, const uint bitmap_shift1, const uint bitmap_shift2, const uint salt_pos, const uint loop_pos, const uint loop_cnt, const uint il_cnt, const uint digests_cnt, const uint digests_offset, const uint combs_mode, const uint gid_max) +__kernel void m02000_mxx (__global void *pws, __global void *rules_buf, __global void *combs_buf, __global void * words_buf_r, __global void *tmps, __global void *hooks, __global void *bitmaps_buf_s1_a, __global void *bitmaps_buf_s1_b, __global void *bitmaps_buf_s1_c, __global void *bitmaps_buf_s1_d, __global void *bitmaps_buf_s2_a, __global void *bitmaps_buf_s2_b, __global void *bitmaps_buf_s2_c, __global void *bitmaps_buf_s2_d, __global void *plains_buf, __global void *digests_buf, __global void *hashes_shown, __global void *salt_bufs, __global const void *esalt_bufs, __global void *d_return_buf, __global void *d_scryptV0_buf, __global void *d_scryptV1_buf, __global void *d_scryptV2_buf, __global void *d_scryptV3_buf, const uint bitmap_mask, const uint bitmap_shift1, const uint bitmap_shift2, const uint salt_pos, const uint loop_pos, const uint loop_cnt, const uint il_cnt, const uint digests_cnt, const uint digests_offset, const uint combs_mode, const uint gid_max) { } -__kernel void m02000_m08 (__global void *pws, __global void *rules_buf, __global void *combs_buf, __global void * words_buf_r, __global void *tmps, __global void *hooks, __global void *bitmaps_buf_s1_a, __global void *bitmaps_buf_s1_b, __global void *bitmaps_buf_s1_c, __global void *bitmaps_buf_s1_d, __global void *bitmaps_buf_s2_a, __global void *bitmaps_buf_s2_b, __global void *bitmaps_buf_s2_c, __global void *bitmaps_buf_s2_d, __global void *plains_buf, __global void *digests_buf, __global void *hashes_shown, __global void *salt_bufs, __global const void *esalt_bufs, __global void *d_return_buf, __global void *d_scryptV0_buf, __global void *d_scryptV1_buf, __global void *d_scryptV2_buf, __global void *d_scryptV3_buf, const uint bitmap_mask, const uint bitmap_shift1, const uint bitmap_shift2, const uint salt_pos, const uint loop_pos, const uint loop_cnt, const uint il_cnt, const uint digests_cnt, const uint digests_offset, const uint combs_mode, const uint gid_max) -{ -} - -__kernel void m02000_m16 (__global void *pws, __global void *rules_buf, __global void *combs_buf, __global void * words_buf_r, __global void *tmps, __global void *hooks, __global void *bitmaps_buf_s1_a, __global void *bitmaps_buf_s1_b, __global void *bitmaps_buf_s1_c, __global void *bitmaps_buf_s1_d, __global void *bitmaps_buf_s2_a, __global void *bitmaps_buf_s2_b, __global void *bitmaps_buf_s2_c, __global void *bitmaps_buf_s2_d, __global void *plains_buf, __global void *digests_buf, __global void *hashes_shown, __global void *salt_bufs, __global const void *esalt_bufs, __global void *d_return_buf, __global void *d_scryptV0_buf, __global void *d_scryptV1_buf, __global void *d_scryptV2_buf, __global void *d_scryptV3_buf, const uint bitmap_mask, const uint bitmap_shift1, const uint bitmap_shift2, const uint salt_pos, const uint loop_pos, const uint loop_cnt, const uint il_cnt, const uint digests_cnt, const uint digests_offset, const uint combs_mode, const uint gid_max) -{ -} - -__kernel void m02000_s04 (__global void *pws, __global void *rules_buf, __global void *combs_buf, __global void * words_buf_r, __global void *tmps, __global void *hooks, __global void *bitmaps_buf_s1_a, __global void *bitmaps_buf_s1_b, __global void *bitmaps_buf_s1_c, __global void *bitmaps_buf_s1_d, __global void *bitmaps_buf_s2_a, __global void *bitmaps_buf_s2_b, __global void *bitmaps_buf_s2_c, __global void *bitmaps_buf_s2_d, __global void *plains_buf, __global void *digests_buf, __global void *hashes_shown, __global void *salt_bufs, __global const void *esalt_bufs, __global void *d_return_buf, __global void *d_scryptV0_buf, __global void *d_scryptV1_buf, __global void *d_scryptV2_buf, __global void *d_scryptV3_buf, const uint bitmap_mask, const uint bitmap_shift1, const uint bitmap_shift2, const uint salt_pos, const uint loop_pos, const uint loop_cnt, const uint il_cnt, const uint digests_cnt, const uint digests_offset, const uint combs_mode, const uint gid_max) -{ -} - -__kernel void m02000_s08 (__global void *pws, __global void *rules_buf, __global void *combs_buf, __global void * words_buf_r, __global void *tmps, __global void *hooks, __global void *bitmaps_buf_s1_a, __global void *bitmaps_buf_s1_b, __global void *bitmaps_buf_s1_c, __global void *bitmaps_buf_s1_d, __global void *bitmaps_buf_s2_a, __global void *bitmaps_buf_s2_b, __global void *bitmaps_buf_s2_c, __global void *bitmaps_buf_s2_d, __global void *plains_buf, __global void *digests_buf, __global void *hashes_shown, __global void *salt_bufs, __global const void *esalt_bufs, __global void *d_return_buf, __global void *d_scryptV0_buf, __global void *d_scryptV1_buf, __global void *d_scryptV2_buf, __global void *d_scryptV3_buf, const uint bitmap_mask, const uint bitmap_shift1, const uint bitmap_shift2, const uint salt_pos, const uint loop_pos, const uint loop_cnt, const uint il_cnt, const uint digests_cnt, const uint digests_offset, const uint combs_mode, const uint gid_max) -{ -} - -__kernel void m02000_s16 (__global void *pws, __global void *rules_buf, __global void *combs_buf, __global void * words_buf_r, __global void *tmps, __global void *hooks, __global void *bitmaps_buf_s1_a, __global void *bitmaps_buf_s1_b, __global void *bitmaps_buf_s1_c, __global void *bitmaps_buf_s1_d, __global void *bitmaps_buf_s2_a, __global void *bitmaps_buf_s2_b, __global void *bitmaps_buf_s2_c, __global void *bitmaps_buf_s2_d, __global void *plains_buf, __global void *digests_buf, __global void *hashes_shown, __global void *salt_bufs, __global const void *esalt_bufs, __global void *d_return_buf, __global void *d_scryptV0_buf, __global void *d_scryptV1_buf, __global void *d_scryptV2_buf, __global void *d_scryptV3_buf, const uint bitmap_mask, const uint bitmap_shift1, const uint bitmap_shift2, const uint salt_pos, const uint loop_pos, const uint loop_cnt, const uint il_cnt, const uint digests_cnt, const uint digests_offset, const uint combs_mode, const uint gid_max) +__kernel void m02000_sxx (__global void *pws, __global void *rules_buf, __global void *combs_buf, __global void * words_buf_r, __global void *tmps, __global void *hooks, __global void *bitmaps_buf_s1_a, __global void *bitmaps_buf_s1_b, __global void *bitmaps_buf_s1_c, __global void *bitmaps_buf_s1_d, __global void *bitmaps_buf_s2_a, __global void *bitmaps_buf_s2_b, __global void *bitmaps_buf_s2_c, __global void *bitmaps_buf_s2_d, __global void *plains_buf, __global void *digests_buf, __global void *hashes_shown, __global void *salt_bufs, __global const void *esalt_bufs, __global void *d_return_buf, __global void *d_scryptV0_buf, __global void *d_scryptV1_buf, __global void *d_scryptV2_buf, __global void *d_scryptV3_buf, const uint bitmap_mask, const uint bitmap_shift1, const uint bitmap_shift2, const uint salt_pos, const uint loop_pos, const uint loop_cnt, const uint il_cnt, const uint digests_cnt, const uint digests_offset, const uint combs_mode, const uint gid_max) { } diff --git a/OpenCL/m02400_a0.cl b/OpenCL/m02400_a0.cl index 2269d8bc8..9bb366b2d 100644 --- a/OpenCL/m02400_a0.cl +++ b/OpenCL/m02400_a0.cl @@ -3,7 +3,7 @@ * License.....: MIT */ -#define NEW_SIMD_CODE +//#define NEW_SIMD_CODE #include "inc_vendor.cl" #include "inc_hash_constants.h" @@ -14,7 +14,7 @@ #include "inc_rp.cl" #include "inc_simd.cl" -__kernel void m02400_m04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 m02400_mxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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) { /** * modifier @@ -152,15 +152,7 @@ __kernel void m02400_m04 (__global pw_t *pws, __global const kernel_rule_t *rule } } -__kernel void m02400_m08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 m02400_m16 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 m02400_s04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 m02400_sxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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) { /** * modifier @@ -312,11 +304,3 @@ __kernel void m02400_s04 (__global pw_t *pws, __global const kernel_rule_t *rule COMPARE_S_SIMD (a, d, c, b); } } - -__kernel void m02400_s08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 m02400_s16 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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) -{ -} diff --git a/OpenCL/m02400_a1.cl b/OpenCL/m02400_a1.cl index a33b08d0b..379fc18fb 100644 --- a/OpenCL/m02400_a1.cl +++ b/OpenCL/m02400_a1.cl @@ -3,7 +3,7 @@ * License.....: MIT */ -#define NEW_SIMD_CODE +//#define NEW_SIMD_CODE #include "inc_vendor.cl" #include "inc_hash_constants.h" @@ -12,7 +12,7 @@ #include "inc_common.cl" #include "inc_simd.cl" -__kernel void m02400_m04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 m02400_mxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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) { /** * modifier @@ -214,15 +214,7 @@ __kernel void m02400_m04 (__global pw_t *pws, __global const kernel_rule_t *rule } } -__kernel void m02400_m08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 m02400_m16 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 m02400_s04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 m02400_sxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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) { /** * modifier @@ -438,11 +430,3 @@ __kernel void m02400_s04 (__global pw_t *pws, __global const kernel_rule_t *rule COMPARE_S_SIMD (a, d, c, b); } } - -__kernel void m02400_s08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 m02400_s16 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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) -{ -} diff --git a/OpenCL/m02400_a3.cl b/OpenCL/m02400_a3.cl index 539b79acc..19a85de3c 100644 --- a/OpenCL/m02400_a3.cl +++ b/OpenCL/m02400_a3.cl @@ -411,7 +411,7 @@ void m02400s (u32 w[16], const u32 pw_len, __global pw_t *pws, __global const ke } } -__kernel void m02400_m04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 m02400_mxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 @@ -449,15 +449,7 @@ __kernel void m02400_m04 (__global pw_t *pws, __global const kernel_rule_t *rule m02400m (w, pw_len, 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 m02400_m08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 m02400_m16 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 m02400_s04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 m02400_sxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 @@ -494,11 +486,3 @@ __kernel void m02400_s04 (__global pw_t *pws, __global const kernel_rule_t *rule m02400s (w, pw_len, 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 m02400_s08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 m02400_s16 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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) -{ -} diff --git a/OpenCL/m02410_a0.cl b/OpenCL/m02410_a0.cl index 169f87e1a..3344f5be6 100644 --- a/OpenCL/m02410_a0.cl +++ b/OpenCL/m02410_a0.cl @@ -3,7 +3,7 @@ * License.....: MIT */ -#define NEW_SIMD_CODE +//#define NEW_SIMD_CODE #include "inc_vendor.cl" #include "inc_hash_constants.h" @@ -14,7 +14,7 @@ #include "inc_rp.cl" #include "inc_simd.cl" -__kernel void m02410_m04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 m02410_mxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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) { /** * modifier @@ -227,15 +227,7 @@ __kernel void m02410_m04 (__global pw_t *pws, __global const kernel_rule_t *rule } } -__kernel void m02410_m08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 m02410_m16 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 m02410_s04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 m02410_sxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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) { /** * modifier @@ -462,11 +454,3 @@ __kernel void m02410_s04 (__global pw_t *pws, __global const kernel_rule_t *rule COMPARE_S_SIMD (a, d, c, b); } } - -__kernel void m02410_s08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 m02410_s16 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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) -{ -} diff --git a/OpenCL/m02410_a1.cl b/OpenCL/m02410_a1.cl index accaacdfa..950e67e69 100644 --- a/OpenCL/m02410_a1.cl +++ b/OpenCL/m02410_a1.cl @@ -3,7 +3,7 @@ * License.....: MIT */ -#define NEW_SIMD_CODE +//#define NEW_SIMD_CODE #include "inc_vendor.cl" #include "inc_hash_constants.h" @@ -12,7 +12,7 @@ #include "inc_common.cl" #include "inc_simd.cl" -__kernel void m02410_m04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 m02410_mxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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) { /** * modifier @@ -285,15 +285,7 @@ __kernel void m02410_m04 (__global pw_t *pws, __global const kernel_rule_t *rule } } -__kernel void m02410_m08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 m02410_m16 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 m02410_s04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 m02410_sxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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) { /** * modifier @@ -580,11 +572,3 @@ __kernel void m02410_s04 (__global pw_t *pws, __global const kernel_rule_t *rule COMPARE_S_SIMD (a, d, c, b); } } - -__kernel void m02410_s08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 m02410_s16 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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) -{ -} diff --git a/OpenCL/m02410_a3.cl b/OpenCL/m02410_a3.cl index 2cc702a46..51c4f1e22 100644 --- a/OpenCL/m02410_a3.cl +++ b/OpenCL/m02410_a3.cl @@ -501,7 +501,7 @@ void m02410s (u32 w[16], const u32 pw_len, __global pw_t *pws, __global const ke } } -__kernel void m02410_m04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 m02410_mxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 @@ -539,15 +539,7 @@ __kernel void m02410_m04 (__global pw_t *pws, __global const kernel_rule_t *rule m02410m (w, pw_len, 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 m02410_m08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 m02410_m16 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 m02410_s04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 m02410_sxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 @@ -584,11 +576,3 @@ __kernel void m02410_s04 (__global pw_t *pws, __global const kernel_rule_t *rule m02410s (w, pw_len, 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 m02410_s08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 m02410_s16 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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) -{ -} diff --git a/OpenCL/m02710_a0.cl b/OpenCL/m02710_a0-optimized.cl similarity index 100% rename from OpenCL/m02710_a0.cl rename to OpenCL/m02710_a0-optimized.cl diff --git a/OpenCL/m02710_a1.cl b/OpenCL/m02710_a1-optimized.cl similarity index 100% rename from OpenCL/m02710_a1.cl rename to OpenCL/m02710_a1-optimized.cl diff --git a/OpenCL/m02710_a3.cl b/OpenCL/m02710_a3-optimized.cl similarity index 100% rename from OpenCL/m02710_a3.cl rename to OpenCL/m02710_a3-optimized.cl diff --git a/OpenCL/m03000_a0.cl b/OpenCL/m03000_a0.cl index 2ef586f3b..f54b03b05 100644 --- a/OpenCL/m03000_a0.cl +++ b/OpenCL/m03000_a0.cl @@ -3,7 +3,7 @@ * License.....: MIT */ -#define NEW_SIMD_CODE +//#define NEW_SIMD_CODE #include "inc_vendor.cl" #include "inc_hash_constants.h" @@ -491,7 +491,7 @@ void transform_netntlmv1_key (const u32x w0, const u32x w1, u32x out[2]) | ((k[7] & 0xff) << 24); } -__kernel void m03000_m04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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_mxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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) { /** * modifier @@ -591,15 +591,7 @@ __kernel void m03000_m04 (__global pw_t *pws, __global const kernel_rule_t *rule } } -__kernel void m03000_m08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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_sxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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) { /** * modifier @@ -710,11 +702,3 @@ __kernel void m03000_s04 (__global pw_t *pws, __global const kernel_rule_t *rule COMPARE_S_SIMD (iv[0], iv[1], z, z); } } - -__kernel void m03000_s08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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) -{ -} diff --git a/OpenCL/m03000_a1.cl b/OpenCL/m03000_a1.cl index 45ed69ffa..4a0fe3365 100644 --- a/OpenCL/m03000_a1.cl +++ b/OpenCL/m03000_a1.cl @@ -3,7 +3,7 @@ * License.....: MIT */ -#define NEW_SIMD_CODE +//#define NEW_SIMD_CODE #include "inc_vendor.cl" #include "inc_hash_constants.h" @@ -489,7 +489,7 @@ void transform_netntlmv1_key (const u32x w0, const u32x w1, u32x out[2]) | ((k[7] & 0xff) << 24); } -__kernel void m03000_m04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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_mxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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) { /** * modifier @@ -653,15 +653,7 @@ __kernel void m03000_m04 (__global pw_t *pws, __global const kernel_rule_t *rule } } -__kernel void m03000_m08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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_sxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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) { /** * modifier @@ -836,11 +828,3 @@ __kernel void m03000_s04 (__global pw_t *pws, __global const kernel_rule_t *rule COMPARE_S_SIMD (iv[0], iv[1], z, z); } } - -__kernel void m03000_s08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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) -{ -} diff --git a/OpenCL/m03000_a3.cl b/OpenCL/m03000_a3.cl index 95f563771..9817c3828 100644 --- a/OpenCL/m03000_a3.cl +++ b/OpenCL/m03000_a3.cl @@ -2481,7 +2481,7 @@ __kernel void m03000_tm (__global u32 *mod, __global bs_word_t *words_buf_r) } } -__kernel void m03000_m04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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_mxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 @@ -2499,15 +2499,7 @@ __kernel void m03000_m04 (__global pw_t *pws, __global const kernel_rule_t *rule 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 const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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_sxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 @@ -2524,11 +2516,3 @@ __kernel void m03000_s04 (__global pw_t *pws, __global const kernel_rule_t *rule 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 const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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 const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global bs_word_t * words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const 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) -{ -} diff --git a/OpenCL/m03100_a0.cl b/OpenCL/m03100_a0-optimized.cl similarity index 100% rename from OpenCL/m03100_a0.cl rename to OpenCL/m03100_a0-optimized.cl diff --git a/OpenCL/m03100_a1.cl b/OpenCL/m03100_a1-optimized.cl similarity index 100% rename from OpenCL/m03100_a1.cl rename to OpenCL/m03100_a1-optimized.cl diff --git a/OpenCL/m03100_a3.cl b/OpenCL/m03100_a3-optimized.cl similarity index 100% rename from OpenCL/m03100_a3.cl rename to OpenCL/m03100_a3-optimized.cl diff --git a/OpenCL/m04400_a0.cl b/OpenCL/m04400_a0-optimized.cl similarity index 100% rename from OpenCL/m04400_a0.cl rename to OpenCL/m04400_a0-optimized.cl diff --git a/OpenCL/m04400_a1.cl b/OpenCL/m04400_a1-optimized.cl similarity index 100% rename from OpenCL/m04400_a1.cl rename to OpenCL/m04400_a1-optimized.cl diff --git a/OpenCL/m04400_a3.cl b/OpenCL/m04400_a3-optimized.cl similarity index 100% rename from OpenCL/m04400_a3.cl rename to OpenCL/m04400_a3-optimized.cl diff --git a/OpenCL/m04500_a0.cl b/OpenCL/m04500_a0-optimized.cl similarity index 100% rename from OpenCL/m04500_a0.cl rename to OpenCL/m04500_a0-optimized.cl diff --git a/OpenCL/m04500_a1.cl b/OpenCL/m04500_a1-optimized.cl similarity index 100% rename from OpenCL/m04500_a1.cl rename to OpenCL/m04500_a1-optimized.cl diff --git a/OpenCL/m04500_a3.cl b/OpenCL/m04500_a3-optimized.cl similarity index 100% rename from OpenCL/m04500_a3.cl rename to OpenCL/m04500_a3-optimized.cl diff --git a/OpenCL/m04520_a0.cl b/OpenCL/m04520_a0-optimized.cl similarity index 100% rename from OpenCL/m04520_a0.cl rename to OpenCL/m04520_a0-optimized.cl diff --git a/OpenCL/m04520_a1.cl b/OpenCL/m04520_a1-optimized.cl similarity index 100% rename from OpenCL/m04520_a1.cl rename to OpenCL/m04520_a1-optimized.cl diff --git a/OpenCL/m04520_a3.cl b/OpenCL/m04520_a3-optimized.cl similarity index 100% rename from OpenCL/m04520_a3.cl rename to OpenCL/m04520_a3-optimized.cl diff --git a/OpenCL/m04700_a0.cl b/OpenCL/m04700_a0-optimized.cl similarity index 100% rename from OpenCL/m04700_a0.cl rename to OpenCL/m04700_a0-optimized.cl diff --git a/OpenCL/m04700_a1.cl b/OpenCL/m04700_a1-optimized.cl similarity index 100% rename from OpenCL/m04700_a1.cl rename to OpenCL/m04700_a1-optimized.cl diff --git a/OpenCL/m04700_a3.cl b/OpenCL/m04700_a3-optimized.cl similarity index 100% rename from OpenCL/m04700_a3.cl rename to OpenCL/m04700_a3-optimized.cl diff --git a/OpenCL/m04800_a0.cl b/OpenCL/m04800_a0-optimized.cl similarity index 100% rename from OpenCL/m04800_a0.cl rename to OpenCL/m04800_a0-optimized.cl diff --git a/OpenCL/m04800_a1.cl b/OpenCL/m04800_a1-optimized.cl similarity index 100% rename from OpenCL/m04800_a1.cl rename to OpenCL/m04800_a1-optimized.cl diff --git a/OpenCL/m04800_a3.cl b/OpenCL/m04800_a3-optimized.cl similarity index 100% rename from OpenCL/m04800_a3.cl rename to OpenCL/m04800_a3-optimized.cl diff --git a/OpenCL/m04900_a0.cl b/OpenCL/m04900_a0-optimized.cl similarity index 100% rename from OpenCL/m04900_a0.cl rename to OpenCL/m04900_a0-optimized.cl diff --git a/OpenCL/m04900_a1.cl b/OpenCL/m04900_a1-optimized.cl similarity index 100% rename from OpenCL/m04900_a1.cl rename to OpenCL/m04900_a1-optimized.cl diff --git a/OpenCL/m04900_a3.cl b/OpenCL/m04900_a3-optimized.cl similarity index 100% rename from OpenCL/m04900_a3.cl rename to OpenCL/m04900_a3-optimized.cl diff --git a/OpenCL/m05000_a0.cl b/OpenCL/m05000_a0-optimized.cl similarity index 100% rename from OpenCL/m05000_a0.cl rename to OpenCL/m05000_a0-optimized.cl diff --git a/OpenCL/m05000_a1.cl b/OpenCL/m05000_a1-optimized.cl similarity index 100% rename from OpenCL/m05000_a1.cl rename to OpenCL/m05000_a1-optimized.cl diff --git a/OpenCL/m05000_a3.cl b/OpenCL/m05000_a3-optimized.cl similarity index 100% rename from OpenCL/m05000_a3.cl rename to OpenCL/m05000_a3-optimized.cl diff --git a/OpenCL/m05100_a0.cl b/OpenCL/m05100_a0-optimized.cl similarity index 100% rename from OpenCL/m05100_a0.cl rename to OpenCL/m05100_a0-optimized.cl diff --git a/OpenCL/m05100_a1.cl b/OpenCL/m05100_a1-optimized.cl similarity index 100% rename from OpenCL/m05100_a1.cl rename to OpenCL/m05100_a1-optimized.cl diff --git a/OpenCL/m05100_a3.cl b/OpenCL/m05100_a3-optimized.cl similarity index 100% rename from OpenCL/m05100_a3.cl rename to OpenCL/m05100_a3-optimized.cl diff --git a/OpenCL/m05300_a0.cl b/OpenCL/m05300_a0-optimized.cl similarity index 100% rename from OpenCL/m05300_a0.cl rename to OpenCL/m05300_a0-optimized.cl diff --git a/OpenCL/m05300_a1.cl b/OpenCL/m05300_a1-optimized.cl similarity index 100% rename from OpenCL/m05300_a1.cl rename to OpenCL/m05300_a1-optimized.cl diff --git a/OpenCL/m05300_a3.cl b/OpenCL/m05300_a3-optimized.cl similarity index 100% rename from OpenCL/m05300_a3.cl rename to OpenCL/m05300_a3-optimized.cl diff --git a/OpenCL/m05400_a0.cl b/OpenCL/m05400_a0-optimized.cl similarity index 100% rename from OpenCL/m05400_a0.cl rename to OpenCL/m05400_a0-optimized.cl diff --git a/OpenCL/m05400_a1.cl b/OpenCL/m05400_a1-optimized.cl similarity index 100% rename from OpenCL/m05400_a1.cl rename to OpenCL/m05400_a1-optimized.cl diff --git a/OpenCL/m05400_a3.cl b/OpenCL/m05400_a3-optimized.cl similarity index 100% rename from OpenCL/m05400_a3.cl rename to OpenCL/m05400_a3-optimized.cl diff --git a/OpenCL/m05500_a0.cl b/OpenCL/m05500_a0-optimized.cl similarity index 100% rename from OpenCL/m05500_a0.cl rename to OpenCL/m05500_a0-optimized.cl diff --git a/OpenCL/m05500_a1.cl b/OpenCL/m05500_a1-optimized.cl similarity index 100% rename from OpenCL/m05500_a1.cl rename to OpenCL/m05500_a1-optimized.cl diff --git a/OpenCL/m05500_a3.cl b/OpenCL/m05500_a3-optimized.cl similarity index 100% rename from OpenCL/m05500_a3.cl rename to OpenCL/m05500_a3-optimized.cl diff --git a/OpenCL/m05600_a0.cl b/OpenCL/m05600_a0-optimized.cl similarity index 100% rename from OpenCL/m05600_a0.cl rename to OpenCL/m05600_a0-optimized.cl diff --git a/OpenCL/m05600_a1.cl b/OpenCL/m05600_a1-optimized.cl similarity index 100% rename from OpenCL/m05600_a1.cl rename to OpenCL/m05600_a1-optimized.cl diff --git a/OpenCL/m05600_a3.cl b/OpenCL/m05600_a3-optimized.cl similarity index 100% rename from OpenCL/m05600_a3.cl rename to OpenCL/m05600_a3-optimized.cl diff --git a/OpenCL/m06000_a0.cl b/OpenCL/m06000_a0-optimized.cl similarity index 100% rename from OpenCL/m06000_a0.cl rename to OpenCL/m06000_a0-optimized.cl diff --git a/OpenCL/m06000_a1.cl b/OpenCL/m06000_a1-optimized.cl similarity index 100% rename from OpenCL/m06000_a1.cl rename to OpenCL/m06000_a1-optimized.cl diff --git a/OpenCL/m06000_a3.cl b/OpenCL/m06000_a3-optimized.cl similarity index 100% rename from OpenCL/m06000_a3.cl rename to OpenCL/m06000_a3-optimized.cl diff --git a/OpenCL/m06900_a0.cl b/OpenCL/m06900_a0-optimized.cl similarity index 100% rename from OpenCL/m06900_a0.cl rename to OpenCL/m06900_a0-optimized.cl diff --git a/OpenCL/m06900_a1.cl b/OpenCL/m06900_a1-optimized.cl similarity index 100% rename from OpenCL/m06900_a1.cl rename to OpenCL/m06900_a1-optimized.cl diff --git a/OpenCL/m06900_a3.cl b/OpenCL/m06900_a3-optimized.cl similarity index 100% rename from OpenCL/m06900_a3.cl rename to OpenCL/m06900_a3-optimized.cl diff --git a/OpenCL/m07000_a0.cl b/OpenCL/m07000_a0-optimized.cl similarity index 100% rename from OpenCL/m07000_a0.cl rename to OpenCL/m07000_a0-optimized.cl diff --git a/OpenCL/m07000_a1.cl b/OpenCL/m07000_a1-optimized.cl similarity index 100% rename from OpenCL/m07000_a1.cl rename to OpenCL/m07000_a1-optimized.cl diff --git a/OpenCL/m07000_a3.cl b/OpenCL/m07000_a3-optimized.cl similarity index 100% rename from OpenCL/m07000_a3.cl rename to OpenCL/m07000_a3-optimized.cl diff --git a/OpenCL/m07300_a0.cl b/OpenCL/m07300_a0-optimized.cl similarity index 100% rename from OpenCL/m07300_a0.cl rename to OpenCL/m07300_a0-optimized.cl diff --git a/OpenCL/m07300_a1.cl b/OpenCL/m07300_a1-optimized.cl similarity index 100% rename from OpenCL/m07300_a1.cl rename to OpenCL/m07300_a1-optimized.cl diff --git a/OpenCL/m07300_a3.cl b/OpenCL/m07300_a3-optimized.cl similarity index 100% rename from OpenCL/m07300_a3.cl rename to OpenCL/m07300_a3-optimized.cl diff --git a/OpenCL/m07500_a0.cl b/OpenCL/m07500_a0-optimized.cl similarity index 100% rename from OpenCL/m07500_a0.cl rename to OpenCL/m07500_a0-optimized.cl diff --git a/OpenCL/m07500_a1.cl b/OpenCL/m07500_a1-optimized.cl similarity index 100% rename from OpenCL/m07500_a1.cl rename to OpenCL/m07500_a1-optimized.cl diff --git a/OpenCL/m07500_a3.cl b/OpenCL/m07500_a3-optimized.cl similarity index 100% rename from OpenCL/m07500_a3.cl rename to OpenCL/m07500_a3-optimized.cl diff --git a/OpenCL/m07700_a0.cl b/OpenCL/m07700_a0-optimized.cl similarity index 100% rename from OpenCL/m07700_a0.cl rename to OpenCL/m07700_a0-optimized.cl diff --git a/OpenCL/m07700_a1.cl b/OpenCL/m07700_a1-optimized.cl similarity index 100% rename from OpenCL/m07700_a1.cl rename to OpenCL/m07700_a1-optimized.cl diff --git a/OpenCL/m07700_a3.cl b/OpenCL/m07700_a3-optimized.cl similarity index 100% rename from OpenCL/m07700_a3.cl rename to OpenCL/m07700_a3-optimized.cl diff --git a/OpenCL/m07800_a0.cl b/OpenCL/m07800_a0-optimized.cl similarity index 100% rename from OpenCL/m07800_a0.cl rename to OpenCL/m07800_a0-optimized.cl diff --git a/OpenCL/m07800_a1.cl b/OpenCL/m07800_a1-optimized.cl similarity index 100% rename from OpenCL/m07800_a1.cl rename to OpenCL/m07800_a1-optimized.cl diff --git a/OpenCL/m07800_a3.cl b/OpenCL/m07800_a3-optimized.cl similarity index 100% rename from OpenCL/m07800_a3.cl rename to OpenCL/m07800_a3-optimized.cl diff --git a/OpenCL/m08000_a0.cl b/OpenCL/m08000_a0-optimized.cl similarity index 100% rename from OpenCL/m08000_a0.cl rename to OpenCL/m08000_a0-optimized.cl diff --git a/OpenCL/m08000_a1.cl b/OpenCL/m08000_a1-optimized.cl similarity index 100% rename from OpenCL/m08000_a1.cl rename to OpenCL/m08000_a1-optimized.cl diff --git a/OpenCL/m08000_a3.cl b/OpenCL/m08000_a3-optimized.cl similarity index 100% rename from OpenCL/m08000_a3.cl rename to OpenCL/m08000_a3-optimized.cl diff --git a/OpenCL/m08100_a0.cl b/OpenCL/m08100_a0-optimized.cl similarity index 100% rename from OpenCL/m08100_a0.cl rename to OpenCL/m08100_a0-optimized.cl diff --git a/OpenCL/m08100_a1.cl b/OpenCL/m08100_a1-optimized.cl similarity index 100% rename from OpenCL/m08100_a1.cl rename to OpenCL/m08100_a1-optimized.cl diff --git a/OpenCL/m08100_a3.cl b/OpenCL/m08100_a3-optimized.cl similarity index 100% rename from OpenCL/m08100_a3.cl rename to OpenCL/m08100_a3-optimized.cl diff --git a/OpenCL/m08300_a0.cl b/OpenCL/m08300_a0-optimized.cl similarity index 100% rename from OpenCL/m08300_a0.cl rename to OpenCL/m08300_a0-optimized.cl diff --git a/OpenCL/m08300_a1.cl b/OpenCL/m08300_a1-optimized.cl similarity index 100% rename from OpenCL/m08300_a1.cl rename to OpenCL/m08300_a1-optimized.cl diff --git a/OpenCL/m08300_a3.cl b/OpenCL/m08300_a3-optimized.cl similarity index 100% rename from OpenCL/m08300_a3.cl rename to OpenCL/m08300_a3-optimized.cl diff --git a/OpenCL/m08400_a0.cl b/OpenCL/m08400_a0-optimized.cl similarity index 100% rename from OpenCL/m08400_a0.cl rename to OpenCL/m08400_a0-optimized.cl diff --git a/OpenCL/m08400_a1.cl b/OpenCL/m08400_a1-optimized.cl similarity index 100% rename from OpenCL/m08400_a1.cl rename to OpenCL/m08400_a1-optimized.cl diff --git a/OpenCL/m08400_a3.cl b/OpenCL/m08400_a3-optimized.cl similarity index 100% rename from OpenCL/m08400_a3.cl rename to OpenCL/m08400_a3-optimized.cl diff --git a/OpenCL/m08500_a0.cl b/OpenCL/m08500_a0-optimized.cl similarity index 100% rename from OpenCL/m08500_a0.cl rename to OpenCL/m08500_a0-optimized.cl diff --git a/OpenCL/m08500_a1.cl b/OpenCL/m08500_a1-optimized.cl similarity index 100% rename from OpenCL/m08500_a1.cl rename to OpenCL/m08500_a1-optimized.cl diff --git a/OpenCL/m08500_a3.cl b/OpenCL/m08500_a3-optimized.cl similarity index 100% rename from OpenCL/m08500_a3.cl rename to OpenCL/m08500_a3-optimized.cl diff --git a/OpenCL/m08600_a0.cl b/OpenCL/m08600_a0-optimized.cl similarity index 100% rename from OpenCL/m08600_a0.cl rename to OpenCL/m08600_a0-optimized.cl diff --git a/OpenCL/m08600_a1.cl b/OpenCL/m08600_a1-optimized.cl similarity index 100% rename from OpenCL/m08600_a1.cl rename to OpenCL/m08600_a1-optimized.cl diff --git a/OpenCL/m08600_a3.cl b/OpenCL/m08600_a3-optimized.cl similarity index 100% rename from OpenCL/m08600_a3.cl rename to OpenCL/m08600_a3-optimized.cl diff --git a/OpenCL/m08700_a0.cl b/OpenCL/m08700_a0-optimized.cl similarity index 100% rename from OpenCL/m08700_a0.cl rename to OpenCL/m08700_a0-optimized.cl diff --git a/OpenCL/m08700_a1.cl b/OpenCL/m08700_a1-optimized.cl similarity index 100% rename from OpenCL/m08700_a1.cl rename to OpenCL/m08700_a1-optimized.cl diff --git a/OpenCL/m08700_a3.cl b/OpenCL/m08700_a3-optimized.cl similarity index 100% rename from OpenCL/m08700_a3.cl rename to OpenCL/m08700_a3-optimized.cl diff --git a/OpenCL/m09700_a0.cl b/OpenCL/m09700_a0-optimized.cl similarity index 100% rename from OpenCL/m09700_a0.cl rename to OpenCL/m09700_a0-optimized.cl diff --git a/OpenCL/m09700_a1.cl b/OpenCL/m09700_a1-optimized.cl similarity index 100% rename from OpenCL/m09700_a1.cl rename to OpenCL/m09700_a1-optimized.cl diff --git a/OpenCL/m09700_a3.cl b/OpenCL/m09700_a3-optimized.cl similarity index 100% rename from OpenCL/m09700_a3.cl rename to OpenCL/m09700_a3-optimized.cl diff --git a/OpenCL/m09710_a0.cl b/OpenCL/m09710_a0-optimized.cl similarity index 100% rename from OpenCL/m09710_a0.cl rename to OpenCL/m09710_a0-optimized.cl diff --git a/OpenCL/m09710_a1.cl b/OpenCL/m09710_a1-optimized.cl similarity index 100% rename from OpenCL/m09710_a1.cl rename to OpenCL/m09710_a1-optimized.cl diff --git a/OpenCL/m09710_a3.cl b/OpenCL/m09710_a3-optimized.cl similarity index 100% rename from OpenCL/m09710_a3.cl rename to OpenCL/m09710_a3-optimized.cl diff --git a/OpenCL/m09720_a0.cl b/OpenCL/m09720_a0-optimized.cl similarity index 100% rename from OpenCL/m09720_a0.cl rename to OpenCL/m09720_a0-optimized.cl diff --git a/OpenCL/m09720_a1.cl b/OpenCL/m09720_a1-optimized.cl similarity index 100% rename from OpenCL/m09720_a1.cl rename to OpenCL/m09720_a1-optimized.cl diff --git a/OpenCL/m09720_a3.cl b/OpenCL/m09720_a3-optimized.cl similarity index 100% rename from OpenCL/m09720_a3.cl rename to OpenCL/m09720_a3-optimized.cl diff --git a/OpenCL/m09800_a0.cl b/OpenCL/m09800_a0-optimized.cl similarity index 100% rename from OpenCL/m09800_a0.cl rename to OpenCL/m09800_a0-optimized.cl diff --git a/OpenCL/m09800_a1.cl b/OpenCL/m09800_a1-optimized.cl similarity index 100% rename from OpenCL/m09800_a1.cl rename to OpenCL/m09800_a1-optimized.cl diff --git a/OpenCL/m09800_a3.cl b/OpenCL/m09800_a3-optimized.cl similarity index 100% rename from OpenCL/m09800_a3.cl rename to OpenCL/m09800_a3-optimized.cl diff --git a/OpenCL/m09810_a0.cl b/OpenCL/m09810_a0-optimized.cl similarity index 100% rename from OpenCL/m09810_a0.cl rename to OpenCL/m09810_a0-optimized.cl diff --git a/OpenCL/m09810_a1.cl b/OpenCL/m09810_a1-optimized.cl similarity index 100% rename from OpenCL/m09810_a1.cl rename to OpenCL/m09810_a1-optimized.cl diff --git a/OpenCL/m09810_a3.cl b/OpenCL/m09810_a3-optimized.cl similarity index 100% rename from OpenCL/m09810_a3.cl rename to OpenCL/m09810_a3-optimized.cl diff --git a/OpenCL/m09820_a0.cl b/OpenCL/m09820_a0-optimized.cl similarity index 100% rename from OpenCL/m09820_a0.cl rename to OpenCL/m09820_a0-optimized.cl diff --git a/OpenCL/m09820_a1.cl b/OpenCL/m09820_a1-optimized.cl similarity index 100% rename from OpenCL/m09820_a1.cl rename to OpenCL/m09820_a1-optimized.cl diff --git a/OpenCL/m09820_a3.cl b/OpenCL/m09820_a3-optimized.cl similarity index 100% rename from OpenCL/m09820_a3.cl rename to OpenCL/m09820_a3-optimized.cl diff --git a/OpenCL/m09900_a0.cl b/OpenCL/m09900_a0-optimized.cl similarity index 100% rename from OpenCL/m09900_a0.cl rename to OpenCL/m09900_a0-optimized.cl diff --git a/OpenCL/m09900_a1.cl b/OpenCL/m09900_a1-optimized.cl similarity index 100% rename from OpenCL/m09900_a1.cl rename to OpenCL/m09900_a1-optimized.cl diff --git a/OpenCL/m09900_a3.cl b/OpenCL/m09900_a3-optimized.cl similarity index 100% rename from OpenCL/m09900_a3.cl rename to OpenCL/m09900_a3-optimized.cl diff --git a/OpenCL/m10100_a0.cl b/OpenCL/m10100_a0-optimized.cl similarity index 100% rename from OpenCL/m10100_a0.cl rename to OpenCL/m10100_a0-optimized.cl diff --git a/OpenCL/m10100_a1.cl b/OpenCL/m10100_a1-optimized.cl similarity index 100% rename from OpenCL/m10100_a1.cl rename to OpenCL/m10100_a1-optimized.cl diff --git a/OpenCL/m10100_a3.cl b/OpenCL/m10100_a3-optimized.cl similarity index 100% rename from OpenCL/m10100_a3.cl rename to OpenCL/m10100_a3-optimized.cl diff --git a/OpenCL/m10400_a0.cl b/OpenCL/m10400_a0-optimized.cl similarity index 100% rename from OpenCL/m10400_a0.cl rename to OpenCL/m10400_a0-optimized.cl diff --git a/OpenCL/m10400_a1.cl b/OpenCL/m10400_a1-optimized.cl similarity index 100% rename from OpenCL/m10400_a1.cl rename to OpenCL/m10400_a1-optimized.cl diff --git a/OpenCL/m10400_a3.cl b/OpenCL/m10400_a3-optimized.cl similarity index 100% rename from OpenCL/m10400_a3.cl rename to OpenCL/m10400_a3-optimized.cl diff --git a/OpenCL/m10410_a0.cl b/OpenCL/m10410_a0-optimized.cl similarity index 100% rename from OpenCL/m10410_a0.cl rename to OpenCL/m10410_a0-optimized.cl diff --git a/OpenCL/m10410_a1.cl b/OpenCL/m10410_a1-optimized.cl similarity index 100% rename from OpenCL/m10410_a1.cl rename to OpenCL/m10410_a1-optimized.cl diff --git a/OpenCL/m10410_a3.cl b/OpenCL/m10410_a3-optimized.cl similarity index 100% rename from OpenCL/m10410_a3.cl rename to OpenCL/m10410_a3-optimized.cl diff --git a/OpenCL/m10420_a0.cl b/OpenCL/m10420_a0-optimized.cl similarity index 100% rename from OpenCL/m10420_a0.cl rename to OpenCL/m10420_a0-optimized.cl diff --git a/OpenCL/m10420_a1.cl b/OpenCL/m10420_a1-optimized.cl similarity index 100% rename from OpenCL/m10420_a1.cl rename to OpenCL/m10420_a1-optimized.cl diff --git a/OpenCL/m10420_a3.cl b/OpenCL/m10420_a3-optimized.cl similarity index 100% rename from OpenCL/m10420_a3.cl rename to OpenCL/m10420_a3-optimized.cl diff --git a/OpenCL/m11000_a0.cl b/OpenCL/m11000_a0-optimized.cl similarity index 100% rename from OpenCL/m11000_a0.cl rename to OpenCL/m11000_a0-optimized.cl diff --git a/OpenCL/m11000_a1.cl b/OpenCL/m11000_a1-optimized.cl similarity index 100% rename from OpenCL/m11000_a1.cl rename to OpenCL/m11000_a1-optimized.cl diff --git a/OpenCL/m11000_a3.cl b/OpenCL/m11000_a3-optimized.cl similarity index 100% rename from OpenCL/m11000_a3.cl rename to OpenCL/m11000_a3-optimized.cl diff --git a/OpenCL/m11100_a0.cl b/OpenCL/m11100_a0-optimized.cl similarity index 100% rename from OpenCL/m11100_a0.cl rename to OpenCL/m11100_a0-optimized.cl diff --git a/OpenCL/m11100_a1.cl b/OpenCL/m11100_a1-optimized.cl similarity index 100% rename from OpenCL/m11100_a1.cl rename to OpenCL/m11100_a1-optimized.cl diff --git a/OpenCL/m11100_a3.cl b/OpenCL/m11100_a3-optimized.cl similarity index 100% rename from OpenCL/m11100_a3.cl rename to OpenCL/m11100_a3-optimized.cl diff --git a/OpenCL/m11200_a0.cl b/OpenCL/m11200_a0-optimized.cl similarity index 100% rename from OpenCL/m11200_a0.cl rename to OpenCL/m11200_a0-optimized.cl diff --git a/OpenCL/m11200_a1.cl b/OpenCL/m11200_a1-optimized.cl similarity index 100% rename from OpenCL/m11200_a1.cl rename to OpenCL/m11200_a1-optimized.cl diff --git a/OpenCL/m11200_a3.cl b/OpenCL/m11200_a3-optimized.cl similarity index 100% rename from OpenCL/m11200_a3.cl rename to OpenCL/m11200_a3-optimized.cl diff --git a/OpenCL/m11400_a0.cl b/OpenCL/m11400_a0-optimized.cl similarity index 100% rename from OpenCL/m11400_a0.cl rename to OpenCL/m11400_a0-optimized.cl diff --git a/OpenCL/m11400_a1.cl b/OpenCL/m11400_a1-optimized.cl similarity index 100% rename from OpenCL/m11400_a1.cl rename to OpenCL/m11400_a1-optimized.cl diff --git a/OpenCL/m11400_a3.cl b/OpenCL/m11400_a3-optimized.cl similarity index 100% rename from OpenCL/m11400_a3.cl rename to OpenCL/m11400_a3-optimized.cl diff --git a/OpenCL/m11500_a0.cl b/OpenCL/m11500_a0-optimized.cl similarity index 100% rename from OpenCL/m11500_a0.cl rename to OpenCL/m11500_a0-optimized.cl diff --git a/OpenCL/m11500_a1.cl b/OpenCL/m11500_a1-optimized.cl similarity index 100% rename from OpenCL/m11500_a1.cl rename to OpenCL/m11500_a1-optimized.cl diff --git a/OpenCL/m11500_a3.cl b/OpenCL/m11500_a3-optimized.cl similarity index 100% rename from OpenCL/m11500_a3.cl rename to OpenCL/m11500_a3-optimized.cl diff --git a/OpenCL/m11700_a0.cl b/OpenCL/m11700_a0-optimized.cl similarity index 100% rename from OpenCL/m11700_a0.cl rename to OpenCL/m11700_a0-optimized.cl diff --git a/OpenCL/m11700_a1.cl b/OpenCL/m11700_a1-optimized.cl similarity index 100% rename from OpenCL/m11700_a1.cl rename to OpenCL/m11700_a1-optimized.cl diff --git a/OpenCL/m11700_a3.cl b/OpenCL/m11700_a3-optimized.cl similarity index 100% rename from OpenCL/m11700_a3.cl rename to OpenCL/m11700_a3-optimized.cl diff --git a/OpenCL/m11800_a0.cl b/OpenCL/m11800_a0-optimized.cl similarity index 100% rename from OpenCL/m11800_a0.cl rename to OpenCL/m11800_a0-optimized.cl diff --git a/OpenCL/m11800_a1.cl b/OpenCL/m11800_a1-optimized.cl similarity index 100% rename from OpenCL/m11800_a1.cl rename to OpenCL/m11800_a1-optimized.cl diff --git a/OpenCL/m11800_a3.cl b/OpenCL/m11800_a3-optimized.cl similarity index 100% rename from OpenCL/m11800_a3.cl rename to OpenCL/m11800_a3-optimized.cl diff --git a/OpenCL/m12600_a0.cl b/OpenCL/m12600_a0-optimized.cl similarity index 100% rename from OpenCL/m12600_a0.cl rename to OpenCL/m12600_a0-optimized.cl diff --git a/OpenCL/m12600_a1.cl b/OpenCL/m12600_a1-optimized.cl similarity index 100% rename from OpenCL/m12600_a1.cl rename to OpenCL/m12600_a1-optimized.cl diff --git a/OpenCL/m12600_a3.cl b/OpenCL/m12600_a3-optimized.cl similarity index 100% rename from OpenCL/m12600_a3.cl rename to OpenCL/m12600_a3-optimized.cl diff --git a/OpenCL/m13100_a0.cl b/OpenCL/m13100_a0-optimized.cl similarity index 100% rename from OpenCL/m13100_a0.cl rename to OpenCL/m13100_a0-optimized.cl diff --git a/OpenCL/m13100_a1.cl b/OpenCL/m13100_a1-optimized.cl similarity index 100% rename from OpenCL/m13100_a1.cl rename to OpenCL/m13100_a1-optimized.cl diff --git a/OpenCL/m13100_a3.cl b/OpenCL/m13100_a3-optimized.cl similarity index 100% rename from OpenCL/m13100_a3.cl rename to OpenCL/m13100_a3-optimized.cl diff --git a/OpenCL/m13300_a0.cl b/OpenCL/m13300_a0-optimized.cl similarity index 100% rename from OpenCL/m13300_a0.cl rename to OpenCL/m13300_a0-optimized.cl diff --git a/OpenCL/m13300_a1.cl b/OpenCL/m13300_a1-optimized.cl similarity index 100% rename from OpenCL/m13300_a1.cl rename to OpenCL/m13300_a1-optimized.cl diff --git a/OpenCL/m13300_a3.cl b/OpenCL/m13300_a3-optimized.cl similarity index 100% rename from OpenCL/m13300_a3.cl rename to OpenCL/m13300_a3-optimized.cl diff --git a/OpenCL/m13500_a0.cl b/OpenCL/m13500_a0-optimized.cl similarity index 100% rename from OpenCL/m13500_a0.cl rename to OpenCL/m13500_a0-optimized.cl diff --git a/OpenCL/m13500_a1.cl b/OpenCL/m13500_a1-optimized.cl similarity index 100% rename from OpenCL/m13500_a1.cl rename to OpenCL/m13500_a1-optimized.cl diff --git a/OpenCL/m13500_a3.cl b/OpenCL/m13500_a3-optimized.cl similarity index 100% rename from OpenCL/m13500_a3.cl rename to OpenCL/m13500_a3-optimized.cl diff --git a/OpenCL/m13800_a0.cl b/OpenCL/m13800_a0-optimized.cl similarity index 100% rename from OpenCL/m13800_a0.cl rename to OpenCL/m13800_a0-optimized.cl diff --git a/OpenCL/m13800_a1.cl b/OpenCL/m13800_a1-optimized.cl similarity index 100% rename from OpenCL/m13800_a1.cl rename to OpenCL/m13800_a1-optimized.cl diff --git a/OpenCL/m13800_a3.cl b/OpenCL/m13800_a3-optimized.cl similarity index 100% rename from OpenCL/m13800_a3.cl rename to OpenCL/m13800_a3-optimized.cl diff --git a/OpenCL/m13900_a0.cl b/OpenCL/m13900_a0-optimized.cl similarity index 100% rename from OpenCL/m13900_a0.cl rename to OpenCL/m13900_a0-optimized.cl diff --git a/OpenCL/m13900_a1.cl b/OpenCL/m13900_a1-optimized.cl similarity index 100% rename from OpenCL/m13900_a1.cl rename to OpenCL/m13900_a1-optimized.cl diff --git a/OpenCL/m13900_a3.cl b/OpenCL/m13900_a3-optimized.cl similarity index 100% rename from OpenCL/m13900_a3.cl rename to OpenCL/m13900_a3-optimized.cl diff --git a/OpenCL/m14000_a0.cl b/OpenCL/m14000_a0-optimized.cl similarity index 100% rename from OpenCL/m14000_a0.cl rename to OpenCL/m14000_a0-optimized.cl diff --git a/OpenCL/m14000_a1.cl b/OpenCL/m14000_a1-optimized.cl similarity index 100% rename from OpenCL/m14000_a1.cl rename to OpenCL/m14000_a1-optimized.cl diff --git a/OpenCL/m14000_a3.cl b/OpenCL/m14000_a3-optimized.cl similarity index 100% rename from OpenCL/m14000_a3.cl rename to OpenCL/m14000_a3-optimized.cl diff --git a/OpenCL/m14100_a0.cl b/OpenCL/m14100_a0-optimized.cl similarity index 100% rename from OpenCL/m14100_a0.cl rename to OpenCL/m14100_a0-optimized.cl diff --git a/OpenCL/m14100_a1.cl b/OpenCL/m14100_a1-optimized.cl similarity index 100% rename from OpenCL/m14100_a1.cl rename to OpenCL/m14100_a1-optimized.cl diff --git a/OpenCL/m14100_a3.cl b/OpenCL/m14100_a3-optimized.cl similarity index 100% rename from OpenCL/m14100_a3.cl rename to OpenCL/m14100_a3-optimized.cl diff --git a/OpenCL/m14400_a0.cl b/OpenCL/m14400_a0-optimized.cl similarity index 100% rename from OpenCL/m14400_a0.cl rename to OpenCL/m14400_a0-optimized.cl diff --git a/OpenCL/m14400_a1.cl b/OpenCL/m14400_a1-optimized.cl similarity index 100% rename from OpenCL/m14400_a1.cl rename to OpenCL/m14400_a1-optimized.cl diff --git a/OpenCL/m14400_a3.cl b/OpenCL/m14400_a3-optimized.cl similarity index 100% rename from OpenCL/m14400_a3.cl rename to OpenCL/m14400_a3-optimized.cl diff --git a/OpenCL/m14900_a0.cl b/OpenCL/m14900_a0-optimized.cl similarity index 100% rename from OpenCL/m14900_a0.cl rename to OpenCL/m14900_a0-optimized.cl diff --git a/OpenCL/m14900_a1.cl b/OpenCL/m14900_a1-optimized.cl similarity index 100% rename from OpenCL/m14900_a1.cl rename to OpenCL/m14900_a1-optimized.cl diff --git a/OpenCL/m14900_a3.cl b/OpenCL/m14900_a3-optimized.cl similarity index 100% rename from OpenCL/m14900_a3.cl rename to OpenCL/m14900_a3-optimized.cl diff --git a/OpenCL/m15000_a0.cl b/OpenCL/m15000_a0-optimized.cl similarity index 100% rename from OpenCL/m15000_a0.cl rename to OpenCL/m15000_a0-optimized.cl diff --git a/OpenCL/m15000_a1.cl b/OpenCL/m15000_a1-optimized.cl similarity index 100% rename from OpenCL/m15000_a1.cl rename to OpenCL/m15000_a1-optimized.cl diff --git a/OpenCL/m15000_a3.cl b/OpenCL/m15000_a3-optimized.cl similarity index 100% rename from OpenCL/m15000_a3.cl rename to OpenCL/m15000_a3-optimized.cl diff --git a/OpenCL/m15400_a0.cl b/OpenCL/m15400_a0-optimized.cl similarity index 100% rename from OpenCL/m15400_a0.cl rename to OpenCL/m15400_a0-optimized.cl diff --git a/OpenCL/m15400_a1.cl b/OpenCL/m15400_a1-optimized.cl similarity index 100% rename from OpenCL/m15400_a1.cl rename to OpenCL/m15400_a1-optimized.cl diff --git a/OpenCL/m15400_a3.cl b/OpenCL/m15400_a3-optimized.cl similarity index 100% rename from OpenCL/m15400_a3.cl rename to OpenCL/m15400_a3-optimized.cl diff --git a/OpenCL/m15500_a0.cl b/OpenCL/m15500_a0-optimized.cl similarity index 100% rename from OpenCL/m15500_a0.cl rename to OpenCL/m15500_a0-optimized.cl diff --git a/OpenCL/m15500_a1.cl b/OpenCL/m15500_a1-optimized.cl similarity index 100% rename from OpenCL/m15500_a1.cl rename to OpenCL/m15500_a1-optimized.cl diff --git a/OpenCL/m15500_a3.cl b/OpenCL/m15500_a3-optimized.cl similarity index 100% rename from OpenCL/m15500_a3.cl rename to OpenCL/m15500_a3-optimized.cl diff --git a/include/opencl.h b/include/opencl.h index 17ec6f74b..a190197f3 100644 --- a/include/opencl.h +++ b/include/opencl.h @@ -66,6 +66,13 @@ int run_kernel_bzero (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_par int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const u32 pws_cnt); int run_cracker (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const u32 pws_cnt); +void generate_source_kernel_filename (const u32 attack_exec, const u32 attack_kern, const u32 kern_type, const u32 opti_type, char *shared_dir, char *source_file); +void generate_cached_kernel_filename (const u32 attack_exec, const u32 attack_kern, const u32 kern_type, const u32 opti_type, char *profile_dir, const char *device_name_chksum, char *cached_file); +void generate_source_kernel_mp_filename (const u32 opti_type, const u64 opts_type, char *shared_dir, char *source_file); +void generate_cached_kernel_mp_filename (const u32 opti_type, const u64 opts_type, char *profile_dir, const char *device_name_chksum, char *cached_file); +void generate_source_kernel_amp_filename (const u32 attack_kern, char *shared_dir, char *source_file); +void generate_cached_kernel_amp_filename (const u32 attack_kern, char *profile_dir, const char *device_name_chksum, char *cached_file); + int opencl_ctx_init (hashcat_ctx_t *hashcat_ctx); void opencl_ctx_destroy (hashcat_ctx_t *hashcat_ctx); diff --git a/include/types.h b/include/types.h index d1ee0df63..07d1095ee 100644 --- a/include/types.h +++ b/include/types.h @@ -326,6 +326,7 @@ typedef enum salt_type typedef enum opti_type { + OPTI_TYPE_OPTIMIZED_KERNEL = (1 << 0), OPTI_TYPE_ZERO_BYTE = (1 << 1), OPTI_TYPE_PRECOMPUTE_INIT = (1 << 2), OPTI_TYPE_PRECOMPUTE_MERKLE = (1 << 3), diff --git a/src/autotune.c b/src/autotune.c index d997e8b7c..652319f59 100644 --- a/src/autotune.c +++ b/src/autotune.c @@ -13,8 +13,7 @@ static double try_run (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const u32 kernel_accel, const u32 kernel_loops) { - hashconfig_t *hashconfig = hashcat_ctx->hashconfig; - user_options_t *user_options = hashcat_ctx->user_options; + hashconfig_t *hashconfig = hashcat_ctx->hashconfig; device_param->kernel_params_buf32[28] = 0; device_param->kernel_params_buf32[29] = kernel_loops; // not a bug, both need to be set @@ -22,7 +21,7 @@ static double try_run (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_par if (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL) { - if (user_options->optimized_kernel_enable == true) + if (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) { const u32 kernel_power_try = device_param->device_processors * device_param->kernel_threads_by_wgs_kernel1 * kernel_accel; diff --git a/src/combinator.c b/src/combinator.c index 5ca42fb87..316a0568a 100644 --- a/src/combinator.c +++ b/src/combinator.c @@ -14,6 +14,7 @@ int combinator_ctx_init (hashcat_ctx_t *hashcat_ctx) { combinator_ctx_t *combinator_ctx = hashcat_ctx->combinator_ctx; + hashconfig_t *hashconfig = hashcat_ctx->hashconfig; user_options_extra_t *user_options_extra = hashcat_ctx->user_options_extra; user_options_t *user_options = hashcat_ctx->user_options; @@ -33,7 +34,7 @@ int combinator_ctx_init (hashcat_ctx_t *hashcat_ctx) combinator_ctx->scratch_buf = (char *) hcmalloc (HCBUFSIZ_LARGE); - if (user_options->optimized_kernel_enable == true) + if (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) { if (user_options->attack_mode == ATTACK_MODE_COMBI) { diff --git a/src/dispatch.c b/src/dispatch.c index f2ff388b3..cd7bfd962 100644 --- a/src/dispatch.c +++ b/src/dispatch.c @@ -313,9 +313,9 @@ static int calc (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param) const u32 attack_mode = user_options->attack_mode; const u32 attack_kern = user_options_extra->attack_kern; - if ((attack_mode == ATTACK_MODE_BF) || ((user_options->optimized_kernel_enable == false) && (attack_mode == ATTACK_MODE_HYBRID2))) + if ((attack_mode == ATTACK_MODE_BF) || (((hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) == 0) && (attack_mode == ATTACK_MODE_HYBRID2))) { - if ((user_options->optimized_kernel_enable == false) && (attack_mode == ATTACK_MODE_HYBRID2)) + if (((hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) == 0) && (attack_mode == ATTACK_MODE_HYBRID2)) { char *dictfile = straight_ctx->dict; diff --git a/src/interface.c b/src/interface.c index d1ffc474a..24fa038ba 100644 --- a/src/interface.c +++ b/src/interface.c @@ -19,6 +19,7 @@ #include "cpu_sha256.h" #include "cpu_blake2.h" #include "shared.h" +#include "opencl.h" #include "interface.h" #include "ext_lzma.h" @@ -19626,6 +19627,7 @@ int ascii_digest (hashcat_ctx_t *hashcat_ctx, char *out_buf, const size_t out_le int hashconfig_init (hashcat_ctx_t *hashcat_ctx) { + folder_config_t *folder_config = hashcat_ctx->folder_config; hashconfig_t *hashconfig = hashcat_ctx->hashconfig; user_options_t *user_options = hashcat_ctx->user_options; user_options_extra_t *user_options_extra = hashcat_ctx->user_options_extra; @@ -24293,7 +24295,26 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx) hashconfig->opts_type |= OPTS_TYPE_PT_NEVERCRACK; } - if (user_options->optimized_kernel_enable == false) + // some kernels do not have an optimized kernel, simply because they do not need them + // or because they are not yet converted, for them we should switch off optimized mode + + if (user_options->optimized_kernel_enable == true) + { + char source_file[256] = { 0 }; + + generate_source_kernel_filename (hashconfig->attack_exec, user_options_extra->attack_kern, hashconfig->kern_type, user_options->optimized_kernel_enable, folder_config->shared_dir, source_file); + + if (hc_path_read (source_file) == false) + { + if (user_options->quiet == false) event_log_warning (hashcat_ctx, "%s: Optimized kernel not found, falling back to pure kernel", source_file); + } + else + { + hashconfig->opti_type |= OPTI_TYPE_OPTIMIZED_KERNEL; + } + } + + if ((hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) == 0) { hashconfig->opts_type &= ~OPTS_TYPE_PT_UTF16LE; hashconfig->opts_type &= ~OPTS_TYPE_PT_UTF16BE; @@ -24556,7 +24577,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx) // pw_max : some algo suffer from support for long passwords, // the user need to add -L to enable support for them - if (user_options->optimized_kernel_enable == true) + if (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) { hashconfig->pw_max = PW_MAX_OLD; @@ -24755,7 +24776,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx) hashconfig->salt_min = SALT_MIN; hashconfig->salt_max = SALT_MAX; - if (user_options->optimized_kernel_enable == true) + if (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) { hashconfig->salt_max = SALT_MAX_OLD; diff --git a/src/mpsp.c b/src/mpsp.c index 83f4ab1af..8cc0cc5c8 100644 --- a/src/mpsp.c +++ b/src/mpsp.c @@ -1174,7 +1174,7 @@ int mask_ctx_update_loop (hashcat_ctx_t *hashcat_ctx) } else if ((user_options->attack_mode == ATTACK_MODE_HYBRID1) || (user_options->attack_mode == ATTACK_MODE_HYBRID2)) { - if ((user_options->optimized_kernel_enable == false) && (user_options->attack_mode == ATTACK_MODE_HYBRID2)) + if (((hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) == 0) && (user_options->attack_mode == ATTACK_MODE_HYBRID2)) { mask_ctx->mask = mask_ctx->masks[mask_ctx->masks_pos]; diff --git a/src/opencl.c b/src/opencl.c index ac6c66c47..3f26e8468 100644 --- a/src/opencl.c +++ b/src/opencl.c @@ -108,112 +108,6 @@ static int ocl_check_dri (MAYBE_UNUSED hashcat_ctx_t *hashcat_ctx) return 0; } -static void generate_source_kernel_filename (const u32 attack_exec, const u32 attack_kern, const u32 kern_type, const bool optimized_kernel_enable, char *shared_dir, char *source_file) -{ - if (optimized_kernel_enable == true) - { - if (attack_exec == ATTACK_EXEC_INSIDE_KERNEL) - { - if (attack_kern == ATTACK_KERN_STRAIGHT) - snprintf (source_file, 255, "%s/OpenCL/m%05d_a0-optimized.cl", shared_dir, (int) kern_type); - else if (attack_kern == ATTACK_KERN_COMBI) - snprintf (source_file, 255, "%s/OpenCL/m%05d_a1-optimized.cl", shared_dir, (int) kern_type); - else if (attack_kern == ATTACK_KERN_BF) - snprintf (source_file, 255, "%s/OpenCL/m%05d_a3-optimized.cl", shared_dir, (int) kern_type); - } - else - { - snprintf (source_file, 255, "%s/OpenCL/m%05d-optimized.cl", shared_dir, (int) kern_type); - } - } - else - { - if (attack_exec == ATTACK_EXEC_INSIDE_KERNEL) - { - if (attack_kern == ATTACK_KERN_STRAIGHT) - snprintf (source_file, 255, "%s/OpenCL/m%05d_a0.cl", shared_dir, (int) kern_type); - else if (attack_kern == ATTACK_KERN_COMBI) - snprintf (source_file, 255, "%s/OpenCL/m%05d_a1.cl", shared_dir, (int) kern_type); - else if (attack_kern == ATTACK_KERN_BF) - snprintf (source_file, 255, "%s/OpenCL/m%05d_a3.cl", shared_dir, (int) kern_type); - } - else - { - snprintf (source_file, 255, "%s/OpenCL/m%05d.cl", shared_dir, (int) kern_type); - } - } -} - -static void generate_cached_kernel_filename (const u32 attack_exec, const u32 attack_kern, const u32 kern_type, const bool optimized_kernel_enable, char *profile_dir, const char *device_name_chksum, char *cached_file) -{ - if (optimized_kernel_enable == true) - { - if (attack_exec == ATTACK_EXEC_INSIDE_KERNEL) - { - if (attack_kern == ATTACK_KERN_STRAIGHT) - snprintf (cached_file, 255, "%s/kernels/m%05d_a0-optimized.%s.kernel", profile_dir, (int) kern_type, device_name_chksum); - else if (attack_kern == ATTACK_KERN_COMBI) - snprintf (cached_file, 255, "%s/kernels/m%05d_a1-optimized.%s.kernel", profile_dir, (int) kern_type, device_name_chksum); - else if (attack_kern == ATTACK_KERN_BF) - snprintf (cached_file, 255, "%s/kernels/m%05d_a3-optimized.%s.kernel", profile_dir, (int) kern_type, device_name_chksum); - } - else - { - snprintf (cached_file, 255, "%s/kernels/m%05d-optimized.%s.kernel", profile_dir, (int) kern_type, device_name_chksum); - } - } - else - { - if (attack_exec == ATTACK_EXEC_INSIDE_KERNEL) - { - if (attack_kern == ATTACK_KERN_STRAIGHT) - snprintf (cached_file, 255, "%s/kernels/m%05d_a0.%s.kernel", profile_dir, (int) kern_type, device_name_chksum); - else if (attack_kern == ATTACK_KERN_COMBI) - snprintf (cached_file, 255, "%s/kernels/m%05d_a1.%s.kernel", profile_dir, (int) kern_type, device_name_chksum); - else if (attack_kern == ATTACK_KERN_BF) - snprintf (cached_file, 255, "%s/kernels/m%05d_a3.%s.kernel", profile_dir, (int) kern_type, device_name_chksum); - } - else - { - snprintf (cached_file, 255, "%s/kernels/m%05d.%s.kernel", profile_dir, (int) kern_type, device_name_chksum); - } - } -} - -static void generate_source_kernel_mp_filename (const u32 opti_type, const u64 opts_type, char *shared_dir, char *source_file) -{ - if ((opti_type & OPTI_TYPE_BRUTE_FORCE) && (opts_type & OPTS_TYPE_PT_GENERATE_BE)) - { - snprintf (source_file, 255, "%s/OpenCL/markov_be.cl", shared_dir); - } - else - { - snprintf (source_file, 255, "%s/OpenCL/markov_le.cl", shared_dir); - } -} - -static void generate_cached_kernel_mp_filename (const u32 opti_type, const u64 opts_type, char *profile_dir, const char *device_name_chksum, char *cached_file) -{ - if ((opti_type & OPTI_TYPE_BRUTE_FORCE) && (opts_type & OPTS_TYPE_PT_GENERATE_BE)) - { - snprintf (cached_file, 255, "%s/kernels/markov_be.%s.kernel", profile_dir, device_name_chksum); - } - else - { - snprintf (cached_file, 255, "%s/kernels/markov_le.%s.kernel", profile_dir, device_name_chksum); - } -} - -static void generate_source_kernel_amp_filename (const u32 attack_kern, char *shared_dir, char *source_file) -{ - snprintf (source_file, 255, "%s/OpenCL/amp_a%u.cl", shared_dir, attack_kern); -} - -static void generate_cached_kernel_amp_filename (const u32 attack_kern, char *profile_dir, const char *device_name_chksum, char *cached_file) -{ - snprintf (cached_file, 255, "%s/kernels/amp_a%u.%s.kernel", profile_dir, attack_kern, device_name_chksum); -} - static int setup_opencl_platforms_filter (hashcat_ctx_t *hashcat_ctx, const char *opencl_platforms, u32 *out) { u32 opencl_platforms_filter = 0; @@ -424,6 +318,112 @@ static int write_kernel_binary (hashcat_ctx_t *hashcat_ctx, char *kernel_file, c return 0; } +void generate_source_kernel_filename (const u32 attack_exec, const u32 attack_kern, const u32 kern_type, const u32 opti_type, char *shared_dir, char *source_file) +{ + if (opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) + { + if (attack_exec == ATTACK_EXEC_INSIDE_KERNEL) + { + if (attack_kern == ATTACK_KERN_STRAIGHT) + snprintf (source_file, 255, "%s/OpenCL/m%05d_a0-optimized.cl", shared_dir, (int) kern_type); + else if (attack_kern == ATTACK_KERN_COMBI) + snprintf (source_file, 255, "%s/OpenCL/m%05d_a1-optimized.cl", shared_dir, (int) kern_type); + else if (attack_kern == ATTACK_KERN_BF) + snprintf (source_file, 255, "%s/OpenCL/m%05d_a3-optimized.cl", shared_dir, (int) kern_type); + } + else + { + snprintf (source_file, 255, "%s/OpenCL/m%05d-optimized.cl", shared_dir, (int) kern_type); + } + } + else + { + if (attack_exec == ATTACK_EXEC_INSIDE_KERNEL) + { + if (attack_kern == ATTACK_KERN_STRAIGHT) + snprintf (source_file, 255, "%s/OpenCL/m%05d_a0.cl", shared_dir, (int) kern_type); + else if (attack_kern == ATTACK_KERN_COMBI) + snprintf (source_file, 255, "%s/OpenCL/m%05d_a1.cl", shared_dir, (int) kern_type); + else if (attack_kern == ATTACK_KERN_BF) + snprintf (source_file, 255, "%s/OpenCL/m%05d_a3.cl", shared_dir, (int) kern_type); + } + else + { + snprintf (source_file, 255, "%s/OpenCL/m%05d.cl", shared_dir, (int) kern_type); + } + } +} + +void generate_cached_kernel_filename (const u32 attack_exec, const u32 attack_kern, const u32 kern_type, const u32 opti_type, char *profile_dir, const char *device_name_chksum, char *cached_file) +{ + if (opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) + { + if (attack_exec == ATTACK_EXEC_INSIDE_KERNEL) + { + if (attack_kern == ATTACK_KERN_STRAIGHT) + snprintf (cached_file, 255, "%s/kernels/m%05d_a0-optimized.%s.kernel", profile_dir, (int) kern_type, device_name_chksum); + else if (attack_kern == ATTACK_KERN_COMBI) + snprintf (cached_file, 255, "%s/kernels/m%05d_a1-optimized.%s.kernel", profile_dir, (int) kern_type, device_name_chksum); + else if (attack_kern == ATTACK_KERN_BF) + snprintf (cached_file, 255, "%s/kernels/m%05d_a3-optimized.%s.kernel", profile_dir, (int) kern_type, device_name_chksum); + } + else + { + snprintf (cached_file, 255, "%s/kernels/m%05d-optimized.%s.kernel", profile_dir, (int) kern_type, device_name_chksum); + } + } + else + { + if (attack_exec == ATTACK_EXEC_INSIDE_KERNEL) + { + if (attack_kern == ATTACK_KERN_STRAIGHT) + snprintf (cached_file, 255, "%s/kernels/m%05d_a0.%s.kernel", profile_dir, (int) kern_type, device_name_chksum); + else if (attack_kern == ATTACK_KERN_COMBI) + snprintf (cached_file, 255, "%s/kernels/m%05d_a1.%s.kernel", profile_dir, (int) kern_type, device_name_chksum); + else if (attack_kern == ATTACK_KERN_BF) + snprintf (cached_file, 255, "%s/kernels/m%05d_a3.%s.kernel", profile_dir, (int) kern_type, device_name_chksum); + } + else + { + snprintf (cached_file, 255, "%s/kernels/m%05d.%s.kernel", profile_dir, (int) kern_type, device_name_chksum); + } + } +} + +void generate_source_kernel_mp_filename (const u32 opti_type, const u64 opts_type, char *shared_dir, char *source_file) +{ + if ((opti_type & OPTI_TYPE_BRUTE_FORCE) && (opts_type & OPTS_TYPE_PT_GENERATE_BE)) + { + snprintf (source_file, 255, "%s/OpenCL/markov_be.cl", shared_dir); + } + else + { + snprintf (source_file, 255, "%s/OpenCL/markov_le.cl", shared_dir); + } +} + +void generate_cached_kernel_mp_filename (const u32 opti_type, const u64 opts_type, char *profile_dir, const char *device_name_chksum, char *cached_file) +{ + if ((opti_type & OPTI_TYPE_BRUTE_FORCE) && (opts_type & OPTS_TYPE_PT_GENERATE_BE)) + { + snprintf (cached_file, 255, "%s/kernels/markov_be.%s.kernel", profile_dir, device_name_chksum); + } + else + { + snprintf (cached_file, 255, "%s/kernels/markov_le.%s.kernel", profile_dir, device_name_chksum); + } +} + +void generate_source_kernel_amp_filename (const u32 attack_kern, char *shared_dir, char *source_file) +{ + snprintf (source_file, 255, "%s/OpenCL/amp_a%u.cl", shared_dir, attack_kern); +} + +void generate_cached_kernel_amp_filename (const u32 attack_kern, char *profile_dir, const char *device_name_chksum, char *cached_file) +{ + snprintf (cached_file, 255, "%s/kernels/amp_a%u.%s.kernel", profile_dir, attack_kern, device_name_chksum); +} + int ocl_init (hashcat_ctx_t *hashcat_ctx) { opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx; @@ -1165,7 +1165,7 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, } } - if (user_options->optimized_kernel_enable == true) + if (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) { if (highest_pw_len < 16) { @@ -1792,7 +1792,7 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const } else if (user_options_extra->attack_kern == ATTACK_KERN_COMBI) { - if (user_options->optimized_kernel_enable == true) + if (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) { if (user_options->attack_mode == ATTACK_MODE_COMBI) { @@ -1954,7 +1954,7 @@ int run_cracker (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, co FILE *combs_fp = device_param->combs_fp; - if ((user_options->attack_mode == ATTACK_MODE_COMBI) || ((user_options->optimized_kernel_enable == false) && (user_options->attack_mode == ATTACK_MODE_HYBRID2))) + if ((user_options->attack_mode == ATTACK_MODE_COMBI) || (((hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) == 0) && (user_options->attack_mode == ATTACK_MODE_HYBRID2))) { rewind (combs_fp); } @@ -2018,7 +2018,7 @@ int run_cracker (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, co } else if (user_options_extra->attack_kern == ATTACK_KERN_COMBI) { - if (user_options->optimized_kernel_enable == true) + if (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) { if (user_options->attack_mode == ATTACK_MODE_COMBI) { @@ -3730,7 +3730,7 @@ int opencl_session_begin (hashcat_ctx_t *hashcat_ctx) // We can't have SIMD in kernels where we have an unknown final password length - if (user_options->optimized_kernel_enable == false) + if ((hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) == 0) { if (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL) { @@ -4264,32 +4264,13 @@ int opencl_session_begin (hashcat_ctx_t *hashcat_ctx) char source_file[256] = { 0 }; - generate_source_kernel_filename (hashconfig->attack_exec, user_options_extra->attack_kern, hashconfig->kern_type, user_options->optimized_kernel_enable, folder_config->shared_dir, source_file); - - if (user_options->optimized_kernel_enable == true) - { - if (hc_path_read (source_file) == false) - { - if (user_options->quiet == false) event_log_warning (hashcat_ctx, "%s: Optimized kernel not found, falling back to pure kernel", source_file); - - generate_source_kernel_filename (hashconfig->attack_exec, user_options_extra->attack_kern, hashconfig->kern_type, false, folder_config->shared_dir, source_file); - } - - if (hc_path_read (source_file) == false) - { - event_log_error (hashcat_ctx, "%s: %s", source_file, strerror (errno)); + generate_source_kernel_filename (hashconfig->attack_exec, user_options_extra->attack_kern, hashconfig->kern_type, hashconfig->opti_type, folder_config->shared_dir, source_file); - return -1; - } - } - else + if (hc_path_read (source_file) == false) { - if (hc_path_read (source_file) == false) - { - event_log_error (hashcat_ctx, "%s: %s", source_file, strerror (errno)); + event_log_error (hashcat_ctx, "%s: %s", source_file, strerror (errno)); - return -1; - } + return -1; } /** @@ -4298,43 +4279,18 @@ int opencl_session_begin (hashcat_ctx_t *hashcat_ctx) char cached_file[256] = { 0 }; - generate_cached_kernel_filename (hashconfig->attack_exec, user_options_extra->attack_kern, hashconfig->kern_type, user_options->optimized_kernel_enable, folder_config->profile_dir, device_name_chksum, cached_file); + generate_cached_kernel_filename (hashconfig->attack_exec, user_options_extra->attack_kern, hashconfig->kern_type, hashconfig->opti_type, folder_config->profile_dir, device_name_chksum, cached_file); bool cached = true; - if (user_options->optimized_kernel_enable == true) + if (hc_path_read (cached_file) == false) { - if (hc_path_read (cached_file) == false) - { - generate_cached_kernel_filename (hashconfig->attack_exec, user_options_extra->attack_kern, hashconfig->kern_type, false, folder_config->profile_dir, device_name_chksum, cached_file); - } - - if (hc_path_is_empty (cached_file) == true) - { - generate_cached_kernel_filename (hashconfig->attack_exec, user_options_extra->attack_kern, hashconfig->kern_type, false, folder_config->profile_dir, device_name_chksum, cached_file); - } - - if (hc_path_read (cached_file) == false) - { - cached = false; - } - - if (hc_path_is_empty (cached_file) == true) - { - cached = false; - } + cached = false; } - else - { - if (hc_path_read (cached_file) == false) - { - cached = false; - } - if (hc_path_is_empty (cached_file) == true) - { - cached = false; - } + if (hc_path_is_empty (cached_file) == true) + { + cached = false; } /** @@ -4967,7 +4923,7 @@ int opencl_session_begin (hashcat_ctx_t *hashcat_ctx) device_param->kernel_params_mp_buf32[7] = 0; device_param->kernel_params_mp_buf32[8] = 0; - if (user_options->optimized_kernel_enable == true) + if (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) { device_param->kernel_params_mp[0] = &device_param->d_combs; } @@ -5059,7 +5015,7 @@ int opencl_session_begin (hashcat_ctx_t *hashcat_ctx) { if (hashconfig->opti_type & OPTI_TYPE_SINGLE_HASH) { - if (user_options->optimized_kernel_enable == true) + if (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) { // kernel1 @@ -5112,7 +5068,7 @@ int opencl_session_begin (hashcat_ctx_t *hashcat_ctx) } else { - if (user_options->optimized_kernel_enable == true) + if (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) { // kernel1 diff --git a/src/outfile.c b/src/outfile.c index a03051fc4..0915b4639 100644 --- a/src/outfile.c +++ b/src/outfile.c @@ -86,7 +86,7 @@ int build_plain (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, pl plain_len += comb_len; - if (user_options->optimized_kernel_enable == true) + if (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) { int pw_max_combi; @@ -152,7 +152,7 @@ int build_plain (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, pl } else if (user_options->attack_mode == ATTACK_MODE_HYBRID2) { - if (user_options->optimized_kernel_enable == true) + if (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) { pw_t pw; diff --git a/src/selftest.c b/src/selftest.c index 65315b88c..0568f1c74 100644 --- a/src/selftest.c +++ b/src/selftest.c @@ -18,7 +18,6 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param hashconfig_t *hashconfig = hashcat_ctx->hashconfig; hashes_t *hashes = hashcat_ctx->hashes; status_ctx_t *status_ctx = hashcat_ctx->status_ctx; - user_options_t *user_options = hashcat_ctx->user_options; user_options_extra_t *user_options_extra = hashcat_ctx->user_options_extra; cl_int CL_err; @@ -306,7 +305,7 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL) { - if (user_options->optimized_kernel_enable == true) + if (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) { if (highest_pw_len < 16) { diff --git a/src/status.c b/src/status.c index 3e664c1f3..26b191b60 100644 --- a/src/status.c +++ b/src/status.c @@ -435,6 +435,7 @@ int status_get_guess_mode (const hashcat_ctx_t *hashcat_ctx) char *status_get_guess_base (const hashcat_ctx_t *hashcat_ctx) { + const hashconfig_t *hashconfig = hashcat_ctx->hashconfig; const user_options_t *user_options = hashcat_ctx->user_options; const user_options_extra_t *user_options_extra = hashcat_ctx->user_options_extra; @@ -474,7 +475,7 @@ char *status_get_guess_base (const hashcat_ctx_t *hashcat_ctx) } else if (user_options->attack_mode == ATTACK_MODE_HYBRID2) { - if (user_options->optimized_kernel_enable == true) + if (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) { const mask_ctx_t *mask_ctx = hashcat_ctx->mask_ctx; @@ -493,6 +494,7 @@ char *status_get_guess_base (const hashcat_ctx_t *hashcat_ctx) int status_get_guess_base_offset (const hashcat_ctx_t *hashcat_ctx) { + const hashconfig_t *hashconfig = hashcat_ctx->hashconfig; const user_options_t *user_options = hashcat_ctx->user_options; if (user_options->attack_mode == ATTACK_MODE_STRAIGHT) @@ -519,7 +521,7 @@ int status_get_guess_base_offset (const hashcat_ctx_t *hashcat_ctx) } else if (user_options->attack_mode == ATTACK_MODE_HYBRID2) { - if (user_options->optimized_kernel_enable == true) + if (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) { const mask_ctx_t *mask_ctx = hashcat_ctx->mask_ctx; @@ -538,6 +540,7 @@ int status_get_guess_base_offset (const hashcat_ctx_t *hashcat_ctx) int status_get_guess_base_count (const hashcat_ctx_t *hashcat_ctx) { + const hashconfig_t *hashconfig = hashcat_ctx->hashconfig; const user_options_t *user_options = hashcat_ctx->user_options; if (user_options->attack_mode == ATTACK_MODE_STRAIGHT) @@ -564,7 +567,7 @@ int status_get_guess_base_count (const hashcat_ctx_t *hashcat_ctx) } else if (user_options->attack_mode == ATTACK_MODE_HYBRID2) { - if (user_options->optimized_kernel_enable == true) + if (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) { const mask_ctx_t *mask_ctx = hashcat_ctx->mask_ctx; @@ -593,6 +596,7 @@ double status_get_guess_base_percent (const hashcat_ctx_t *hashcat_ctx) char *status_get_guess_mod (const hashcat_ctx_t *hashcat_ctx) { + const hashconfig_t *hashconfig = hashcat_ctx->hashconfig; const user_options_t *user_options = hashcat_ctx->user_options; if (user_options->attack_mode == ATTACK_MODE_STRAIGHT) @@ -624,7 +628,7 @@ char *status_get_guess_mod (const hashcat_ctx_t *hashcat_ctx) } else if (user_options->attack_mode == ATTACK_MODE_HYBRID2) { - if (user_options->optimized_kernel_enable == true) + if (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) { const straight_ctx_t *straight_ctx = hashcat_ctx->straight_ctx; @@ -643,6 +647,7 @@ char *status_get_guess_mod (const hashcat_ctx_t *hashcat_ctx) int status_get_guess_mod_offset (const hashcat_ctx_t *hashcat_ctx) { + const hashconfig_t *hashconfig = hashcat_ctx->hashconfig; const user_options_t *user_options = hashcat_ctx->user_options; if (user_options->attack_mode == ATTACK_MODE_STRAIGHT) @@ -665,7 +670,7 @@ int status_get_guess_mod_offset (const hashcat_ctx_t *hashcat_ctx) } else if (user_options->attack_mode == ATTACK_MODE_HYBRID2) { - if (user_options->optimized_kernel_enable == true) + if (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) { const straight_ctx_t *straight_ctx = hashcat_ctx->straight_ctx; @@ -684,6 +689,7 @@ int status_get_guess_mod_offset (const hashcat_ctx_t *hashcat_ctx) int status_get_guess_mod_count (const hashcat_ctx_t *hashcat_ctx) { + const hashconfig_t *hashconfig = hashcat_ctx->hashconfig; const user_options_t *user_options = hashcat_ctx->user_options; if (user_options->attack_mode == ATTACK_MODE_STRAIGHT) @@ -706,7 +712,7 @@ int status_get_guess_mod_count (const hashcat_ctx_t *hashcat_ctx) } else if (user_options->attack_mode == ATTACK_MODE_HYBRID2) { - if (user_options->optimized_kernel_enable == true) + if (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) { const straight_ctx_t *straight_ctx = hashcat_ctx->straight_ctx; diff --git a/src/terminal.c b/src/terminal.c index 4264cf7f9..8588f7643 100644 --- a/src/terminal.c +++ b/src/terminal.c @@ -699,6 +699,7 @@ void status_display_machine_readable (hashcat_ctx_t *hashcat_ctx) void status_display (hashcat_ctx_t *hashcat_ctx) { + const hashconfig_t *hashconfig = hashcat_ctx->hashconfig; const user_options_t *user_options = hashcat_ctx->user_options; if (user_options->machine_readable == true) @@ -888,7 +889,7 @@ void status_display (hashcat_ctx_t *hashcat_ctx) case GUESS_MODE_HYBRID2: - if (user_options->optimized_kernel_enable == true) + if (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) { event_log_info (hashcat_ctx, "Guess.Base.......: Mask (%s) [%d], Left Side", @@ -915,7 +916,7 @@ void status_display (hashcat_ctx_t *hashcat_ctx) case GUESS_MODE_HYBRID2_CS: - if (user_options->optimized_kernel_enable == true) + if (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) { event_log_info (hashcat_ctx, "Guess.Base.......: Mask (%s) [%d], Left Side", diff --git a/src/weak_hash.c b/src/weak_hash.c index 9396fbfe1..f844e951f 100644 --- a/src/weak_hash.c +++ b/src/weak_hash.c @@ -15,7 +15,6 @@ int weak_hash_check (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param hashconfig_t *hashconfig = hashcat_ctx->hashconfig; hashes_t *hashes = hashcat_ctx->hashes; straight_ctx_t *straight_ctx = hashcat_ctx->straight_ctx; - user_options_t *user_options = hashcat_ctx->user_options; salt_t *salt_buf = &hashes->salts_buf[salt_pos]; @@ -38,7 +37,7 @@ int weak_hash_check (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL) { - if (user_options->optimized_kernel_enable == true) + if (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) { CL_rc = run_kernel (hashcat_ctx, device_param, KERN_RUN_1, 1, false, 0); diff --git a/src/wordlist.c b/src/wordlist.c index 5a5ecfbac..a95b4fb74 100644 --- a/src/wordlist.c +++ b/src/wordlist.c @@ -261,6 +261,7 @@ void pw_add (hc_device_param_t *device_param, const u8 *pw_buf, const int pw_len int count_words (hashcat_ctx_t *hashcat_ctx, FILE *fd, const char *dictfile, u64 *result) { combinator_ctx_t *combinator_ctx = hashcat_ctx->combinator_ctx; + hashconfig_t *hashconfig = hashcat_ctx->hashconfig; straight_ctx_t *straight_ctx = hashcat_ctx->straight_ctx; mask_ctx_t *mask_ctx = hashcat_ctx->mask_ctx; user_options_extra_t *user_options_extra = hashcat_ctx->user_options_extra; @@ -321,7 +322,7 @@ int count_words (hashcat_ctx_t *hashcat_ctx, FILE *fd, const char *dictfile, u64 } else if (user_options_extra->attack_kern == ATTACK_KERN_COMBI) { - if ((user_options->optimized_kernel_enable == false) && (user_options->attack_mode == ATTACK_MODE_HYBRID2)) + if (((hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) == 0) && (user_options->attack_mode == ATTACK_MODE_HYBRID2)) { if (overflow_check_u64_mul (keyspace, mask_ctx->bfs_cnt) == false) return -1; @@ -424,7 +425,7 @@ int count_words (hashcat_ctx_t *hashcat_ctx, FILE *fd, const char *dictfile, u64 } else if (user_options_extra->attack_kern == ATTACK_KERN_COMBI) { - if ((user_options->optimized_kernel_enable == false) && (user_options->attack_mode == ATTACK_MODE_HYBRID2)) + if (((hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) == 0) && (user_options->attack_mode == ATTACK_MODE_HYBRID2)) { if (overflow_check_u64_add (cnt, mask_ctx->bfs_cnt) == false) return -1; diff --git a/tools/test.sh b/tools/test.sh index 403633cec..1606864c1 100755 --- a/tools/test.sh +++ b/tools/test.sh @@ -24,7 +24,7 @@ NEVER_CRACK="11600 14900" SLOW_ALGOS="400 500 501 1600 1800 2100 2500 3200 5200 5800 6211 6212 6213 6221 6222 6223 6231 6232 6233 6241 6242 6243 6251 6261 6271 6281 6300 6400 6500 6600 6700 6800 7100 7200 7400 7900 8200 8800 8900 9000 9100 9200 9300 9400 9500 9600 10000 10300 10500 10700 10900 11300 11600 11900 12000 12001 12100 12200 12300 12400 12500 12700 12800 12900 13000 13200 13400 13600 14600 14700 14800 15100 15200 15300 15600 15700" -OPTS="--quiet --force --potfile-disable --runtime 400 --gpu-temp-disable --weak-hash-threshold=0" +OPTS="--quiet --force --potfile-disable --runtime 400 --gpu-temp-disable --weak-hash-threshold=0 -O" OUTD="test_$(date +%s)"