X-Git-Url: https://www.flypig.org.uk/git/?a=blobdiff_plain;f=OpenCL%2Fmarkov_le.cl;h=27294464d6919a61014b5ba3bd0092d13bfade37;hb=161a6eb4bc643d8e636e96eda613f5137d30da59;hp=e56a9097ccbd72a18da62457bf20ff2926c71e1b;hpb=e4465d7dcb88426532f6b0488cf9bc1cd3b06cc6;p=hashcat.git diff --git a/OpenCL/markov_le.cl b/OpenCL/markov_le.cl index e56a909..2729446 100644 --- a/OpenCL/markov_le.cl +++ b/OpenCL/markov_le.cl @@ -3,13 +3,13 @@ * License.....: MIT */ -#include "include/kernel_vendor.h" +#include "inc_vendor.cl" #define CHARSIZ 256 -#include "types_ocl.c" +#include "inc_types.cl" -static void generate_pw (u32 pw_buf[16], __global cs_t *root_css_buf, __global cs_t *markov_css_buf, const u32 pw_l_len, const u32 pw_r_len, const u32 mask80, const u32 bits14, const u32 bits15, u64 val) +inline void generate_pw (u32 pw_buf[16], __global cs_t *root_css_buf, __global cs_t *markov_css_buf, const u32 pw_l_len, const u32 pw_r_len, const u32 mask80, const u32 bits14, const u32 bits15, u64 val) { pw_buf[ 0] = 0; pw_buf[ 1] = 0; @@ -61,7 +61,7 @@ static void generate_pw (u32 pw_buf[16], __global cs_t *root_css_buf, __global c if (bits15) pw_buf[15] = (pw_l_len + pw_r_len) * 8; } -__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) l_markov (__global pw_t *pws_buf_l, __global cs_t *root_css_buf, __global cs_t *markov_css_buf, const u64 off, const u32 pw_l_len, const u32 pw_r_len, const u32 mask80, const u32 bits14, const u32 bits15, const u32 gid_max) +__kernel void l_markov (__global pw_t *pws_buf_l, __global cs_t *root_css_buf, __global cs_t *markov_css_buf, const u64 off, const u32 pw_l_len, const u32 pw_r_len, const u32 mask80, const u32 bits14, const u32 bits15, const u32 gid_max) { const u32 gid = get_global_id (0); @@ -91,7 +91,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) l_markov (__globa pws_buf_l[gid].pw_len = pw_l_len + pw_r_len; } -__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) r_markov (__global bf_t *pws_buf_r, __global cs_t *root_css_buf, __global cs_t *markov_css_buf, const u64 off, const u32 pw_r_len, const u32 mask80, const u32 bits14, const u32 bits15, const u32 gid_max) +__kernel void r_markov (__global bf_t *pws_buf_r, __global cs_t *root_css_buf, __global cs_t *markov_css_buf, const u64 off, const u32 pw_r_len, const u32 mask80, const u32 bits14, const u32 bits15, const u32 gid_max) { const u32 gid = get_global_id (0); @@ -104,7 +104,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) r_markov (__globa pws_buf_r[gid].i = pw_buf[0]; } -__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) C_markov (__global comb_t *pws_buf, __global cs_t *root_css_buf, __global cs_t *markov_css_buf, const u64 off, const u32 pw_len, const u32 mask80, const u32 bits14, const u32 bits15, const u32 gid_max) +__kernel void C_markov (__global comb_t *pws_buf, __global cs_t *root_css_buf, __global cs_t *markov_css_buf, const u64 off, const u32 pw_len, const u32 mask80, const u32 bits14, const u32 bits15, const u32 gid_max) { const u32 gid = get_global_id (0);