Move files from include/ to OpenCL/ if they are used within kernels

Rename includes in OpenCL so that it's easier to recognize them as such
pull/352/merge
Jens Steube 8 years ago
parent 083c8ed515
commit 2899f53a15

@ -3,12 +3,12 @@
* License.....: MIT * License.....: MIT
*/ */
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "include/rp_kernel.h" #include "inc_rp.h"
#include "OpenCL/rp.c" #include "inc_rp.cl"
__kernel void amp (__global pw_t *pws, __global pw_t *pws_amp, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, const u32 combs_mode, const u32 gid_max) __kernel void amp (__global pw_t *pws, __global pw_t *pws_amp, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, const u32 combs_mode, const u32 gid_max)
{ {

@ -3,9 +3,9 @@
* License.....: MIT * License.....: MIT
*/ */
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
inline void switch_buffer_by_offset_le (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 offset) inline void switch_buffer_by_offset_le (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 offset)
{ {

@ -3,9 +3,9 @@
* License.....: MIT * License.....: MIT
*/ */
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
__kernel void amp (__global pw_t *pws, __global pw_t *pws_amp, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, const u32 combs_mode, const u32 gid_max) __kernel void amp (__global pw_t *pws, __global pw_t *pws_amp, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, const u32 combs_mode, const u32 gid_max)
{ {

@ -7,20 +7,20 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 0 #define DGST_R0 0
#define DGST_R1 3 #define DGST_R1 3
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "include/rp_kernel.h" #include "inc_rp.h"
#include "OpenCL/rp.c" #include "inc_rp.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
__kernel void m00000_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) __kernel void m00000_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{ {

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 0 #define DGST_R0 0
#define DGST_R1 3 #define DGST_R1 3
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
__kernel void m00000_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) __kernel void m00000_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{ {

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 0 #define DGST_R0 0
#define DGST_R1 3 #define DGST_R1 3
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
#define MD5_STEP_REV(f,a,b,c,d,x,t,s) \ #define MD5_STEP_REV(f,a,b,c,d,x,t,s) \
{ \ { \

@ -7,20 +7,20 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 0 #define DGST_R0 0
#define DGST_R1 3 #define DGST_R1 3
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "include/rp_kernel.h" #include "inc_rp.h"
#include "OpenCL/rp.c" #include "inc_rp.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
__kernel void m00010_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) __kernel void m00010_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{ {

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 0 #define DGST_R0 0
#define DGST_R1 3 #define DGST_R1 3
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
__kernel void m00010_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) __kernel void m00010_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{ {

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 0 #define DGST_R0 0
#define DGST_R1 3 #define DGST_R1 3
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
#define MD5_STEP_REV(f,a,b,c,d,x,t,s) \ #define MD5_STEP_REV(f,a,b,c,d,x,t,s) \
{ \ { \

@ -7,20 +7,20 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 0 #define DGST_R0 0
#define DGST_R1 3 #define DGST_R1 3
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "include/rp_kernel.h" #include "inc_rp.h"
#include "OpenCL/rp.c" #include "inc_rp.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
__kernel void m00020_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) __kernel void m00020_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{ {

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 0 #define DGST_R0 0
#define DGST_R1 3 #define DGST_R1 3
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
__kernel void m00020_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) __kernel void m00020_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{ {

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 0 #define DGST_R0 0
#define DGST_R1 3 #define DGST_R1 3
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
void m00020m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset) void m00020m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset)
{ {

@ -7,20 +7,20 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 0 #define DGST_R0 0
#define DGST_R1 3 #define DGST_R1 3
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "include/rp_kernel.h" #include "inc_rp.h"
#include "OpenCL/rp.c" #include "inc_rp.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
// no unicode yet // no unicode yet

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 0 #define DGST_R0 0
#define DGST_R1 3 #define DGST_R1 3
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
// no unicode yet // no unicode yet

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 0 #define DGST_R0 0
#define DGST_R1 3 #define DGST_R1 3
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
#define MD5_STEP_REV(f,a,b,c,d,x,t,s) \ #define MD5_STEP_REV(f,a,b,c,d,x,t,s) \
{ \ { \

@ -7,20 +7,20 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 0 #define DGST_R0 0
#define DGST_R1 3 #define DGST_R1 3
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "include/rp_kernel.h" #include "inc_rp.h"
#include "OpenCL/rp.c" #include "inc_rp.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
__kernel void m00040_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) __kernel void m00040_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{ {

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 0 #define DGST_R0 0
#define DGST_R1 3 #define DGST_R1 3
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
// no unicode yet // no unicode yet

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 0 #define DGST_R0 0
#define DGST_R1 3 #define DGST_R1 3
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
void m00040m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset) void m00040m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset)
{ {

@ -7,20 +7,20 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 0 #define DGST_R0 0
#define DGST_R1 3 #define DGST_R1 3
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "include/rp_kernel.h" #include "inc_rp.h"
#include "OpenCL/rp.c" #include "inc_rp.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
void md5_transform (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[4]) void md5_transform (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[4])
{ {

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 0 #define DGST_R0 0
#define DGST_R1 3 #define DGST_R1 3
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
void md5_transform (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[4]) void md5_transform (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[4])
{ {

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 0 #define DGST_R0 0
#define DGST_R1 3 #define DGST_R1 3
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
void md5_transform (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[4]) void md5_transform (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[4])
{ {

@ -7,20 +7,20 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 0 #define DGST_R0 0
#define DGST_R1 3 #define DGST_R1 3
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "include/rp_kernel.h" #include "inc_rp.h"
#include "OpenCL/rp.c" #include "inc_rp.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
void md5_transform (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[4]) void md5_transform (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[4])
{ {

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 0 #define DGST_R0 0
#define DGST_R1 3 #define DGST_R1 3
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
void md5_transform (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[4]) void md5_transform (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[4])
{ {

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 0 #define DGST_R0 0
#define DGST_R1 3 #define DGST_R1 3
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
void md5_transform (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[4]) void md5_transform (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[4])
{ {

@ -7,20 +7,20 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 3 #define DGST_R0 3
#define DGST_R1 4 #define DGST_R1 4
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "include/rp_kernel.h" #include "inc_rp.h"
#include "OpenCL/rp.c" #include "inc_rp.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
__kernel void m00100_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) __kernel void m00100_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{ {

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 3 #define DGST_R0 3
#define DGST_R1 4 #define DGST_R1 4
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
__kernel void m00100_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) __kernel void m00100_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{ {

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 3 #define DGST_R0 3
#define DGST_R1 4 #define DGST_R1 4
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
void m00100m (u32 w[16], const u32 pw_len, __global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32x * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset) void m00100m (u32 w[16], const u32 pw_len, __global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32x * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset)
{ {

@ -7,20 +7,20 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 3 #define DGST_R0 3
#define DGST_R1 4 #define DGST_R1 4
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "include/rp_kernel.h" #include "inc_rp.h"
#include "OpenCL/rp.c" #include "inc_rp.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
__kernel void m00110_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) __kernel void m00110_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{ {

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 3 #define DGST_R0 3
#define DGST_R1 4 #define DGST_R1 4
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
__kernel void m00110_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) __kernel void m00110_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{ {

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 3 #define DGST_R0 3
#define DGST_R1 4 #define DGST_R1 4
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
void m00110m (u32 w[16], const u32 pw_len, __global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32x * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset) void m00110m (u32 w[16], const u32 pw_len, __global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32x * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset)
{ {

@ -7,20 +7,20 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 3 #define DGST_R0 3
#define DGST_R1 4 #define DGST_R1 4
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "include/rp_kernel.h" #include "inc_rp.h"
#include "OpenCL/rp.c" #include "inc_rp.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
__kernel void m00120_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) __kernel void m00120_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{ {

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 3 #define DGST_R0 3
#define DGST_R1 4 #define DGST_R1 4
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
__kernel void m00120_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) __kernel void m00120_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{ {

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 3 #define DGST_R0 3
#define DGST_R1 4 #define DGST_R1 4
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
void m00120m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset) void m00120m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset)
{ {

@ -7,20 +7,20 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 3 #define DGST_R0 3
#define DGST_R1 4 #define DGST_R1 4
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "include/rp_kernel.h" #include "inc_rp.h"
#include "OpenCL/rp.c" #include "inc_rp.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
// no unicode yet // no unicode yet

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 3 #define DGST_R0 3
#define DGST_R1 4 #define DGST_R1 4
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
// no unicode yet // no unicode yet

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 3 #define DGST_R0 3
#define DGST_R1 4 #define DGST_R1 4
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
void m00130m (u32 w[16], const u32 pw_len, __global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32x * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset) void m00130m (u32 w[16], const u32 pw_len, __global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32x * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset)
{ {

@ -7,20 +7,20 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 3 #define DGST_R0 3
#define DGST_R1 4 #define DGST_R1 4
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "include/rp_kernel.h" #include "inc_rp.h"
#include "OpenCL/rp.c" #include "inc_rp.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
// no unicode yet // no unicode yet

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 3 #define DGST_R0 3
#define DGST_R1 4 #define DGST_R1 4
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
// no unicode yet // no unicode yet

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 3 #define DGST_R0 3
#define DGST_R1 4 #define DGST_R1 4
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
void m00140m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset) void m00140m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset)
{ {

@ -7,20 +7,20 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 3 #define DGST_R0 3
#define DGST_R1 4 #define DGST_R1 4
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "include/rp_kernel.h" #include "inc_rp.h"
#include "OpenCL/rp.c" #include "inc_rp.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
void sha1_transform (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[5]) void sha1_transform (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[5])
{ {

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 3 #define DGST_R0 3
#define DGST_R1 4 #define DGST_R1 4
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
void sha1_transform (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[5]) void sha1_transform (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[5])
{ {

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 3 #define DGST_R0 3
#define DGST_R1 4 #define DGST_R1 4
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
void sha1_transform (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[5]) void sha1_transform (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[5])
{ {

@ -7,20 +7,20 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 3 #define DGST_R0 3
#define DGST_R1 4 #define DGST_R1 4
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "include/rp_kernel.h" #include "inc_rp.h"
#include "OpenCL/rp.c" #include "inc_rp.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
void sha1_transform (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[5]) void sha1_transform (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[5])
{ {

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 3 #define DGST_R0 3
#define DGST_R1 4 #define DGST_R1 4
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
void sha1_transform (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[5]) void sha1_transform (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[5])
{ {

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 3 #define DGST_R0 3
#define DGST_R1 4 #define DGST_R1 4
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
void sha1_transform (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[5]) void sha1_transform (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[5])
{ {

@ -7,20 +7,20 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 0 #define DGST_R0 0
#define DGST_R1 4 #define DGST_R1 4
#define DGST_R2 3 #define DGST_R2 3
#define DGST_R3 2 #define DGST_R3 2
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "include/rp_kernel.h" #include "inc_rp.h"
#include "OpenCL/rp.c" #include "inc_rp.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
__kernel void m00190_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) __kernel void m00190_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{ {

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 0 #define DGST_R0 0
#define DGST_R1 4 #define DGST_R1 4
#define DGST_R2 3 #define DGST_R2 3
#define DGST_R3 2 #define DGST_R3 2
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
__kernel void m00190_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) __kernel void m00190_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{ {

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 0 #define DGST_R0 0
#define DGST_R1 4 #define DGST_R1 4
#define DGST_R2 3 #define DGST_R2 3
#define DGST_R3 2 #define DGST_R3 2
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
void m00190m (u32 w[16], const u32 pw_len, __global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32x * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset) void m00190m (u32 w[16], const u32 pw_len, __global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32x * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset)
{ {

@ -8,20 +8,20 @@
//incompatible //incompatible
//#define NEW_SIMD_CODE //#define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 0 #define DGST_R0 0
#define DGST_R1 1 #define DGST_R1 1
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 3 #define DGST_R3 3
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "include/rp_kernel.h" #include "inc_rp.h"
#include "OpenCL/rp.c" #include "inc_rp.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
__kernel void m00200_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) __kernel void m00200_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{ {

@ -8,18 +8,18 @@
//incompatible //incompatible
//#define NEW_SIMD_CODE //#define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 0 #define DGST_R0 0
#define DGST_R1 1 #define DGST_R1 1
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 3 #define DGST_R3 3
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
__kernel void m00200_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) __kernel void m00200_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{ {

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 0 #define DGST_R0 0
#define DGST_R1 1 #define DGST_R1 1
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 3 #define DGST_R3 3
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
#define ROUND(v) \ #define ROUND(v) \
{ \ { \

@ -7,20 +7,20 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 3 #define DGST_R0 3
#define DGST_R1 4 #define DGST_R1 4
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "include/rp_kernel.h" #include "inc_rp.h"
#include "OpenCL/rp.c" #include "inc_rp.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
__kernel void m00300_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) __kernel void m00300_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{ {

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 3 #define DGST_R0 3
#define DGST_R1 4 #define DGST_R1 4
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
__kernel void m00300_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) __kernel void m00300_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{ {

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 3 #define DGST_R0 3
#define DGST_R1 4 #define DGST_R1 4
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
void m00300m (u32 w[16], const u32 pw_len, __global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32x * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset) void m00300m (u32 w[16], const u32 pw_len, __global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32x * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset)
{ {

@ -7,21 +7,21 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 0 #define DGST_R0 0
#define DGST_R1 1 #define DGST_R1 1
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 3 #define DGST_R3 3
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
#define COMPARE_S "OpenCL/check_single_comp4.c" #define COMPARE_S "inc_comp_single.cl"
#define COMPARE_M "OpenCL/check_multi_comp4.c" #define COMPARE_M "inc_comp_multi.cl"
void md5_transform_S (const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], u32 digest[4]) void md5_transform_S (const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], u32 digest[4])
{ {

@ -5,20 +5,20 @@
#define _MD5_ #define _MD5_
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 0 #define DGST_R0 0
#define DGST_R1 1 #define DGST_R1 1
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 3 #define DGST_R3 3
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#define COMPARE_S "OpenCL/check_single_comp4.c" #define COMPARE_S "inc_comp_single.cl"
#define COMPARE_M "OpenCL/check_multi_comp4.c" #define COMPARE_M "inc_comp_multi.cl"
#define md5crypt_magic 0x00243124u #define md5crypt_magic 0x00243124u

@ -7,20 +7,20 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 0 #define DGST_R0 0
#define DGST_R1 3 #define DGST_R1 3
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "include/rp_kernel.h" #include "inc_rp.h"
#include "OpenCL/rp.c" #include "inc_rp.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
__kernel void m00900_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) __kernel void m00900_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{ {

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 0 #define DGST_R0 0
#define DGST_R1 3 #define DGST_R1 3
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
__kernel void m00900_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) __kernel void m00900_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{ {

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 0 #define DGST_R0 0
#define DGST_R1 3 #define DGST_R1 3
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
#define MD4_STEP_REV(f,a,b,c,d,x,t,s) \ #define MD4_STEP_REV(f,a,b,c,d,x,t,s) \
{ \ { \

@ -7,20 +7,20 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 0 #define DGST_R0 0
#define DGST_R1 3 #define DGST_R1 3
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "include/rp_kernel.h" #include "inc_rp.h"
#include "OpenCL/rp.c" #include "inc_rp.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
__kernel void m01000_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) __kernel void m01000_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{ {

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 0 #define DGST_R0 0
#define DGST_R1 3 #define DGST_R1 3
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
__kernel void m01000_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) __kernel void m01000_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{ {

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 0 #define DGST_R0 0
#define DGST_R1 3 #define DGST_R1 3
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
#define MD4_STEP_REV(f,a,b,c,d,x,t,s) \ #define MD4_STEP_REV(f,a,b,c,d,x,t,s) \
{ \ { \

@ -7,20 +7,20 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 0 #define DGST_R0 0
#define DGST_R1 3 #define DGST_R1 3
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "include/rp_kernel.h" #include "inc_rp.h"
#include "OpenCL/rp.c" #include "inc_rp.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
__kernel void m01100_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) __kernel void m01100_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{ {

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 0 #define DGST_R0 0
#define DGST_R1 3 #define DGST_R1 3
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
__kernel void m01100_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) __kernel void m01100_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{ {

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 0 #define DGST_R0 0
#define DGST_R1 3 #define DGST_R1 3
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 1 #define DGST_R3 1
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
void m01100m (u32 w[16], const u32 pw_len, __global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32x * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset) void m01100m (u32 w[16], const u32 pw_len, __global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32x * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset)
{ {

@ -7,20 +7,20 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 3 #define DGST_R0 3
#define DGST_R1 7 #define DGST_R1 7
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 6 #define DGST_R3 6
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "include/rp_kernel.h" #include "inc_rp.h"
#include "OpenCL/rp.c" #include "inc_rp.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
__kernel void m01400_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) __kernel void m01400_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{ {

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 3 #define DGST_R0 3
#define DGST_R1 7 #define DGST_R1 7
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 6 #define DGST_R3 6
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
__kernel void m01400_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) __kernel void m01400_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{ {

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 3 #define DGST_R0 3
#define DGST_R1 7 #define DGST_R1 7
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 6 #define DGST_R3 6
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
void m01400m (u32 w[16], const u32 pw_len, __global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32x * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset) void m01400m (u32 w[16], const u32 pw_len, __global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32x * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset)
{ {

@ -7,20 +7,20 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 3 #define DGST_R0 3
#define DGST_R1 7 #define DGST_R1 7
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 6 #define DGST_R3 6
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "include/rp_kernel.h" #include "inc_rp.h"
#include "OpenCL/rp.c" #include "inc_rp.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
__kernel void m01410_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) __kernel void m01410_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{ {

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 3 #define DGST_R0 3
#define DGST_R1 7 #define DGST_R1 7
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 6 #define DGST_R3 6
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
__kernel void m01410_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) __kernel void m01410_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{ {

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 3 #define DGST_R0 3
#define DGST_R1 7 #define DGST_R1 7
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 6 #define DGST_R3 6
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
void m01410m (u32 w[16], const u32 pw_len, __global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32x * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset) void m01410m (u32 w[16], const u32 pw_len, __global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32x * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset)
{ {

@ -7,20 +7,20 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 3 #define DGST_R0 3
#define DGST_R1 7 #define DGST_R1 7
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 6 #define DGST_R3 6
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "include/rp_kernel.h" #include "inc_rp.h"
#include "OpenCL/rp.c" #include "inc_rp.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
__kernel void m01420_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) __kernel void m01420_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{ {

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 3 #define DGST_R0 3
#define DGST_R1 7 #define DGST_R1 7
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 6 #define DGST_R3 6
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
__kernel void m01420_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) __kernel void m01420_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{ {

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 3 #define DGST_R0 3
#define DGST_R1 7 #define DGST_R1 7
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 6 #define DGST_R3 6
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
void m01420m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset) void m01420m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset)
{ {

@ -7,20 +7,20 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 3 #define DGST_R0 3
#define DGST_R1 7 #define DGST_R1 7
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 6 #define DGST_R3 6
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "include/rp_kernel.h" #include "inc_rp.h"
#include "OpenCL/rp.c" #include "inc_rp.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
__kernel void m01430_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) __kernel void m01430_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{ {

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 3 #define DGST_R0 3
#define DGST_R1 7 #define DGST_R1 7
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 6 #define DGST_R3 6
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
__kernel void m01430_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) __kernel void m01430_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{ {

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 3 #define DGST_R0 3
#define DGST_R1 7 #define DGST_R1 7
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 6 #define DGST_R3 6
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
void m01430m (u32 w[16], const u32 pw_len, __global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32x * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset) void m01430m (u32 w[16], const u32 pw_len, __global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32x * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset)
{ {

@ -7,20 +7,20 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 3 #define DGST_R0 3
#define DGST_R1 7 #define DGST_R1 7
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 6 #define DGST_R3 6
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "include/rp_kernel.h" #include "inc_rp.h"
#include "OpenCL/rp.c" #include "inc_rp.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
__kernel void m01440_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) __kernel void m01440_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{ {

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 3 #define DGST_R0 3
#define DGST_R1 7 #define DGST_R1 7
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 6 #define DGST_R3 6
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
__kernel void m01440_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) __kernel void m01440_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{ {

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 3 #define DGST_R0 3
#define DGST_R1 7 #define DGST_R1 7
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 6 #define DGST_R3 6
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
void m01440m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset) void m01440m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset)
{ {

@ -7,20 +7,20 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 3 #define DGST_R0 3
#define DGST_R1 7 #define DGST_R1 7
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 6 #define DGST_R3 6
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "include/rp_kernel.h" #include "inc_rp.h"
#include "OpenCL/rp.c" #include "inc_rp.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
__constant u32 k_sha256[64] = __constant u32 k_sha256[64] =
{ {

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 3 #define DGST_R0 3
#define DGST_R1 7 #define DGST_R1 7
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 6 #define DGST_R3 6
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
__constant u32 k_sha256[64] = __constant u32 k_sha256[64] =
{ {

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 3 #define DGST_R0 3
#define DGST_R1 7 #define DGST_R1 7
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 6 #define DGST_R3 6
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
__constant u32 k_sha256[64] = __constant u32 k_sha256[64] =
{ {

@ -7,20 +7,20 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 3 #define DGST_R0 3
#define DGST_R1 7 #define DGST_R1 7
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 6 #define DGST_R3 6
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "include/rp_kernel.h" #include "inc_rp.h"
#include "OpenCL/rp.c" #include "inc_rp.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
__constant u32 k_sha256[64] = __constant u32 k_sha256[64] =
{ {

@ -7,18 +7,18 @@
#define NEW_SIMD_CODE #define NEW_SIMD_CODE
#include "include/constants.h" #include "inc_hash_constants.h"
#include "include/kernel_vendor.h" #include "inc_vendor.cl"
#define DGST_R0 3 #define DGST_R0 3
#define DGST_R1 7 #define DGST_R1 7
#define DGST_R2 2 #define DGST_R2 2
#define DGST_R3 6 #define DGST_R3 6
#include "include/kernel_functions.c" #include "inc_hash_functions.cl"
#include "OpenCL/types_ocl.c" #include "inc_types.cl"
#include "OpenCL/common.c" #include "inc_common.cl"
#include "OpenCL/simd.c" #include "inc_simd.cl"
__constant u32 k_sha256[64] = __constant u32 k_sha256[64] =
{ {

Some files were not shown because too many files have changed in this diff Show More

Loading…
Cancel
Save