2 * Author......: Jens Steube <jens.steube@gmail.com>
6 #include "include/kernel_vendor.h"
12 #include "types_amd.c"
14 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)
33 __global cs_t *cs = &root_css_buf[pw_r_len];
38 for (i = 0, j = pw_r_len; i < pw_l_len; i++, j++)
40 const u32 len = cs->cs_len;
42 const u64 next = val / len;
43 const u64 pos = val % len;
47 const u32 key = cs->cs_buf[pos];
49 const u32 jd4 = j / 4;
50 const u32 jm4 = j % 4;
52 pw_buf[jd4] |= key << ((3 - jm4) * 8);
54 cs = &markov_css_buf[(j * CHARSIZ) + key];
57 const u32 jd4 = j / 4;
58 const u32 jm4 = j % 4;
60 pw_buf[jd4] |= (0xff << ((3 - jm4) * 8)) & mask80;
62 if (bits14) pw_buf[14] = (pw_l_len + pw_r_len) * 8;
63 if (bits15) pw_buf[15] = (pw_l_len + pw_r_len) * 8;
66 __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)
68 const u32 gid = get_global_id (0);
70 if (gid >= gid_max) return;
72 const u32 gid2 = gid * 2;
77 generate_pw (pw_buf0, root_css_buf, markov_css_buf, pw_l_len, pw_r_len, mask80, bits14, bits15, off + gid2 + 0);
78 generate_pw (pw_buf1, root_css_buf, markov_css_buf, pw_l_len, pw_r_len, mask80, bits14, bits15, off + gid2 + 1);
81 for (int i = 0; i < 16; i++)
83 pws_buf_l[gid].i[i].s0 = pw_buf0[i];
84 pws_buf_l[gid].i[i].s1 = pw_buf1[i];
87 pws_buf_l[gid].pw_len = pw_l_len + pw_r_len;
90 __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)
92 const u32 gid = get_global_id (0);
94 if (gid >= gid_max) return;
96 const u32 gid2 = gid * 2;
100 generate_pw (pw_buf, root_css_buf, markov_css_buf, pw_r_len, 0, 0, 0, 0, off + gid2 + 0);
102 pws_buf_r[gid2 + 0].i = pw_buf[0];
104 generate_pw (pw_buf, root_css_buf, markov_css_buf, pw_r_len, 0, 0, 0, 0, off + gid2 + 1);
106 pws_buf_r[gid2 + 1].i = pw_buf[0];
108 generate_pw (pw_buf, root_css_buf, markov_css_buf, pw_r_len, 0, 0, 0, 0, off + gid2 + 2);
110 pws_buf_r[gid2 + 2].i = pw_buf[0];
112 generate_pw (pw_buf, root_css_buf, markov_css_buf, pw_r_len, 0, 0, 0, 0, off + gid2 + 3);
114 pws_buf_r[gid2 + 3].i = pw_buf[0];
117 __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)
119 const u32 gid = get_global_id (0);
121 if (gid >= gid_max) return;
123 const u32 gid2 = gid * 2;
127 generate_pw (pw_buf, root_css_buf, markov_css_buf, pw_len, 0, mask80, bits14, bits15, off + gid2 + 0);
129 pws_buf[gid2 + 0].i[ 0] = pw_buf[ 0];
130 pws_buf[gid2 + 0].i[ 1] = pw_buf[ 1];
131 pws_buf[gid2 + 0].i[ 2] = pw_buf[ 2];
132 pws_buf[gid2 + 0].i[ 3] = pw_buf[ 3];
133 pws_buf[gid2 + 0].i[ 4] = pw_buf[ 4];
134 pws_buf[gid2 + 0].i[ 5] = pw_buf[ 5];
135 pws_buf[gid2 + 0].i[ 6] = pw_buf[ 6];
136 pws_buf[gid2 + 0].i[ 7] = pw_buf[ 7];
138 pws_buf[gid2 + 0].pw_len = pw_len;
140 generate_pw (pw_buf, root_css_buf, markov_css_buf, pw_len, 0, mask80, bits14, bits15, off + gid2 + 1);
142 pws_buf[gid2 + 1].i[ 0] = pw_buf[ 0];
143 pws_buf[gid2 + 1].i[ 1] = pw_buf[ 1];
144 pws_buf[gid2 + 1].i[ 2] = pw_buf[ 2];
145 pws_buf[gid2 + 1].i[ 3] = pw_buf[ 3];
146 pws_buf[gid2 + 1].i[ 4] = pw_buf[ 4];
147 pws_buf[gid2 + 1].i[ 5] = pw_buf[ 5];
148 pws_buf[gid2 + 1].i[ 6] = pw_buf[ 6];
149 pws_buf[gid2 + 1].i[ 7] = pw_buf[ 7];
151 pws_buf[gid2 + 1].pw_len = pw_len;
153 generate_pw (pw_buf, root_css_buf, markov_css_buf, pw_len, 0, mask80, bits14, bits15, off + gid2 + 2);
155 pws_buf[gid2 + 2].i[ 0] = pw_buf[ 0];
156 pws_buf[gid2 + 2].i[ 1] = pw_buf[ 1];
157 pws_buf[gid2 + 2].i[ 2] = pw_buf[ 2];
158 pws_buf[gid2 + 2].i[ 3] = pw_buf[ 3];
159 pws_buf[gid2 + 2].i[ 4] = pw_buf[ 4];
160 pws_buf[gid2 + 2].i[ 5] = pw_buf[ 5];
161 pws_buf[gid2 + 2].i[ 6] = pw_buf[ 6];
162 pws_buf[gid2 + 2].i[ 7] = pw_buf[ 7];
164 pws_buf[gid2 + 2].pw_len = pw_len;
166 generate_pw (pw_buf, root_css_buf, markov_css_buf, pw_len, 0, mask80, bits14, bits15, off + gid2 + 3);
168 pws_buf[gid2 + 3].i[ 0] = pw_buf[ 0];
169 pws_buf[gid2 + 3].i[ 1] = pw_buf[ 1];
170 pws_buf[gid2 + 3].i[ 2] = pw_buf[ 2];
171 pws_buf[gid2 + 3].i[ 3] = pw_buf[ 3];
172 pws_buf[gid2 + 3].i[ 4] = pw_buf[ 4];
173 pws_buf[gid2 + 3].i[ 5] = pw_buf[ 5];
174 pws_buf[gid2 + 3].i[ 6] = pw_buf[ 6];
175 pws_buf[gid2 + 3].i[ 7] = pw_buf[ 7];
177 pws_buf[gid2 + 3].pw_len = pw_len;