source: GPU/pipeline_xmlwf.cl @ 4532

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

Pipelined Parabix on GPU

File size: 14.5 KB
RevLine 
[1759]1#pragma OPENCL EXTENSION cl_amd_printf:enable
2#define CARRY_BIT_MASK 0x8000000000000000
3#define WORK_GROUP_SIZE 64
4#define ELEMENTS 1024*128
5#define MAX_CARRY 2
6#define ENTRIES 8
7
8/******************Library Functions 64-bit********************/
9#define BitBlock unsigned long long
10#define simd_or(b1, b2) (b1 | b2)
11#define simd_and(b1, b2) (b1 & b2)
12#define simd_xor(b1, b2) (b1 ^ b2)
13#define simd_andc(b1, b2) (b1 & (~b2))
14#define simd_not(b) (~b)
15#define simd_nor(a,b) (simd_not(simd_or(a,b)))
16
17#define simd_himask_2 0xAAAAAAAAAAAAAAAA
18#define simd_himask_4 0xCCCCCCCCCCCCCCCC
19#define simd_himask_8 0xF0F0F0F0F0F0F0F0
20/* Little-endian */
21#define simd_himask_16 0xFF00FF00FF00FF00
22#define simd_himask_32 0xFFFF0000FFFF0000
23
24#define bitblock_has_bit(v) (v!=0)
25
26#define simd_add_32(a,b) (a+b)
27#define simd_sub_32(a,b) (a-b)
28
29#define sisd_slli(r, shft) (r<<shft)
30#define sisd_srli(r, shft) (r>>shft)
31
32#define sisd_from_int(n) n
33
34#define simd_const_1(x)\
35  ((x)==0 ? 0 : 0xFFFFFFFFFFFFFFFF)
36
37#define simd_if(cond, then_val, else_val) \
38  simd_or(simd_and(then_val, cond), simd_andc(else_val, cond))
39
40#define simd_pack_16_hh(b, a)\
41(a & 0xFF00000000000000)|((a & 0x0000FF0000000000)<<8)|((a & 0X00000000FF000000)<<16)|((a & 0X000000000000FF00)<<24)|\
42((b & 0xFF00000000000000)>>32)|((b & 0x0000FF0000000000)>>24)|((b & 0X00000000FF000000)>>16)|((b & 0X000000000000FF00)>>8)
43
44#define simd_pack_16_ll(b, a)\
45((a & 0x00FF000000000000)<<8)|((a & 0x000000FF00000000)<<16)|((a & 0X0000000000FF0000)<<24)|((a & 0X00000000000000FF)<<32)|\
46((b & 0x00FF000000000000)>>24)|((b & 0x000000FF00000000)>>16)|((b & 0X0000000000FF0000)>>8)|(b & 0X00000000000000FF)
47
48#define s2p_step(s0, s1, hi_mask, shift, p0, p1)  \
49  do {\
50        BitBlock t0,t1;\
51        t0 = simd_pack_16_hh(s0, s1);\
52        t1 = simd_pack_16_ll(s0, s1);\
53        p0 = simd_if(hi_mask, t0, (t1 >> shift));\
54        p1 = simd_if(hi_mask, (t0 << shift), t1);\
55  } while(0)
56
57#define s2p_bytepack(s0, s1, s2, s3, s4, s5, s6, s7, p0, p1, p2, p3, p4, p5, p6, p7) \
58  do {\
59        BitBlock bit00224466_0,bit00224466_1,bit00224466_2,bit00224466_3;\
60        BitBlock bit11335577_0,bit11335577_1,bit11335577_2,bit11335577_3;\
61        BitBlock bit00004444_0,bit22226666_0,bit00004444_1,bit22226666_1;\
62        BitBlock bit11115555_0,bit33337777_0,bit11115555_1,bit33337777_1;\
63        s2p_step(s0,s1,simd_himask_2,1,bit00224466_0,bit11335577_0);\
64        s2p_step(s2,s3,simd_himask_2,1,bit00224466_1,bit11335577_1);\
65        s2p_step(s4,s5,simd_himask_2,1,bit00224466_2,bit11335577_2);\
66        s2p_step(s6,s7,simd_himask_2,1,bit00224466_3,bit11335577_3);\
67        s2p_step(bit00224466_0,bit00224466_1,simd_himask_4,2,bit00004444_0,bit22226666_0);\
68        s2p_step(bit00224466_2,bit00224466_3,simd_himask_4,2,bit00004444_1,bit22226666_1);\
69        s2p_step(bit11335577_0,bit11335577_1,simd_himask_4,2,bit11115555_0,bit33337777_0);\
70        s2p_step(bit11335577_2,bit11335577_3,simd_himask_4,2,bit11115555_1,bit33337777_1);\
71        s2p_step(bit00004444_0,bit00004444_1,simd_himask_8,4,p0,p4);\
72        s2p_step(bit11115555_0,bit11115555_1,simd_himask_8,4,p1,p5);\
73        s2p_step(bit22226666_0,bit22226666_1,simd_himask_8,4,p2,p6);\
74        s2p_step(bit33337777_0,bit33337777_1,simd_himask_8,4,p3,p7);\
75  } while(0)
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
189struct Pdata{
190          int cur_stage[WORK_GROUP_SIZE];
191          struct Basis_bits basis_bits[WORK_GROUP_SIZE];
192          struct Lex lex[WORK_GROUP_SIZE];
193};
194
195/**************Parabix Functions***************/
196#define classify_bytes(idx, basis_bits, lex, carry)\
197 do {\
198                BitBlock temp1, temp2, temp3, temp4, temp5, temp6, temp7, temp8, temp9, temp10;\
199                BitBlock temp11, temp12, temp13, temp14, temp15, temp16, temp17, temp18;\
200                BitBlock temp19, temp20, temp21, temp22, temp23, temp24, temp25, temp26;\
201                BitBlock temp27, temp28, temp29, temp30, temp31, temp32, temp33, temp34;\
202                BitBlock temp35, temp36, temp37, temp38, temp39, temp40, temp41, temp42;\
203                BitBlock temp43, temp44, temp45, temp46, temp47, temp48, x00_x1F, temp49;\
204                BitBlock temp50, temp51, temp52, temp53, temp54, temp55, temp56, temp57;\
205                BitBlock temp58, temp59, temp60, temp61, temp62, temp63, temp64, temp65;\
206\
207        temp1 = simd_or(basis_bits.bit_0, basis_bits.bit_1);\
208        temp2 = simd_andc(basis_bits.bit_2, basis_bits.bit_3);\
209        temp3 = simd_andc(temp2, temp1);\
210        temp4 = simd_andc(basis_bits.bit_5, basis_bits.bit_4);\
211        temp5 = simd_andc(basis_bits.bit_6, basis_bits.bit_7);\
212        temp6 = simd_and(temp4, temp5);\
213        lex.RefStart = simd_and(temp3, temp6);\
214        temp7 = simd_and(basis_bits.bit_2, basis_bits.bit_3);\
215        temp8 = simd_andc(temp7, temp1);\
216        temp9 = simd_andc(basis_bits.bit_4, basis_bits.bit_5);\
217        temp10 = simd_and(basis_bits.bit_6, basis_bits.bit_7);\
218        temp11 = simd_and(temp9, temp10);\
219        lex.Semicolon = simd_and(temp8, temp11);\
220        temp12 = simd_and(basis_bits.bit_4, basis_bits.bit_5);\
221        temp13 = simd_or(basis_bits.bit_6, basis_bits.bit_7);\
222        temp14 = simd_andc(temp12, temp13);\
223        lex.LAngle = simd_and(temp8, temp14);\
224        temp15 = simd_and(temp12, temp5);\
225        lex.RAngle = simd_and(temp8, temp15);\
226        temp16 = simd_andc(basis_bits.bit_1, basis_bits.bit_0);\
227        temp17 = simd_andc(basis_bits.bit_3, basis_bits.bit_2);\
228        temp18 = simd_and(temp16, temp17);\
229        lex.LBracket = simd_and(temp18, temp11);\
230        temp19 = simd_andc(basis_bits.bit_7, basis_bits.bit_6);\
231        temp20 = simd_and(temp12, temp19);\
232        lex.RBracket = simd_and(temp18, temp20);\
233        temp21 = simd_or(basis_bits.bit_4, basis_bits.bit_5);\
234        temp22 = simd_andc(temp19, temp21);\
235        lex.Exclam = simd_and(temp3, temp22);\
236        temp23 = simd_and(temp12, temp10);\
237        lex.QMark = simd_and(temp8, temp23);\
238        lex.Hyphen = simd_and(temp3, temp20);\
239        lex.Equals = simd_and(temp8, temp20);\
240        temp24 = simd_and(temp4, temp10);\
241        lex.SQuote = simd_and(temp3, temp24);\
242        temp25 = simd_andc(temp5, temp21);\
243        lex.DQuote = simd_and(temp3, temp25);\
244        lex.Slash = simd_and(temp3, temp23);\
245        temp26 = simd_andc(temp10, temp21);\
246        lex.Hash = simd_and(temp3, temp26);\
247        temp27 = simd_and(temp16, temp7);\
248        temp28 = simd_andc(temp9, temp13);\
249        lex.x = simd_and(temp27, temp28);\
250        temp29 = simd_and(temp9, temp5);\
251        lex.Colon = simd_and(temp8, temp29);\
252        temp30 = simd_and(temp18, temp23);\
253        temp31 = simd_or(temp30, lex.Colon);\
254        temp32 = simd_andc(temp16, basis_bits.bit_2);\
255        temp33 = simd_or(basis_bits.bit_5, temp10);\
256        temp34 = simd_and(basis_bits.bit_4, temp33);\
257        temp35 = simd_not(temp34);\
258        temp36 = simd_or(temp21, temp13);\
259        temp37 = simd_or(simd_and(basis_bits.bit_3, temp35), simd_andc(temp36, basis_bits.bit_3));\
260        temp38 = simd_and(temp32, temp37);\
261        temp39 = simd_or(temp31, temp38);\
262        temp40 = simd_and(temp16, basis_bits.bit_2);\
263        temp41 = simd_and(temp40, temp37);\
264        lex.ASCII_name_start = simd_or(temp39, temp41);\
265        temp42 = simd_or(temp30, lex.Hyphen);\
266        temp43 = simd_and(temp3, temp15);\
267        temp44 = simd_or(temp42, temp43);\
268        temp45 = simd_andc(temp8, temp34);\
269        temp46 = simd_or(temp44, temp45);\
270        temp47 = simd_or(temp46, temp38);\
271        lex.ASCII_name_char = simd_or(temp47, temp41);\
272        lex.NameScan = simd_or(lex.ASCII_name_char, basis_bits.bit_0);\
273        temp48 = simd_or(temp1, basis_bits.bit_2);\
274        x00_x1F = simd_not(temp48);\
275        temp49 = simd_or(basis_bits.bit_2, basis_bits.bit_3);\
276        temp50 = simd_or(temp1, temp49);\
277        lex.CR = simd_andc(temp20, temp50);\
278        lex.LF = simd_andc(temp29, temp50);\
279        temp51 = simd_and(temp9, temp19);\
280        lex.HT = simd_andc(temp51, temp50);\
281        lex.SP = simd_andc(temp3, temp36);\
282        temp52 = simd_or(temp20, temp29);\
283        temp53 = simd_or(temp52, temp51);\
284        temp54 = simd_andc(temp53, temp50);\
285        lex.WS = simd_or(temp54, lex.SP);\
286        temp55 = simd_or(basis_bits.bit_5, basis_bits.bit_6);\
287        temp56 = simd_and(basis_bits.bit_4, temp55);\
288        lex.Digit = simd_andc(temp8, temp56);\
289        temp57 = simd_andc(temp16, temp49);\
290        temp58 = simd_andc(temp57, basis_bits.bit_4);\
291        temp59 = simd_not(temp10);\
292        temp60 = simd_or(simd_and(basis_bits.bit_5, temp59), simd_andc(temp13, basis_bits.bit_5));\
293        temp61 = simd_and(temp58, temp60);\
294        temp62 = simd_or(lex.Digit, temp61);\
295        temp63 = simd_and(temp16, temp2);\
296        temp64 = simd_andc(temp63, basis_bits.bit_4);\
297        temp65 = simd_and(temp64, temp60);\
298        lex.Hex = simd_or(temp62, temp65);\
299        lex.error = simd_andc(x00_x1F, lex.WS);\
300  } while (0)
301
302/*************Kernel Function****************/
303
304__kernel
305void tag_parsing(__global BitBlock *s, 
306                   __global BitBlock *err_pos,
307                        __local BitBlock *carry,
308                        __local BitBlock *bubble,
309                        __global struct Pdata *pdata) {
310
311   
312        int gid = get_global_id(0);
313        int i = gid % WORK_GROUP_SIZE;
314        int stage = gid / WORK_GROUP_SIZE;
315   
316        BitBlock group_carryQ[MAX_CARRY];
317        int entry_no = 0;
318        switch(stage){
319                case 0:
320                        for(int idx=i; idx<ELEMENTS;idx=idx+WORK_GROUP_SIZE, entry_no=(entry_no+1)%ENTRIES){
321                                while(pdata[entry_no].cur_stage[i]!=0)
322                                        barrier(CLK_LOCAL_MEM_FENCE);
323                                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], 
324                                        pdata[entry_no].basis_bits[i].bit_0, pdata[entry_no].basis_bits[i].bit_1, pdata[entry_no].basis_bits[i].bit_2, pdata[entry_no].basis_bits[i].bit_3, 
325                                        pdata[entry_no].basis_bits[i].bit_4, pdata[entry_no].basis_bits[i].bit_5, pdata[entry_no].basis_bits[i].bit_6, pdata[entry_no].basis_bits[i].bit_7);
326                                pdata[entry_no].cur_stage[i]=1;
327                        }
328                break;
329                case 1:
330                        for(int idx=i; idx<ELEMENTS;idx=idx+WORK_GROUP_SIZE, entry_no=(entry_no+1)%ENTRIES){
331                                while(pdata[entry_no].cur_stage[i]!=1)
332                                        barrier(CLK_LOCAL_MEM_FENCE);
333                                classify_bytes(i, pdata[entry_no].basis_bits[i], pdata[entry_no].lex[i], carry);
334                                pdata[entry_no].cur_stage[i]=2;
335                        }
336                break;
337                case 2:         
338                        for(int i=0; i<MAX_CARRY; i++)
339                                group_carryQ[i] = 0;
340                        for(int idx=i; idx<ELEMENTS;idx=idx+WORK_GROUP_SIZE, entry_no=(entry_no+1)%ENTRIES){
341                                while(pdata[entry_no].cur_stage[i]!=2)
342                                        barrier(CLK_LOCAL_MEM_FENCE);
343
344                                BitBlock LA_scope,StartTag_NameStart,StartTag_End;
345                                advance(i,pdata[entry_no].lex[i].LAngle, LA_scope, carry, group_carryQ[0]);
346                                StartTag_NameStart = simd_andc(LA_scope,pdata[entry_no].lex[i].Slash);
347                                adc(i,StartTag_NameStart,pdata[entry_no].lex[i].NameScan,StartTag_End,carry,bubble, group_carryQ[1]);
348                                StartTag_End = simd_andc(StartTag_End, pdata[entry_no].lex[i].NameScan);
349                                err_pos[idx] = simd_andc(StartTag_End,pdata[entry_no].lex[i].RAngle);
350
351                                pdata[entry_no].cur_stage[i]=0;
352                        }
353                break;
354   }
355
356/*     
357        for(int i=0; i<MAX_CARRY; i++)
358                group_carryQ[i] = 0;
359
360        for(int idx=i; idx<ELEMENTS;idx=idx+WORK_GROUP_SIZE){
361                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], 
362                pdata[0].basis_bits[i].bit_0, pdata[0].basis_bits[i].bit_1, pdata[0].basis_bits[i].bit_2, pdata[0].basis_bits[i].bit_3, 
363                pdata[0].basis_bits[i].bit_4, pdata[0].basis_bits[i].bit_5, pdata[0].basis_bits[i].bit_6, pdata[0].basis_bits[i].bit_7);
364               
365                classify_bytes(i, pdata[0].basis_bits[i], pdata[0].lex[i], carry);
366                BitBlock LA_scope,StartTag_NameStart,StartTag_End;
367                advance(i,pdata[0].lex[i].LAngle, LA_scope, carry, group_carryQ[0]);
368                StartTag_NameStart = simd_andc(LA_scope,pdata[0].lex[i].Slash);
369                adc(i,StartTag_NameStart,pdata[0].lex[i].NameScan,StartTag_End,carry,bubble, group_carryQ[1]);
370                StartTag_End = simd_andc(StartTag_End, pdata[0].lex[i].NameScan);
371                err_pos[idx] = simd_andc(StartTag_End,pdata[0].lex[i].RAngle);
372
373        }
374*/     
375}
Note: See TracBrowser for help on using the repository browser.