|
|
|
@ -9,25 +9,8 @@
|
|
|
|
|
|
|
|
|
|
#include "inc_types.cl"
|
|
|
|
|
|
|
|
|
|
inline void generate_pw (u32 pw_buf[16], __global const cs_t *root_css_buf, __global const 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[64], __global const cs_t *root_css_buf, __global const 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;
|
|
|
|
|
pw_buf[ 2] = 0;
|
|
|
|
|
pw_buf[ 3] = 0;
|
|
|
|
|
pw_buf[ 4] = 0;
|
|
|
|
|
pw_buf[ 5] = 0;
|
|
|
|
|
pw_buf[ 6] = 0;
|
|
|
|
|
pw_buf[ 7] = 0;
|
|
|
|
|
pw_buf[ 8] = 0;
|
|
|
|
|
pw_buf[ 9] = 0;
|
|
|
|
|
pw_buf[10] = 0;
|
|
|
|
|
pw_buf[11] = 0;
|
|
|
|
|
pw_buf[12] = 0;
|
|
|
|
|
pw_buf[13] = 0;
|
|
|
|
|
pw_buf[14] = 0;
|
|
|
|
|
pw_buf[15] = 0;
|
|
|
|
|
|
|
|
|
|
__global const cs_t *cs = &root_css_buf[pw_r_len];
|
|
|
|
|
|
|
|
|
|
u32 i;
|
|
|
|
@ -67,26 +50,15 @@ __kernel void l_markov (__global pw_t *pws_buf_l, __global const cs_t *root_css_
|
|
|
|
|
|
|
|
|
|
if (gid >= gid_max) return;
|
|
|
|
|
|
|
|
|
|
u32 pw_buf[16];
|
|
|
|
|
u32 pw_buf[64] = { 0 };
|
|
|
|
|
|
|
|
|
|
generate_pw (pw_buf, root_css_buf, markov_css_buf, pw_l_len, pw_r_len, mask80, bits14, bits15, off + gid);
|
|
|
|
|
|
|
|
|
|
pws_buf_l[gid].i[ 0] = pw_buf[ 0];
|
|
|
|
|
pws_buf_l[gid].i[ 1] = pw_buf[ 1];
|
|
|
|
|
pws_buf_l[gid].i[ 2] = pw_buf[ 2];
|
|
|
|
|
pws_buf_l[gid].i[ 3] = pw_buf[ 3];
|
|
|
|
|
pws_buf_l[gid].i[ 4] = pw_buf[ 4];
|
|
|
|
|
pws_buf_l[gid].i[ 5] = pw_buf[ 5];
|
|
|
|
|
pws_buf_l[gid].i[ 6] = pw_buf[ 6];
|
|
|
|
|
pws_buf_l[gid].i[ 7] = pw_buf[ 7];
|
|
|
|
|
pws_buf_l[gid].i[ 8] = pw_buf[ 8];
|
|
|
|
|
pws_buf_l[gid].i[ 9] = pw_buf[ 9];
|
|
|
|
|
pws_buf_l[gid].i[10] = pw_buf[10];
|
|
|
|
|
pws_buf_l[gid].i[11] = pw_buf[11];
|
|
|
|
|
pws_buf_l[gid].i[12] = pw_buf[12];
|
|
|
|
|
pws_buf_l[gid].i[13] = pw_buf[13];
|
|
|
|
|
pws_buf_l[gid].i[14] = pw_buf[14];
|
|
|
|
|
pws_buf_l[gid].i[15] = pw_buf[15];
|
|
|
|
|
#pragma unroll
|
|
|
|
|
for (int idx = 0; idx < 64; idx++)
|
|
|
|
|
{
|
|
|
|
|
pws_buf_l[gid].i[idx] = pw_buf[idx];
|
|
|
|
|
}
|
|
|
|
|
|
|
|
|
|
pws_buf_l[gid].pw_len = pw_l_len + pw_r_len;
|
|
|
|
|
}
|
|
|
|
@ -97,7 +69,7 @@ __kernel void r_markov (__global bf_t *pws_buf_r, __global const cs_t *root_css_
|
|
|
|
|
|
|
|
|
|
if (gid >= gid_max) return;
|
|
|
|
|
|
|
|
|
|
u32 pw_buf[16];
|
|
|
|
|
u32 pw_buf[64] = { 0 };
|
|
|
|
|
|
|
|
|
|
generate_pw (pw_buf, root_css_buf, markov_css_buf, pw_r_len, 0, 0, 0, 0, off + gid);
|
|
|
|
|
|
|
|
|
@ -110,18 +82,15 @@ __kernel void C_markov (__global comb_t *pws_buf, __global const cs_t *root_css_
|
|
|
|
|
|
|
|
|
|
if (gid >= gid_max) return;
|
|
|
|
|
|
|
|
|
|
u32 pw_buf[16];
|
|
|
|
|
u32 pw_buf[64] = { 0 };
|
|
|
|
|
|
|
|
|
|
generate_pw (pw_buf, root_css_buf, markov_css_buf, pw_len, 0, mask80, bits14, bits15, off + gid);
|
|
|
|
|
|
|
|
|
|
pws_buf[gid].i[ 0] = pw_buf[ 0];
|
|
|
|
|
pws_buf[gid].i[ 1] = pw_buf[ 1];
|
|
|
|
|
pws_buf[gid].i[ 2] = pw_buf[ 2];
|
|
|
|
|
pws_buf[gid].i[ 3] = pw_buf[ 3];
|
|
|
|
|
pws_buf[gid].i[ 4] = pw_buf[ 4];
|
|
|
|
|
pws_buf[gid].i[ 5] = pw_buf[ 5];
|
|
|
|
|
pws_buf[gid].i[ 6] = pw_buf[ 6];
|
|
|
|
|
pws_buf[gid].i[ 7] = pw_buf[ 7];
|
|
|
|
|
#pragma unroll
|
|
|
|
|
for (int idx = 0; idx < 64; idx++)
|
|
|
|
|
{
|
|
|
|
|
pws_buf[gid].i[idx] = pw_buf[idx];
|
|
|
|
|
}
|
|
|
|
|
|
|
|
|
|
pws_buf[gid].pw_len = pw_len;
|
|
|
|
|
}
|
|
|
|
|