[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] | 78 | inline 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] | 105 | inline 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 | |
---|
| 132 | inline 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] | 142 | inline 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] | 146 | inline 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] | 150 | inline 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 | |
---|
| 158 | inline void CarryInit(BitBlock * carry, int count){ |
---|
| 159 | for (int j=0; j < count; j++) |
---|
| 160 | carry[j] = 0; |
---|
| 161 | } |
---|
| 162 | |
---|
[1673] | 163 | inline 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] | 171 | inline 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 |
---|
| 715 | void 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 | } |
---|