source: GPU/xmlwf.cl @ 1668

Last change on this file since 1668 was 1668, checked in by lindanl, 8 years ago

Parabix on GPU : start tag parsing

File size: 12.4 KB
Line 
1
2#define CARRY_BIT_MASK 0x8000000000000000
3#define WORK_GROUP_SIZE 64
4#define ELEMENTS 1024*128
5#define MAX_CARRY 2
6
7/******************Library Functions 64-bit********************/
8#define BitBlock unsigned long long
9#define simd_or(b1, b2) (b1 | b2)
10#define simd_and(b1, b2) (b1 & b2)
11#define simd_xor(b1, b2) (b1 ^ b2)
12#define simd_andc(b1, b2) (b1 & (~b2))
13#define simd_not(b) (~b)
14#define simd_nor(a,b) (simd_not(simd_or(a,b)))
15
16#define simd_himask_2 0xAAAAAAAAAAAAAAAA
17#define simd_himask_4 0xCCCCCCCCCCCCCCCC
18#define simd_himask_8 0xF0F0F0F0F0F0F0F0
19/* Little-endian */
20#define simd_himask_16 0xFF00FF00FF00FF00
21#define simd_himask_32 0xFFFF0000FFFF0000
22
23#define bitblock_has_bit(v) (v!=0)
24
25#define simd_add_32(a,b) (a+b)
26#define simd_sub_32(a,b) (a-b)
27
28#define sisd_slli(r, shft) (r<<shft)
29#define sisd_srli(r, shft) (r>>shft)
30
31#define sisd_from_int(n) n
32
33#define simd_const_1(x)\
34  ((x)==0 ? 0 : 0xFFFFFFFFFFFFFFFF)
35
36#define simd_if(cond, then_val, else_val) \
37  simd_or(simd_and(then_val, cond), simd_andc(else_val, cond))
38
39#define simd_pack_16_hh(b, a)\
40(a & 0xFF00000000000000)|((a & 0x0000FF0000000000)<<8)|((a & 0X00000000FF000000)<<16)|((a & 0X000000000000FF00)<<24)|\
41((b & 0xFF00000000000000)>>32)|((b & 0x0000FF0000000000)>>24)|((b & 0X00000000FF000000)>>16)|((b & 0X000000000000FF00)>>8)
42
43#define simd_pack_16_ll(b, a)\
44((a & 0x00FF000000000000)<<8)|((a & 0x000000FF00000000)<<16)|((a & 0X0000000000FF0000)<<24)|((a & 0X00000000000000FF)<<32)|\
45((b & 0x00FF000000000000)>>24)|((b & 0x000000FF00000000)>>16)|((b & 0X0000000000FF0000)>>8)|(b & 0X00000000000000FF)
46
47#define s2p_step(s0, s1, hi_mask, shift, p0, p1)  \
48  do {\
49        BitBlock t0,t1;\
50        t0 = simd_pack_16_hh(s0, s1);\
51        t1 = simd_pack_16_ll(s0, s1);\
52        p0 = simd_if(hi_mask, t0, (t1 >> shift));\
53        p1 = simd_if(hi_mask, (t0 << shift), t1);\
54  } while(0)
55
56#define s2p_bytepack(s0, s1, s2, s3, s4, s5, s6, s7, p0, p1, p2, p3, p4, p5, p6, p7) \
57  do {\
58        BitBlock bit00224466_0,bit00224466_1,bit00224466_2,bit00224466_3;\
59        BitBlock bit11335577_0,bit11335577_1,bit11335577_2,bit11335577_3;\
60        BitBlock bit00004444_0,bit22226666_0,bit00004444_1,bit22226666_1;\
61        BitBlock bit11115555_0,bit33337777_0,bit11115555_1,bit33337777_1;\
62        s2p_step(s0,s1,simd_himask_2,1,bit00224466_0,bit11335577_0);\
63        s2p_step(s2,s3,simd_himask_2,1,bit00224466_1,bit11335577_1);\
64        s2p_step(s4,s5,simd_himask_2,1,bit00224466_2,bit11335577_2);\
65        s2p_step(s6,s7,simd_himask_2,1,bit00224466_3,bit11335577_3);\
66        s2p_step(bit00224466_0,bit00224466_1,simd_himask_4,2,bit00004444_0,bit22226666_0);\
67        s2p_step(bit00224466_2,bit00224466_3,simd_himask_4,2,bit00004444_1,bit22226666_1);\
68        s2p_step(bit11335577_0,bit11335577_1,simd_himask_4,2,bit11115555_0,bit33337777_0);\
69        s2p_step(bit11335577_2,bit11335577_3,simd_himask_4,2,bit11115555_1,bit33337777_1);\
70        s2p_step(bit00004444_0,bit00004444_1,simd_himask_8,4,p0,p4);\
71        s2p_step(bit11115555_0,bit11115555_1,simd_himask_8,4,p1,p5);\
72        s2p_step(bit22226666_0,bit22226666_1,simd_himask_8,4,p2,p6);\
73        s2p_step(bit33337777_0,bit33337777_1,simd_himask_8,4,p3,p7);\
74  } while(0)
75
76
77#define adc(idx, a, b, c, carry, bubble, group_carry) \
78do {\
79\
80   BitBlock carry_mask;\
81   BitBlock bubble_mask;\
82\
83   BitBlock partial_sum = a+b;\
84   BitBlock gen = a&b;\
85   BitBlock prop = a^b;\
86   carry[idx] = ((gen | (prop & ~partial_sum))&CARRY_BIT_MASK)>>(WORK_GROUP_SIZE-1-idx);\
87   bubble[idx] = (partial_sum + 1)? 0:(((long long)1)<<idx);\
88 \
89   barrier(CLK_LOCAL_MEM_FENCE);\
90   for(int offset=WORK_GROUP_SIZE/2; offset>0; offset=offset>>1){\
91                carry[idx] = carry[idx]|carry[idx^offset];\
92                bubble[idx] = bubble[idx]|bubble[idx^offset];\
93                barrier(CLK_LOCAL_MEM_FENCE);\
94   }\
95\
96   carry_mask = (carry[0]<<1)|group_carry;\
97   bubble_mask = bubble[0];\
98\
99   BitBlock s = (carry_mask + bubble_mask) & ~bubble_mask;\
100   BitBlock inc = s | (s-carry_mask);\
101   c = partial_sum + ((inc>>idx)&0x1);\
102   group_carry = (carry[0]|(bubble_mask & inc))>>63;\
103\
104}while(0)
105
106#define sbb(idx, a, b, c, carry, bubble, group_carry) \
107do {\
108\
109   BitBlock carry_mask;\
110   BitBlock bubble_mask;\
111\
112   BitBlock partial_diff = a-b;\
113   BitBlock gen = b & ~a;\
114   BitBlock prop = ~(a^b);\
115   carry[idx] = ((gen | (prop & partial_diff))&CARRY_BIT_MASK)>>(63-idx);\
116   bubble[idx] = (partial_diff)? 0:(((long long)1)<<idx);\
117 \
118   barrier(CLK_LOCAL_MEM_FENCE);\
119   for(int offset=WORK_GROUP_SIZE/2; offset>0; offset=offset>>1){\
120                carry[idx] = carry[idx]|carry[idx^offset];\
121                bubble[idx] = bubble[idx]|bubble[idx^offset];\
122                barrier(CLK_LOCAL_MEM_FENCE);\
123   }\
124\
125   carry_mask = (carry[0]<<1)|group_carry;\
126   bubble_mask = bubble[0];\
127\
128   BitBlock s = (carry_mask + bubble_mask) & ~bubble_mask;\
129   BitBlock dec = s | (s-carry_mask);\
130   c = partial_diff - ((dec>>idx)&0x1);\
131   group_carry = (carry[0]|(bubble_mask & dec))>>63;\
132\
133}while(0)
134
135#define advance(idx, a, c, carry, group_carry) \
136do {\
137  carry[0] = group_carry;\
138  c = a<<1;\
139  carry[idx+1] = (a & CARRY_BIT_MASK)>>63;\
140  barrier(CLK_LOCAL_MEM_FENCE);\
141  group_carry = carry[WORK_GROUP_SIZE];\
142  c = c | carry[idx];\
143\
144}while(0)
145
146/**************Parabix Structs***************/
147  struct Basis_bits {
148  BitBlock bit_0;
149  BitBlock bit_1;
150  BitBlock bit_2;
151  BitBlock bit_3;
152  BitBlock bit_4;
153  BitBlock bit_5;
154  BitBlock bit_6;
155  BitBlock bit_7;
156};
157
158  struct Lex {
159  BitBlock CR;
160  BitBlock LF;
161  BitBlock HT;
162  BitBlock SP;
163  BitBlock CRLF;
164  BitBlock RefStart;
165  BitBlock Semicolon;
166  BitBlock Colon;
167  BitBlock LAngle;
168  BitBlock RAngle;
169  BitBlock LBracket;
170  BitBlock RBracket;
171  BitBlock Exclam;
172  BitBlock QMark;
173  BitBlock Hyphen;
174  BitBlock Equals;
175  BitBlock SQuote;
176  BitBlock DQuote;
177  BitBlock Slash;
178  BitBlock Hash;
179  BitBlock x;
180  BitBlock ASCII_name_start;
181  BitBlock ASCII_name_char;
182  BitBlock NameScan;
183  BitBlock Digit;
184  BitBlock Hex;
185  BitBlock WS;
186  BitBlock error;
187};
188
189/**************Parabix Functions***************/
190#define classify_bytes(idx, basis_bits, lex, carry)\
191 do {\
192                BitBlock temp1, temp2, temp3, temp4, temp5, temp6, temp7, temp8, temp9, temp10;\
193                BitBlock temp11, temp12, temp13, temp14, temp15, temp16, temp17, temp18;\
194                BitBlock temp19, temp20, temp21, temp22, temp23, temp24, temp25, temp26;\
195                BitBlock temp27, temp28, temp29, temp30, temp31, temp32, temp33, temp34;\
196                BitBlock temp35, temp36, temp37, temp38, temp39, temp40, temp41, temp42;\
197                BitBlock temp43, temp44, temp45, temp46, temp47, temp48, x00_x1F, temp49;\
198                BitBlock temp50, temp51, temp52, temp53, temp54, temp55, temp56, temp57;\
199                BitBlock temp58, temp59, temp60, temp61, temp62, temp63, temp64, temp65;\
200\
201        temp1 = simd_or(basis_bits.bit_0, basis_bits.bit_1);\
202        temp2 = simd_andc(basis_bits.bit_2, basis_bits.bit_3);\
203        temp3 = simd_andc(temp2, temp1);\
204        temp4 = simd_andc(basis_bits.bit_5, basis_bits.bit_4);\
205        temp5 = simd_andc(basis_bits.bit_6, basis_bits.bit_7);\
206        temp6 = simd_and(temp4, temp5);\
207        lex.RefStart = simd_and(temp3, temp6);\
208        temp7 = simd_and(basis_bits.bit_2, basis_bits.bit_3);\
209        temp8 = simd_andc(temp7, temp1);\
210        temp9 = simd_andc(basis_bits.bit_4, basis_bits.bit_5);\
211        temp10 = simd_and(basis_bits.bit_6, basis_bits.bit_7);\
212        temp11 = simd_and(temp9, temp10);\
213        lex.Semicolon = simd_and(temp8, temp11);\
214        temp12 = simd_and(basis_bits.bit_4, basis_bits.bit_5);\
215        temp13 = simd_or(basis_bits.bit_6, basis_bits.bit_7);\
216        temp14 = simd_andc(temp12, temp13);\
217        lex.LAngle = simd_and(temp8, temp14);\
218        temp15 = simd_and(temp12, temp5);\
219        lex.RAngle = simd_and(temp8, temp15);\
220        temp16 = simd_andc(basis_bits.bit_1, basis_bits.bit_0);\
221        temp17 = simd_andc(basis_bits.bit_3, basis_bits.bit_2);\
222        temp18 = simd_and(temp16, temp17);\
223        lex.LBracket = simd_and(temp18, temp11);\
224        temp19 = simd_andc(basis_bits.bit_7, basis_bits.bit_6);\
225        temp20 = simd_and(temp12, temp19);\
226        lex.RBracket = simd_and(temp18, temp20);\
227        temp21 = simd_or(basis_bits.bit_4, basis_bits.bit_5);\
228        temp22 = simd_andc(temp19, temp21);\
229        lex.Exclam = simd_and(temp3, temp22);\
230        temp23 = simd_and(temp12, temp10);\
231        lex.QMark = simd_and(temp8, temp23);\
232        lex.Hyphen = simd_and(temp3, temp20);\
233        lex.Equals = simd_and(temp8, temp20);\
234        temp24 = simd_and(temp4, temp10);\
235        lex.SQuote = simd_and(temp3, temp24);\
236        temp25 = simd_andc(temp5, temp21);\
237        lex.DQuote = simd_and(temp3, temp25);\
238        lex.Slash = simd_and(temp3, temp23);\
239        temp26 = simd_andc(temp10, temp21);\
240        lex.Hash = simd_and(temp3, temp26);\
241        temp27 = simd_and(temp16, temp7);\
242        temp28 = simd_andc(temp9, temp13);\
243        lex.x = simd_and(temp27, temp28);\
244        temp29 = simd_and(temp9, temp5);\
245        lex.Colon = simd_and(temp8, temp29);\
246        temp30 = simd_and(temp18, temp23);\
247        temp31 = simd_or(temp30, lex.Colon);\
248        temp32 = simd_andc(temp16, basis_bits.bit_2);\
249        temp33 = simd_or(basis_bits.bit_5, temp10);\
250        temp34 = simd_and(basis_bits.bit_4, temp33);\
251        temp35 = simd_not(temp34);\
252        temp36 = simd_or(temp21, temp13);\
253        temp37 = simd_or(simd_and(basis_bits.bit_3, temp35), simd_andc(temp36, basis_bits.bit_3));\
254        temp38 = simd_and(temp32, temp37);\
255        temp39 = simd_or(temp31, temp38);\
256        temp40 = simd_and(temp16, basis_bits.bit_2);\
257        temp41 = simd_and(temp40, temp37);\
258        lex.ASCII_name_start = simd_or(temp39, temp41);\
259        temp42 = simd_or(temp30, lex.Hyphen);\
260        temp43 = simd_and(temp3, temp15);\
261        temp44 = simd_or(temp42, temp43);\
262        temp45 = simd_andc(temp8, temp34);\
263        temp46 = simd_or(temp44, temp45);\
264        temp47 = simd_or(temp46, temp38);\
265        lex.ASCII_name_char = simd_or(temp47, temp41);\
266        lex.NameScan = simd_or(lex.ASCII_name_char, basis_bits.bit_0);\
267        temp48 = simd_or(temp1, basis_bits.bit_2);\
268        x00_x1F = simd_not(temp48);\
269        temp49 = simd_or(basis_bits.bit_2, basis_bits.bit_3);\
270        temp50 = simd_or(temp1, temp49);\
271        lex.CR = simd_andc(temp20, temp50);\
272        lex.LF = simd_andc(temp29, temp50);\
273        temp51 = simd_and(temp9, temp19);\
274        lex.HT = simd_andc(temp51, temp50);\
275        lex.SP = simd_andc(temp3, temp36);\
276        temp52 = simd_or(temp20, temp29);\
277        temp53 = simd_or(temp52, temp51);\
278        temp54 = simd_andc(temp53, temp50);\
279        lex.WS = simd_or(temp54, lex.SP);\
280        temp55 = simd_or(basis_bits.bit_5, basis_bits.bit_6);\
281        temp56 = simd_and(basis_bits.bit_4, temp55);\
282        lex.Digit = simd_andc(temp8, temp56);\
283        temp57 = simd_andc(temp16, temp49);\
284        temp58 = simd_andc(temp57, basis_bits.bit_4);\
285        temp59 = simd_not(temp10);\
286        temp60 = simd_or(simd_and(basis_bits.bit_5, temp59), simd_andc(temp13, basis_bits.bit_5));\
287        temp61 = simd_and(temp58, temp60);\
288        temp62 = simd_or(lex.Digit, temp61);\
289        temp63 = simd_and(temp16, temp2);\
290        temp64 = simd_andc(temp63, basis_bits.bit_4);\
291        temp65 = simd_and(temp64, temp60);\
292        lex.Hex = simd_or(temp62, temp65);\
293        lex.error = simd_andc(x00_x1F, lex.WS);\
294  } while (0)
295
296/*************Kernel Function****************/
297__kernel
298void xmlwf(__global BitBlock *s, 
299                   __global BitBlock *err_pos,
300                        __local BitBlock *carry,
301                        __local BitBlock *bubble) {
302
303   int i = get_local_id(0);
304   BitBlock group_carryQ[MAX_CARRY];
305   struct Basis_bits basis_bits;
306   struct Lex lex;
307   struct U8 u8;
308   struct Scope1 scope1;
309   struct CtCDPI_Callouts ctCDPI_Callouts;
310
311   for(int i=0; i<MAX_CARRY; i++)
312                group_carryQ[i] = 0;
313
314   for(int idx=i; idx<ELEMENTS;idx=idx+WORK_GROUP_SIZE){
315                s2p_bytepack(s[idx*8], s[idx*8+1], s[idx*8+2], s[idx*8+3], s[idx*8+4], s[idx*8+5], s[idx*8+6], s[idx*8+7], 
316                basis_bits.bit_0, basis_bits.bit_1, basis_bits.bit_2, basis_bits.bit_3, 
317                basis_bits.bit_4, basis_bits.bit_5, basis_bits.bit_6, basis_bits.bit_7);
318               
319                classify_bytes(i, basis_bits, lex, carry);
320                BitBlock LA_scope,StartTag_NameStart,StartTag_End;
321                advance(i,lex.LAngle, LA_scope, carry, group_carryQ[0]);
322                StartTag_NameStart = simd_andc(LA_scope,lex.Slash);
323                adc(i,StartTag_NameStart,lex.NameScan,StartTag_End,carry,bubble, group_carryQ[1]);
324                StartTag_End = simd_andc(StartTag_End, lex.NameScan);
325                err_pos[idx] = simd_andc(StartTag_End,lex.RAngle);
326
327        }
328
329}
Note: See TracBrowser for help on using the repository browser.