source: GPU/xmlwf.cl @ 2205

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

Parabix on GPU : Add ref, tag, name and check functions.

File size: 38.9 KB
RevLine 
[1668]1
2#define CARRY_BIT_MASK 0x8000000000000000
3#define WORK_GROUP_SIZE 64
4#define ELEMENTS 1024*128
[1669]5#define MAX_CARRY 10
6#define EOF_mask 0xFFFFFFFFFFFFFFFF
[1668]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
[1669]78inline BitBlock adc(int idx, BitBlock a, BitBlock b, __local BitBlock *carry, __local BitBlock *bubble, BitBlock *group_carry, const int carryno){
79        BitBlock carry_mask;
80        BitBlock bubble_mask;
[1668]81
[1669]82        BitBlock partial_sum = a+b;
83        BitBlock gen = a&b;
84        BitBlock prop = a^b;
85        carry[idx] = ((gen | (prop & ~partial_sum))&CARRY_BIT_MASK)>>(WORK_GROUP_SIZE-1-idx);
86        bubble[idx] = (partial_sum + 1)? 0:(((long long)1)<<idx);
87       
88        barrier(CLK_LOCAL_MEM_FENCE);
89        for(int offset=WORK_GROUP_SIZE/2; offset>0; offset=offset>>1){
90                carry[idx] = carry[idx]|carry[idx^offset];
91                bubble[idx] = bubble[idx]|bubble[idx^offset];
92                barrier(CLK_LOCAL_MEM_FENCE);
93        }
94       
95        carry_mask = (carry[0]<<1)|group_carry[carryno];
96        bubble_mask = bubble[0];
97       
98        BitBlock s = (carry_mask + bubble_mask) & ~bubble_mask;
99        BitBlock inc = s | (s-carry_mask);
100        BitBlock rslt = partial_sum + ((inc>>idx)&0x1);
101        group_carry[carryno] = (carry[0]|(bubble_mask & inc))>>63;
102        return rslt;
103}
[1668]104
[1669]105inline BitBlock sbb(int idx, BitBlock a, BitBlock b, __local BitBlock *carry, __local BitBlock *bubble, BitBlock *group_carry, const int carryno){
106        BitBlock carry_mask;
107        BitBlock bubble_mask;
[1668]108
[1669]109        BitBlock partial_diff = a-b;
110        BitBlock gen = b & ~a;
111        BitBlock prop = ~(a^b);
112        carry[idx] = ((gen | (prop & partial_diff))&CARRY_BIT_MASK)>>(63-idx);
113        bubble[idx] = (partial_diff)? 0:(((long long)1)<<idx);
114       
115        barrier(CLK_LOCAL_MEM_FENCE);
116        for(int offset=WORK_GROUP_SIZE/2; offset>0; offset=offset>>1){
117                carry[idx] = carry[idx]|carry[idx^offset];
118                bubble[idx] = bubble[idx]|bubble[idx^offset];
119                barrier(CLK_LOCAL_MEM_FENCE);
120        }
121       
122        carry_mask = (carry[0]<<1)|group_carry[carryno];
123        bubble_mask = bubble[0];
124       
125        BitBlock s = (carry_mask + bubble_mask) & ~bubble_mask;
126        BitBlock dec = s | (s-carry_mask);
127        BitBlock rslt = partial_diff - ((dec>>idx)&0x1);
128        group_carry[carryno] = (carry[0]|(bubble_mask & dec))>>63;
129        return rslt;
130}
131
132inline BitBlock advance(int idx, BitBlock a, __local BitBlock *carry, BitBlock *group_carry, const int carryno){
133        carry[0] = group_carry[carryno];
134        BitBlock c = a<<1;
135        carry[idx+1] = (a & CARRY_BIT_MASK)>>63;
136        barrier(CLK_LOCAL_MEM_FENCE);
137        group_carry[carryno] = carry[WORK_GROUP_SIZE];
138        c = c | carry[idx];
139        return c;
140}
141
[1673]142inline BitBlock scanthru(int idx, BitBlock markers, BitBlock charclass, __local BitBlock *carry, __local BitBlock *bubble, BitBlock *group_carry, const int carryno) {
[1669]143        return simd_andc(adc(idx, markers, charclass, carry, bubble, group_carry, carryno), charclass);
144}
145
[1673]146inline BitBlock scanto(int idx, BitBlock markers, BitBlock charclass, __local BitBlock *carry, __local BitBlock *bubble, BitBlock *group_carry, const int carryno) {
[1669]147        return simd_and(adc(idx, markers, simd_not(charclass), carry, bubble, group_carry, carryno), charclass);
148}
149
[1673]150inline BitBlock scantofirst(int idx, BitBlock charclass, __local BitBlock *carry, __local BitBlock *bubble, BitBlock *group_carry, const int carryno) {
[1669]151        BitBlock marker;
152        group_carry[carryno] = simd_xor(group_carry[carryno], 1);
153    marker = adc(idx, 0, simd_not(charclass), carry, bubble, group_carry, carryno);
154    group_carry[carryno] = simd_xor(group_carry[carryno], 1);
155    return simd_and(marker, charclass);
156}
157
158inline void CarryInit(BitBlock * carry, int count){
159    for (int j=0; j < count; j++)
160         carry[j] = 0;
161}
162
[1673]163inline int CarryTest(BitBlock * cq, const int carryno, const int carry_count) {
[1669]164  BitBlock c1 = cq[carryno];
165  for (int i = carryno + 1; i < carryno + carry_count; i++) {
166    c1 = c1 | cq[i];
167  }
168  return bitblock_has_bit(c1);
169}
170
[1673]171inline void CarryCombine(BitBlock * cq, BitBlock * local_cq, const int carryno, const int carry_count) {
[1669]172  for (int i = 0; i < carry_count; i++) {
173    cq[carryno+i] = simd_or(cq[carryno+i], local_cq[i]);
174  }
175}
[1668]176/**************Parabix Structs***************/
177  struct Basis_bits {
178  BitBlock bit_0;
179  BitBlock bit_1;
180  BitBlock bit_2;
181  BitBlock bit_3;
182  BitBlock bit_4;
183  BitBlock bit_5;
184  BitBlock bit_6;
185  BitBlock bit_7;
186};
187
188  struct Lex {
189  BitBlock CR;
190  BitBlock LF;
191  BitBlock HT;
192  BitBlock SP;
193  BitBlock CRLF;
194  BitBlock RefStart;
195  BitBlock Semicolon;
196  BitBlock Colon;
197  BitBlock LAngle;
198  BitBlock RAngle;
199  BitBlock LBracket;
200  BitBlock RBracket;
201  BitBlock Exclam;
202  BitBlock QMark;
203  BitBlock Hyphen;
204  BitBlock Equals;
205  BitBlock SQuote;
206  BitBlock DQuote;
207  BitBlock Slash;
208  BitBlock Hash;
209  BitBlock x;
210  BitBlock ASCII_name_start;
211  BitBlock ASCII_name_char;
212  BitBlock NameScan;
213  BitBlock Digit;
214  BitBlock Hex;
215  BitBlock WS;
216  BitBlock error;
217};
218
[1669]219  struct U8 {
220  BitBlock suffix;
221  BitBlock FFFE_FFFF;
222  BitBlock error;
223};
224
225  struct Scope1 {
226  BitBlock RefStart;
227  BitBlock LAngle;
228  BitBlock Hyphen;
229  BitBlock QMark;
230  BitBlock RBracket;
231};
232
233  struct CtCDPI_Callouts {
234  BitBlock CD_end;
235  BitBlock Ct_starts;
236  BitBlock Ct_ends;
237  BitBlock CD_starts;
238  BitBlock CD_ends;
239  BitBlock PI_starts;
240  BitBlock PI_name_starts;
241  BitBlock PI_name_ends;
242  BitBlock PI_ends;
243  BitBlock CtCDPI_mask;
244  BitBlock error;
245};
246
[1673]247  struct Ref_Callouts {
248  BitBlock GenRef_starts;
249  BitBlock GenRef_ends;
250  BitBlock DecRef_starts;
251  BitBlock DecRef_ends;
252  BitBlock HexRef_starts;
253  BitBlock HexRef_ends;
254  BitBlock error;
255};
256
257  struct Tag_Callouts {
258  BitBlock ElemName_starts;
259  BitBlock ElemName_ends;
260  BitBlock AttName_starts;
261  BitBlock AttName_ends;
262  BitBlock AttVal_starts;
263  BitBlock AttVal_ends;
264  BitBlock AttVal_spans;
265  BitBlock EmptyTag_marks;
266  BitBlock EndTag_marks;
267  BitBlock LAngleFollow;
268  BitBlock error;
269};
270
271  struct Check_streams {
272  BitBlock non_ascii_name_starts;
273  BitBlock non_ascii_names;
274  BitBlock tag_marks;
275  BitBlock name_follows;
276  BitBlock att_refs;
277  BitBlock error_mask;
278};
279
280  struct Xml_names {
281  BitBlock namespace_error;
282};
283
[1668]284/**************Parabix Functions***************/
[1669]285#define classify_bytes_Validate_utf8(idx, basis_bits, lex, u8, carry, group_carry)\
[1668]286 do {\
287                BitBlock temp1, temp2, temp3, temp4, temp5, temp6, temp7, temp8, temp9, temp10;\
288                BitBlock temp11, temp12, temp13, temp14, temp15, temp16, temp17, temp18;\
289                BitBlock temp19, temp20, temp21, temp22, temp23, temp24, temp25, temp26;\
290                BitBlock temp27, temp28, temp29, temp30, temp31, temp32, temp33, temp34;\
291                BitBlock temp35, temp36, temp37, temp38, temp39, temp40, temp41, temp42;\
292                BitBlock temp43, temp44, temp45, temp46, temp47, temp48, x00_x1F, temp49;\
293                BitBlock temp50, temp51, temp52, temp53, temp54, temp55, temp56, temp57;\
294                BitBlock temp58, temp59, temp60, temp61, temp62, temp63, temp64, temp65;\
[1669]295                BitBlock u8anyscope, temp66, temp67, temp68, temp69, temp70, temp71;\
296                BitBlock unibyte, prefix, prefix2, prefix3, prefix4, badprefix;\
297                BitBlock xE0, xED, xF0, xF4, xA0_xBF, x80_x9F, x90_xBF, x80_x8F, xEF, xBF, xBE;\
298                BitBlock scope22, scope32, scope33, scope42, scope43, scope44, xE0_scope, xED_scope, xF0_scope, xF4_scope, xEF_scope;\
299                BitBlock temp72, temp73, u8lastscope, u8error1, u8error2;\
300                BitBlock u8error3, u8error4, EF_BF_pending, u8mismatch;\
[1668]301\
302        temp1 = simd_or(basis_bits.bit_0, basis_bits.bit_1);\
303        temp2 = simd_andc(basis_bits.bit_2, basis_bits.bit_3);\
304        temp3 = simd_andc(temp2, temp1);\
305        temp4 = simd_andc(basis_bits.bit_5, basis_bits.bit_4);\
306        temp5 = simd_andc(basis_bits.bit_6, basis_bits.bit_7);\
307        temp6 = simd_and(temp4, temp5);\
308        lex.RefStart = simd_and(temp3, temp6);\
309        temp7 = simd_and(basis_bits.bit_2, basis_bits.bit_3);\
310        temp8 = simd_andc(temp7, temp1);\
311        temp9 = simd_andc(basis_bits.bit_4, basis_bits.bit_5);\
312        temp10 = simd_and(basis_bits.bit_6, basis_bits.bit_7);\
313        temp11 = simd_and(temp9, temp10);\
314        lex.Semicolon = simd_and(temp8, temp11);\
315        temp12 = simd_and(basis_bits.bit_4, basis_bits.bit_5);\
316        temp13 = simd_or(basis_bits.bit_6, basis_bits.bit_7);\
317        temp14 = simd_andc(temp12, temp13);\
318        lex.LAngle = simd_and(temp8, temp14);\
319        temp15 = simd_and(temp12, temp5);\
320        lex.RAngle = simd_and(temp8, temp15);\
321        temp16 = simd_andc(basis_bits.bit_1, basis_bits.bit_0);\
322        temp17 = simd_andc(basis_bits.bit_3, basis_bits.bit_2);\
323        temp18 = simd_and(temp16, temp17);\
324        lex.LBracket = simd_and(temp18, temp11);\
325        temp19 = simd_andc(basis_bits.bit_7, basis_bits.bit_6);\
326        temp20 = simd_and(temp12, temp19);\
327        lex.RBracket = simd_and(temp18, temp20);\
328        temp21 = simd_or(basis_bits.bit_4, basis_bits.bit_5);\
329        temp22 = simd_andc(temp19, temp21);\
330        lex.Exclam = simd_and(temp3, temp22);\
331        temp23 = simd_and(temp12, temp10);\
332        lex.QMark = simd_and(temp8, temp23);\
333        lex.Hyphen = simd_and(temp3, temp20);\
334        lex.Equals = simd_and(temp8, temp20);\
335        temp24 = simd_and(temp4, temp10);\
336        lex.SQuote = simd_and(temp3, temp24);\
337        temp25 = simd_andc(temp5, temp21);\
338        lex.DQuote = simd_and(temp3, temp25);\
339        lex.Slash = simd_and(temp3, temp23);\
340        temp26 = simd_andc(temp10, temp21);\
341        lex.Hash = simd_and(temp3, temp26);\
342        temp27 = simd_and(temp16, temp7);\
343        temp28 = simd_andc(temp9, temp13);\
344        lex.x = simd_and(temp27, temp28);\
345        temp29 = simd_and(temp9, temp5);\
346        lex.Colon = simd_and(temp8, temp29);\
347        temp30 = simd_and(temp18, temp23);\
348        temp31 = simd_or(temp30, lex.Colon);\
349        temp32 = simd_andc(temp16, basis_bits.bit_2);\
350        temp33 = simd_or(basis_bits.bit_5, temp10);\
351        temp34 = simd_and(basis_bits.bit_4, temp33);\
352        temp35 = simd_not(temp34);\
353        temp36 = simd_or(temp21, temp13);\
354        temp37 = simd_or(simd_and(basis_bits.bit_3, temp35), simd_andc(temp36, basis_bits.bit_3));\
355        temp38 = simd_and(temp32, temp37);\
356        temp39 = simd_or(temp31, temp38);\
357        temp40 = simd_and(temp16, basis_bits.bit_2);\
358        temp41 = simd_and(temp40, temp37);\
359        lex.ASCII_name_start = simd_or(temp39, temp41);\
360        temp42 = simd_or(temp30, lex.Hyphen);\
361        temp43 = simd_and(temp3, temp15);\
362        temp44 = simd_or(temp42, temp43);\
363        temp45 = simd_andc(temp8, temp34);\
364        temp46 = simd_or(temp44, temp45);\
365        temp47 = simd_or(temp46, temp38);\
366        lex.ASCII_name_char = simd_or(temp47, temp41);\
367        lex.NameScan = simd_or(lex.ASCII_name_char, basis_bits.bit_0);\
368        temp48 = simd_or(temp1, basis_bits.bit_2);\
369        x00_x1F = simd_not(temp48);\
370        temp49 = simd_or(basis_bits.bit_2, basis_bits.bit_3);\
371        temp50 = simd_or(temp1, temp49);\
372        lex.CR = simd_andc(temp20, temp50);\
373        lex.LF = simd_andc(temp29, temp50);\
374        temp51 = simd_and(temp9, temp19);\
375        lex.HT = simd_andc(temp51, temp50);\
376        lex.SP = simd_andc(temp3, temp36);\
377        temp52 = simd_or(temp20, temp29);\
378        temp53 = simd_or(temp52, temp51);\
379        temp54 = simd_andc(temp53, temp50);\
380        lex.WS = simd_or(temp54, lex.SP);\
381        temp55 = simd_or(basis_bits.bit_5, basis_bits.bit_6);\
382        temp56 = simd_and(basis_bits.bit_4, temp55);\
383        lex.Digit = simd_andc(temp8, temp56);\
384        temp57 = simd_andc(temp16, temp49);\
385        temp58 = simd_andc(temp57, basis_bits.bit_4);\
386        temp59 = simd_not(temp10);\
387        temp60 = simd_or(simd_and(basis_bits.bit_5, temp59), simd_andc(temp13, basis_bits.bit_5));\
388        temp61 = simd_and(temp58, temp60);\
389        temp62 = simd_or(lex.Digit, temp61);\
390        temp63 = simd_and(temp16, temp2);\
391        temp64 = simd_andc(temp63, basis_bits.bit_4);\
392        temp65 = simd_and(temp64, temp60);\
393        lex.Hex = simd_or(temp62, temp65);\
394        lex.error = simd_andc(x00_x1F, lex.WS);\
[1669]395                unibyte = simd_not(basis_bits.bit_0);\
396        u8.suffix = simd_const_1(0);\
397        u8.error = simd_const_1(0);\
398        u8.FFFE_FFFF = simd_const_1(0);\
399        u8anyscope = simd_const_1(0);\
400        if ((bitblock_has_bit(basis_bits.bit_0) || CarryTest(group_carry, 0, 11))) {\
401          prefix = simd_and(basis_bits.bit_0, basis_bits.bit_1);\
402          prefix2 = simd_andc(prefix, basis_bits.bit_2);\
403          prefix3 = simd_and(prefix, temp2);\
404          prefix4 = simd_and(prefix, temp7);\
405          u8.suffix = simd_andc(basis_bits.bit_0, basis_bits.bit_1);\
406          temp66 = simd_andc(prefix, temp49);\
407          temp67 = simd_or(temp21, basis_bits.bit_6);\
408          temp68 = simd_andc(temp66, temp67);\
409          temp69 = simd_and(basis_bits.bit_5, temp13);\
410          temp70 = simd_or(basis_bits.bit_4, temp69);\
411          temp71 = simd_and(prefix4, temp70);\
412          badprefix = simd_or(temp68, temp71);\
413          u8.error = badprefix;\
414          scope22 = advance(idx, prefix2, carry, group_carry, 0);\
415          u8anyscope = scope22;\
416          if (bitblock_has_bit(simd_or(prefix3, prefix4))) {\
417            xE0 = simd_andc(prefix3, temp36);\
418            xED = simd_and(prefix3, temp20);\
419            xF0 = simd_andc(prefix4, temp36);\
420            temp72 = simd_andc(temp4, temp13);\
421            xF4 = simd_and(prefix4, temp72);\
422            xA0_xBF = simd_and(u8.suffix, basis_bits.bit_2);\
423            x80_x9F = simd_andc(u8.suffix, basis_bits.bit_2);\
424            x90_xBF = simd_and(u8.suffix, temp49);\
425            x80_x8F = simd_andc(u8.suffix, temp49);\
426            xEF = simd_and(prefix3, temp23);\
427            temp73 = simd_and(u8.suffix, temp7);\
428            xBF = simd_and(temp73, temp23);\
429            xBE = simd_and(temp73, temp15);\
430            xE0_scope = advance(idx, xE0, carry, group_carry, 1);\
431            xED_scope = advance(idx, xED, carry, group_carry, 2);\
432            xF0_scope = advance(idx, xF0, carry, group_carry, 3);\
433            xF4_scope = advance(idx, xF4, carry, group_carry, 4);\
434            xEF_scope = advance(idx, xEF, carry, group_carry, 5);\
435            scope32 = advance(idx, prefix3, carry, group_carry, 6);\
436            scope33 = advance(idx, scope32, carry, group_carry, 7);\
437            scope42 = advance(idx, prefix4, carry, group_carry, 8);\
438            scope43 = advance(idx, scope42, carry, group_carry, 9);\
439            scope44 = advance(idx, scope43, carry, group_carry, 10);\
440            u8lastscope = simd_or(simd_or(scope22, scope33), scope44);\
441            u8anyscope = simd_or(simd_or(simd_or(u8lastscope, scope32), scope42), scope43);\
442            u8error1 = simd_and(xE0_scope, x80_x9F);\
443            u8error2 = simd_and(xED_scope, xA0_xBF);\
444            u8error3 = simd_and(xF0_scope, x80_x8F);\
445            u8error4 = simd_and(xF4_scope, x90_xBF);\
446            u8.error = simd_or(u8.error, simd_or(simd_or(simd_or(u8error1, u8error2), u8error3), u8error4));\
447            EF_BF_pending = advance(idx, simd_and(xEF_scope, xBF), carry, group_carry, 11);\
448            u8.FFFE_FFFF = simd_and(EF_BF_pending, simd_or(xBE, xBF));\
449          }\
450        }\
451        u8mismatch = simd_xor(u8anyscope, u8.suffix);\
452        u8.error = simd_or(u8.error, u8mismatch);\
[1668]453  } while (0)
454
[1669]455 #define add_scope_streams(idx, lex, scope1, carry, group_carry)\
456 do {\
457                BitBlock v, w, v1, w1;\
458\
459        v = simd_or(lex.LAngle, lex.Hyphen);\
460        w = simd_or(lex.Hyphen, lex.QMark);\
461        v1 = advance(idx, v, carry, group_carry, 0);\
462        w1 = advance(idx, w, carry, group_carry, 1);\
463        scope1.LAngle = simd_andc(v1, w1);\
464        scope1.Hyphen = simd_and(v1, w1);\
465        scope1.QMark = simd_andc(w1, v1);\
466  } while (0)
467
468 #define parse_CtCDPI(idx, ctCDPI_Callouts, lex, scope1, carry, bubble, group_carry)\
469 do {\
470                BitBlock CtCDPI_starts, Ct_errors, PI_start, CtCD_start, CtCDPI_start;\
471                BitBlock DoubleHyphen, PI_end, CtCDPI_Cursor, PI_Cursor, CD_Ct_Cursor;\
472                BitBlock CD_Cursor, Ct_Cursor, PI_name_end;\
473\
474        ctCDPI_Callouts.CD_end = simd_const_1(0);\
475        ctCDPI_Callouts.Ct_starts = simd_const_1(0);\
476        ctCDPI_Callouts.Ct_ends = simd_const_1(0);\
477        ctCDPI_Callouts.CD_starts = simd_const_1(0);\
478        ctCDPI_Callouts.CD_ends = simd_const_1(0);\
479        ctCDPI_Callouts.PI_starts = simd_const_1(0);\
480        ctCDPI_Callouts.PI_name_starts = simd_const_1(0);\
481        ctCDPI_Callouts.PI_name_ends = simd_const_1(0);\
482        ctCDPI_Callouts.PI_ends = simd_const_1(0);\
483        ctCDPI_Callouts.CtCDPI_mask = simd_const_1(0);\
484        ctCDPI_Callouts.error = simd_const_1(0);\
485        CtCDPI_starts = simd_const_1(0);\
486        Ct_errors = simd_const_1(0);\
487        if ((bitblock_has_bit(lex.RBracket) || CarryTest(group_carry, 0, 2))) {\
488          scope1.RBracket = advance(idx, lex.RBracket, carry, group_carry, 0);\
489          ctCDPI_Callouts.CD_end = simd_and(advance(idx, simd_and(scope1.RBracket, lex.RBracket), carry, group_carry, 1), lex.RAngle);\
490        }\
491        PI_start = simd_and(scope1.LAngle, lex.QMark);\
492        CtCD_start = simd_and(scope1.LAngle, lex.Exclam);\
493        CtCDPI_start = simd_or(PI_start, CtCD_start);\
494        DoubleHyphen = simd_and(scope1.Hyphen, lex.Hyphen);\
495        PI_end = simd_and(scope1.QMark, lex.RAngle);\
496        CtCDPI_Cursor = scantofirst(idx, CtCDPI_start, carry, bubble, group_carry, 2);\
497        if ((bitblock_has_bit(CtCDPI_Cursor) || CarryTest(group_carry, 3, 14))) {\
498          CtCDPI_starts = simd_or(CtCDPI_starts, CtCDPI_Cursor);\
499          PI_Cursor = simd_and(CtCDPI_Cursor, PI_start);\
500          CD_Ct_Cursor = advance(idx, simd_andc(CtCDPI_Cursor, PI_Cursor), carry, group_carry, 3);\
501          CD_Cursor = simd_and(CD_Ct_Cursor, lex.LBracket);\
502          Ct_Cursor = simd_and(CD_Ct_Cursor, lex.Hyphen);\
503          ctCDPI_Callouts.PI_starts = simd_or(ctCDPI_Callouts.PI_starts, PI_Cursor);\
504          ctCDPI_Callouts.CD_starts = simd_or(ctCDPI_Callouts.CD_starts, CD_Cursor);\
505          ctCDPI_Callouts.Ct_starts = simd_or(ctCDPI_Callouts.Ct_starts, Ct_Cursor);\
506          Ct_Cursor = advance(idx, Ct_Cursor, carry, group_carry, 4);\
507          Ct_errors = simd_or(Ct_errors, simd_andc(Ct_Cursor, lex.Hyphen));\
508          Ct_Cursor = advance(idx, advance(idx, Ct_Cursor, carry, group_carry, 5), carry, group_carry, 6);\
509          PI_Cursor = advance(idx, PI_Cursor, carry, group_carry, 7);\
510          ctCDPI_Callouts.PI_name_starts = simd_or(ctCDPI_Callouts.PI_name_starts, PI_Cursor);\
511          PI_name_end = scanthru(idx, PI_Cursor, lex.NameScan, carry, bubble, group_carry, 8);\
512          ctCDPI_Callouts.PI_name_ends = simd_or(ctCDPI_Callouts.PI_name_ends, PI_name_end);\
513          PI_Cursor = scanto(idx, PI_name_end, PI_end, carry, bubble, group_carry, 9);\
514          CD_Cursor = scanto(idx, CD_Cursor, ctCDPI_Callouts.CD_end, carry, bubble, group_carry, 10);\
515          Ct_Cursor = advance(idx, scanto(idx, Ct_Cursor, DoubleHyphen, carry, bubble, group_carry, 11), carry, group_carry, 12);\
516          ctCDPI_Callouts.PI_ends = simd_or(ctCDPI_Callouts.PI_ends, PI_Cursor);\
517          ctCDPI_Callouts.CD_ends = simd_or(ctCDPI_Callouts.CD_ends, CD_Cursor);\
518          ctCDPI_Callouts.Ct_ends = simd_or(ctCDPI_Callouts.Ct_ends, Ct_Cursor);\
519          CtCDPI_Cursor = simd_or(simd_or(PI_Cursor, CD_Cursor), Ct_Cursor);\
520          CtCDPI_Cursor = scanto(idx, CtCDPI_Cursor, CtCDPI_start, carry, bubble, group_carry, 13);\
521          ctCDPI_Callouts.CtCDPI_mask = sbb(idx, advance(idx, simd_or(simd_or(ctCDPI_Callouts.CD_ends, ctCDPI_Callouts.Ct_ends), ctCDPI_Callouts.PI_ends), carry, group_carry, 14), CtCDPI_starts, carry, bubble, group_carry, 15);\
522          ctCDPI_Callouts.error = simd_or(Ct_errors, simd_andc(ctCDPI_Callouts.Ct_ends, lex.RAngle));\
523          ctCDPI_Callouts.error = simd_or(ctCDPI_Callouts.error, simd_andc(advance(idx, simd_andc(ctCDPI_Callouts.PI_name_ends, lex.WS), carry, group_carry, 16), PI_end));\
524          ctCDPI_Callouts.error = simd_or(ctCDPI_Callouts.error, simd_and(ctCDPI_Callouts.PI_name_starts, ctCDPI_Callouts.PI_name_ends));\
525          ctCDPI_Callouts.error = simd_or(ctCDPI_Callouts.error, simd_andc(ctCDPI_Callouts.CtCDPI_mask, EOF_mask));\
526          while (bitblock_has_bit(CtCDPI_Cursor)) {\
527            BitBlock subcarryQ[14];\
528            CarryInit(subcarryQ, 14);\
529            CtCDPI_starts = simd_or(CtCDPI_starts, CtCDPI_Cursor);\
530            PI_Cursor = simd_and(CtCDPI_Cursor, PI_start);\
531            CD_Ct_Cursor = advance(idx, simd_andc(CtCDPI_Cursor, PI_Cursor), carry, subcarryQ, 0);\
532            CD_Cursor = simd_and(CD_Ct_Cursor, lex.LBracket);\
533            Ct_Cursor = simd_and(CD_Ct_Cursor, lex.Hyphen);\
534            ctCDPI_Callouts.PI_starts = simd_or(ctCDPI_Callouts.PI_starts, PI_Cursor);\
535            ctCDPI_Callouts.CD_starts = simd_or(ctCDPI_Callouts.CD_starts, CD_Cursor);\
536            ctCDPI_Callouts.Ct_starts = simd_or(ctCDPI_Callouts.Ct_starts, Ct_Cursor);\
537            Ct_Cursor = advance(idx, Ct_Cursor, carry, subcarryQ, 1);\
538            Ct_errors = simd_or(Ct_errors, simd_andc(Ct_Cursor, lex.Hyphen));\
539            Ct_Cursor = advance(idx, advance(idx, Ct_Cursor, carry, subcarryQ, 2), carry, subcarryQ, 3);\
540            PI_Cursor = advance(idx, PI_Cursor, carry, subcarryQ, 4);\
541            ctCDPI_Callouts.PI_name_starts = simd_or(ctCDPI_Callouts.PI_name_starts, PI_Cursor);\
542            PI_name_end = scanthru(idx, PI_Cursor, lex.NameScan, carry, bubble, subcarryQ, 5);\
543            ctCDPI_Callouts.PI_name_ends = simd_or(ctCDPI_Callouts.PI_name_ends, PI_name_end);\
544            PI_Cursor = scanto(idx, PI_name_end, PI_end, carry, bubble, subcarryQ, 6);\
545            CD_Cursor = scanto(idx, CD_Cursor, ctCDPI_Callouts.CD_end, carry, bubble, subcarryQ, 7);\
546            Ct_Cursor = advance(idx, scanto(idx, Ct_Cursor, DoubleHyphen, carry, bubble, subcarryQ, 8), carry, subcarryQ, 9);\
547            ctCDPI_Callouts.PI_ends = simd_or(ctCDPI_Callouts.PI_ends, PI_Cursor);\
548            ctCDPI_Callouts.CD_ends = simd_or(ctCDPI_Callouts.CD_ends, CD_Cursor);\
549            ctCDPI_Callouts.Ct_ends = simd_or(ctCDPI_Callouts.Ct_ends, Ct_Cursor);\
550            CtCDPI_Cursor = simd_or(simd_or(PI_Cursor, CD_Cursor), Ct_Cursor);\
551            CtCDPI_Cursor = scanto(idx, CtCDPI_Cursor, CtCDPI_start, carry, bubble, subcarryQ, 10);\
552            ctCDPI_Callouts.CtCDPI_mask = sbb(idx, advance(idx, simd_or(simd_or(ctCDPI_Callouts.CD_ends, ctCDPI_Callouts.Ct_ends), ctCDPI_Callouts.PI_ends), carry, subcarryQ, 11), CtCDPI_starts, carry, bubble, subcarryQ, 12);\
553            ctCDPI_Callouts.error = simd_or(Ct_errors, simd_andc(ctCDPI_Callouts.Ct_ends, lex.RAngle));\
554            ctCDPI_Callouts.error = simd_or(ctCDPI_Callouts.error, simd_andc(advance(idx, simd_andc(ctCDPI_Callouts.PI_name_ends, lex.WS), carry, subcarryQ, 13), PI_end));\
555            ctCDPI_Callouts.error = simd_or(ctCDPI_Callouts.error, simd_and(ctCDPI_Callouts.PI_name_starts, ctCDPI_Callouts.PI_name_ends));\
556            ctCDPI_Callouts.error = simd_or(ctCDPI_Callouts.error, simd_andc(ctCDPI_Callouts.CtCDPI_mask, EOF_mask));\
557            CarryCombine(group_carry, subcarryQ, 3, 14);\
558          }\
559        }\
560  } while (0)
561
[1673]562#define parse_refs(idx, lex, scope1, ctCDPI_Callouts, ref_Callouts, carry, bubble, group_carry)\
563 do {\
564                BitBlock Ref1, NumRef2, NumRef3, HexRef3, ref_error1, ref_error2, ref_ends;\
565                BitBlock ref_error3;\
566\
567        ref_Callouts.GenRef_starts = simd_const_1(0);\
568        ref_Callouts.GenRef_ends = simd_const_1(0);\
569        ref_Callouts.DecRef_starts = simd_const_1(0);\
570        ref_Callouts.DecRef_ends = simd_const_1(0);\
571        ref_Callouts.HexRef_starts = simd_const_1(0);\
572        ref_Callouts.HexRef_ends = simd_const_1(0);\
573        ref_Callouts.error = simd_const_1(0);\
574        Ref1 = simd_andc(lex.RefStart, ctCDPI_Callouts.CtCDPI_mask);\
575        if ((bitblock_has_bit(Ref1) || CarryTest(group_carry, 0, 6))) {\
576          scope1.RefStart = advance(idx, Ref1, carry, group_carry, 0);\
577          NumRef2 = simd_and(scope1.RefStart, lex.Hash);\
578          ref_Callouts.GenRef_starts = simd_andc(scope1.RefStart, lex.Hash);\
579          NumRef3 = advance(idx, NumRef2, carry, group_carry, 1);\
580          HexRef3 = simd_and(NumRef3, lex.x);\
581          ref_Callouts.DecRef_starts = simd_andc(NumRef3, lex.x);\
582          ref_Callouts.HexRef_starts = advance(idx, HexRef3, carry, group_carry, 2);\
583          ref_Callouts.GenRef_ends = scanthru(idx, ref_Callouts.GenRef_starts, lex.NameScan, carry, bubble, group_carry, 3);\
584          ref_Callouts.DecRef_ends = scanthru(idx, ref_Callouts.DecRef_starts, lex.Digit, carry, bubble, group_carry, 4);\
585          ref_Callouts.HexRef_ends = scanthru(idx, ref_Callouts.HexRef_starts, lex.Hex, carry, bubble, group_carry, 5);\
586          ref_error1 = simd_andc(ref_Callouts.DecRef_starts, lex.Digit);\
587          ref_error2 = simd_andc(ref_Callouts.HexRef_starts, lex.Hex);\
588          ref_ends = simd_or(simd_or(ref_Callouts.GenRef_ends, ref_Callouts.DecRef_ends), ref_Callouts.HexRef_ends);\
589          ref_error3 = simd_andc(ref_ends, lex.Semicolon);\
590          ref_Callouts.error = simd_or(simd_or(ref_error1, ref_error2), ref_error3);\
591        }\
592  } while (0)
593
594 #define parse_tags(idx, lex, scope1, ctCDPI_Callouts, tag_Callouts, carry, bubble, group_carry)\
595 do {\
596                BitBlock DQuoteDelim, SQuoteDelim, AttListDelim, ParseError, EqToCheck;\
597                BitBlock AttValEnds, AfterWS, AttListEnd, AttNameStart, AttNameFollow;\
598                BitBlock EqExpected, AttValPos, DQuoteAttVal, SQuoteAttVal, DQuoteAttEnd;\
599                BitBlock SQuoteAttEnd, AttValEnd, AttValFollow, STagEnds, EndTagEnds;\
600\
601        DQuoteDelim = simd_or(lex.DQuote, lex.LAngle);\
602        SQuoteDelim = simd_or(lex.SQuote, lex.LAngle);\
603        AttListDelim = simd_or(lex.Slash, lex.RAngle);\
604        tag_Callouts.LAngleFollow = simd_andc(scope1.LAngle, ctCDPI_Callouts.CtCDPI_mask);\
605        tag_Callouts.ElemName_starts = simd_andc(tag_Callouts.LAngleFollow, lex.Slash);\
606        tag_Callouts.EndTag_marks = simd_and(tag_Callouts.LAngleFollow, lex.Slash);\
607        tag_Callouts.ElemName_ends = scanthru(idx, tag_Callouts.ElemName_starts, lex.NameScan, carry, bubble, group_carry, 0);\
608        ParseError = simd_and(tag_Callouts.ElemName_starts, tag_Callouts.ElemName_ends);\
609        tag_Callouts.AttName_starts = simd_const_1(0);\
610        tag_Callouts.AttName_ends = simd_const_1(0);\
611        EqToCheck = simd_const_1(0);\
612        tag_Callouts.AttVal_starts = simd_const_1(0);\
613        AttValEnds = simd_const_1(0);\
614        tag_Callouts.AttVal_ends = simd_const_1(0);\
615        AfterWS = scanthru(idx, tag_Callouts.ElemName_ends, lex.WS, carry, bubble, group_carry, 1);\
616        AttListEnd = simd_and(AfterWS, AttListDelim);\
617        AttNameStart = simd_andc(AfterWS, AttListDelim);\
618        ParseError = simd_or(ParseError, simd_and(tag_Callouts.ElemName_ends, AttNameStart));\
619        if ((bitblock_has_bit(AttNameStart) || CarryTest(group_carry, 2, 7))) {\
620          tag_Callouts.AttName_starts = simd_or(tag_Callouts.AttName_starts, AttNameStart);\
621          AttNameFollow = scanthru(idx, AttNameStart, lex.NameScan, carry, bubble, group_carry, 2);\
622          tag_Callouts.AttName_ends = simd_or(tag_Callouts.AttName_ends, AttNameFollow);\
623          EqExpected = scanthru(idx, AttNameFollow, lex.WS, carry, bubble, group_carry, 3);\
624          EqToCheck = simd_or(EqToCheck, EqExpected);\
625          AttValPos = scanthru(idx, EqExpected, simd_or(EqExpected, lex.WS), carry, bubble, group_carry, 4);\
626          tag_Callouts.AttVal_starts = simd_or(tag_Callouts.AttVal_starts, AttValPos);\
627          DQuoteAttVal = simd_and(AttValPos, lex.DQuote);\
628          SQuoteAttVal = simd_and(AttValPos, lex.SQuote);\
629          DQuoteAttEnd = scanto(idx, DQuoteAttVal, simd_andc(DQuoteDelim, DQuoteAttVal), carry, bubble, group_carry, 5);\
630          SQuoteAttEnd = scanto(idx, SQuoteAttVal, simd_andc(SQuoteDelim, SQuoteAttVal), carry, bubble, group_carry, 6);\
631          AttValEnd = simd_or(DQuoteAttEnd, SQuoteAttEnd);\
632          AttValEnds = simd_or(AttValEnds, AttValEnd);\
633          AttValFollow = advance(idx, AttValEnd, carry, group_carry, 7);\
634          tag_Callouts.AttVal_ends = simd_or(tag_Callouts.AttVal_ends, AttValFollow);\
635          AfterWS = scanthru(idx, AttValFollow, lex.WS, carry, bubble, group_carry, 8);\
636          AttListEnd = simd_or(AttListEnd, simd_and(AfterWS, AttListDelim));\
637          AttNameStart = simd_andc(AfterWS, AttListDelim);\
638          while (bitblock_has_bit(AttNameStart)) {\
639            BitBlock subcarryQ[7];\
640            CarryInit(subcarryQ, 7);\
641            tag_Callouts.AttName_starts = simd_or(tag_Callouts.AttName_starts, AttNameStart);\
642            AttNameFollow = scanthru(idx, AttNameStart, lex.NameScan, carry, bubble, subcarryQ, 0);\
643            tag_Callouts.AttName_ends = simd_or(tag_Callouts.AttName_ends, AttNameFollow);\
644            EqExpected = scanthru(idx, AttNameFollow, lex.WS,  carry, bubble, subcarryQ, 1);\
645            EqToCheck = simd_or(EqToCheck, EqExpected);\
646            AttValPos = scanthru(idx, EqExpected, simd_or(EqExpected, lex.WS), carry, bubble, subcarryQ, 2);\
647            tag_Callouts.AttVal_starts = simd_or(tag_Callouts.AttVal_starts, AttValPos);\
648            DQuoteAttVal = simd_and(AttValPos, lex.DQuote);\
649            SQuoteAttVal = simd_and(AttValPos, lex.SQuote);\
650            DQuoteAttEnd = scanto(idx, DQuoteAttVal, simd_andc(DQuoteDelim, DQuoteAttVal), carry, bubble, subcarryQ, 3);\
651            SQuoteAttEnd = scanto(idx, SQuoteAttVal, simd_andc(SQuoteDelim, SQuoteAttVal), carry, bubble, subcarryQ, 4);\
652            AttValEnd = simd_or(DQuoteAttEnd, SQuoteAttEnd);\
653            AttValEnds = simd_or(AttValEnds, AttValEnd);\
654            AttValFollow = advance(idx, AttValEnd,  carry, subcarryQ, 5);\
655            tag_Callouts.AttVal_ends = simd_or(tag_Callouts.AttVal_ends, AttValFollow);\
656            AfterWS = scanthru(idx, AttValFollow, lex.WS,  carry, bubble, subcarryQ, 6);\
657            AttListEnd = simd_or(AttListEnd, simd_and(AfterWS, AttListDelim));\
658            AttNameStart = simd_andc(AfterWS, AttListDelim);\
659            CarryCombine(group_carry, subcarryQ, 2, 7);\
660          }\
661        }\
662        STagEnds = simd_and(AttListEnd, lex.RAngle);\
663        tag_Callouts.EmptyTag_marks = advance(idx, simd_and(AttListEnd, lex.Slash), carry, group_carry, 9);\
664        ParseError = simd_or(ParseError, simd_and(tag_Callouts.AttVal_ends, tag_Callouts.AttName_starts));\
665        ParseError = simd_or(ParseError, simd_and(tag_Callouts.AttName_starts, tag_Callouts.AttName_ends));\
666        ParseError = simd_or(ParseError, simd_andc(EqToCheck, lex.Equals));\
667        ParseError = simd_or(ParseError, simd_andc(tag_Callouts.AttVal_starts, simd_or(lex.DQuote, lex.SQuote)));\
668        ParseError = simd_or(ParseError, simd_andc(AttValEnds, simd_or(lex.DQuote, lex.SQuote)));\
669        ParseError = simd_or(ParseError, simd_andc(tag_Callouts.EmptyTag_marks, lex.RAngle));\
670        EndTagEnds = scanthru(idx, scanthru(idx, tag_Callouts.EndTag_marks, simd_or(tag_Callouts.EndTag_marks, lex.NameScan), carry, bubble, group_carry, 10), lex.WS, carry, bubble, group_carry, 11);\
671        ParseError = simd_or(ParseError, simd_andc(EndTagEnds, lex.RAngle));\
672        tag_Callouts.error = ParseError;\
673        tag_Callouts.AttVal_spans = sbb(idx, tag_Callouts.AttVal_ends, tag_Callouts.AttVal_starts, carry, bubble, group_carry, 12);\
674  } while (0)
675
676   #define validate_xml_names(idx, ctCDPI_Callouts, ref_Callouts, tag_Callouts, lex, u8, xml_names, check_streams, carry, bubble, group_carry)\
677 do {\
678                BitBlock PI_names, GenRefs, ElemNames, AttNames, qname_stream, ncname_stream;\
679                BitBlock name_stream, name_start, name_cursor, void_prefix_err, namespace_sep;\
680                BitBlock local_part_start, local_part_err, colon2_err, ncname_err;\
681\
682        PI_names = sbb(idx, ctCDPI_Callouts.PI_name_ends, ctCDPI_Callouts.PI_name_starts, carry, bubble, group_carry, 0);\
683        GenRefs = sbb(idx, ref_Callouts.GenRef_ends, ref_Callouts.GenRef_starts, carry, bubble, group_carry, 1);\
684        ElemNames = sbb(idx, tag_Callouts.ElemName_ends, tag_Callouts.ElemName_starts, carry, bubble, group_carry, 2);\
685        AttNames = sbb(idx, tag_Callouts.AttName_ends, tag_Callouts.AttName_starts, carry, bubble, group_carry, 3);\
686        qname_stream = simd_or(ElemNames, AttNames);\
687        ncname_stream = simd_or(PI_names, GenRefs);\
688        name_stream = simd_or(qname_stream, ncname_stream);\
689        name_start = simd_andc(name_stream, advance(idx, name_stream, carry, group_carry, 4));\
690        name_cursor = simd_andc(name_stream, advance(idx, name_stream, carry, group_carry, 5));\
691        void_prefix_err = simd_and(name_cursor, lex.Colon);\
692        namespace_sep = simd_and(scanthru(idx, name_cursor, simd_andc(lex.NameScan, lex.Colon), carry, bubble, group_carry, 6), lex.Colon);\
693        local_part_start = advance(idx, namespace_sep, carry, group_carry, 7);\
694        local_part_err = simd_andc(local_part_start, lex.NameScan);\
695        colon2_err = simd_and(scanthru(idx, local_part_start, simd_andc(lex.NameScan, lex.Colon), carry, bubble, group_carry, 8), lex.Colon);\
696        ncname_err = simd_and(ncname_stream, lex.Colon);\
697        xml_names.namespace_error = simd_or(simd_or(simd_or(void_prefix_err, local_part_err), colon2_err), ncname_err);\
698        check_streams.non_ascii_name_starts = simd_andc(name_start, lex.ASCII_name_start);\
699        check_streams.non_ascii_names = simd_andc(simd_andc(simd_andc(name_stream, name_start), lex.ASCII_name_char), u8.suffix);\
700  } while (0)
701
702  #define do_check_streams(do_check_streams, ctCDPI_Callouts, tag_Callouts, lex, u8, scope1, ref_Callouts, xml_names, check_streams)\
703 do {\
704                BitBlock CD_end_error;\
705\
706        CD_end_error = simd_andc(ctCDPI_Callouts.CD_end, simd_or(ctCDPI_Callouts.CtCDPI_mask, tag_Callouts.AttVal_spans));\
707        check_streams.error_mask = simd_or(simd_or(simd_or(simd_or(simd_or(simd_or(simd_or(simd_and(lex.error, EOF_mask), u8.error), u8.FFFE_FFFF), ctCDPI_Callouts.error), tag_Callouts.error), CD_end_error), ref_Callouts.error), xml_names.namespace_error);\
708        check_streams.tag_marks = simd_or(simd_or(tag_Callouts.EmptyTag_marks, tag_Callouts.LAngleFollow), tag_Callouts.AttName_starts);\
709        check_streams.name_follows = simd_or(tag_Callouts.ElemName_ends, tag_Callouts.AttName_ends);\
710        check_streams.att_refs = simd_and(tag_Callouts.AttVal_spans, scope1.RefStart);\
711  } while (0)
712
[1668]713/*************Kernel Function****************/
714__kernel
715void xmlwf(__global BitBlock *s, 
716                   __global BitBlock *err_pos,
717                        __local BitBlock *carry,
718                        __local BitBlock *bubble) {
719
720   int i = get_local_id(0);
721   struct Basis_bits basis_bits;
722   struct Lex lex;
723   struct U8 u8;
724   struct Scope1 scope1;
725   struct CtCDPI_Callouts ctCDPI_Callouts;
[1673]726   struct Ref_Callouts ref_Callouts;
727   struct Tag_Callouts tag_Callouts;
728   struct Check_streams check_streams;
729   struct Xml_names xml_names;
[1668]730
[1669]731   BitBlock classify_bytes_Validate_utf8_carryQ[12];
732   CarryInit(classify_bytes_Validate_utf8_carryQ, 12);
733   BitBlock add_scope_streams_carry[2];
734   CarryInit(add_scope_streams_carry, 2);
735   BitBlock parse_CtCDPI_carryQ[17];
736   CarryInit(parse_CtCDPI_carryQ, 17);
[1673]737   BitBlock  parse_refs_carryQ[6];
738   CarryInit(parse_refs_carryQ, 6);
739   BitBlock  parse_tags_carryQ[13];
740   CarryInit(parse_tags_carryQ, 13);
741   BitBlock  validate_xml_names_carryQ[9];
742   CarryInit(validate_xml_names_carryQ, 9);
[1669]743   
[1668]744   for(int idx=i; idx<ELEMENTS;idx=idx+WORK_GROUP_SIZE){
745                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], 
746                basis_bits.bit_0, basis_bits.bit_1, basis_bits.bit_2, basis_bits.bit_3, 
747                basis_bits.bit_4, basis_bits.bit_5, basis_bits.bit_6, basis_bits.bit_7);
748               
[1669]749                classify_bytes_Validate_utf8(i, basis_bits, lex, u8, carry, classify_bytes_Validate_utf8_carryQ);
750                add_scope_streams(i, lex, scope1, carry, add_scope_streams_carry);
[1673]751                parse_CtCDPI(i, ctCDPI_Callouts, lex, scope1, carry, bubble, parse_CtCDPI_carryQ);
752                parse_refs(i, lex, scope1, ctCDPI_Callouts, ref_Callouts, carry, bubble, parse_refs_carryQ);
753                parse_tags(i, lex, scope1, ctCDPI_Callouts, tag_Callouts, carry, bubble, parse_tags_carryQ);
754                validate_xml_names(idx, ctCDPI_Callouts, ref_Callouts, tag_Callouts, lex, u8, xml_names, check_streams, carry, bubble, validate_xml_names_carryQ);
755                do_check_streams(do_check_streams, ctCDPI_Callouts, tag_Callouts, lex, u8, scope1, ref_Callouts, xml_names, check_streams);
756
757                //err_pos[idx] = check_streams.error_mask;
[1668]758        }
759}
Note: See TracBrowser for help on using the repository browser.