-
Notifications
You must be signed in to change notification settings - Fork 0
/
nfa.cu
443 lines (399 loc) · 11.5 KB
/
nfa.cu
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
277
278
279
280
281
282
283
284
285
286
287
288
289
290
291
292
293
294
295
296
297
298
299
300
301
302
303
304
305
306
307
308
309
310
311
312
313
314
315
316
317
318
319
320
321
322
323
324
325
326
327
328
329
330
331
332
333
334
335
336
337
338
339
340
341
342
343
344
345
346
347
348
349
350
351
352
353
354
355
356
357
358
359
360
361
362
363
364
365
366
367
368
369
370
371
372
373
374
375
376
377
378
379
380
381
382
383
384
385
386
387
388
389
390
391
392
393
394
395
396
397
398
399
400
401
402
403
404
405
406
407
408
409
410
411
412
413
414
415
416
417
418
419
420
421
422
423
424
425
426
427
428
429
430
431
432
433
434
435
436
437
438
439
440
441
442
443
/*
*
* Mini regex-module inspired by Rob Pike's regex code described in:
*
* http://www.cs.princeton.edu/courses/archive/spr09/cos333/beautiful.html
*
*
*
* Supports:
* ---------
* '.' Dot, matches any character
* '^' Start anchor, matches beginning of string
* '$' End anchor, matches end of string
* '*' Asterisk, match zero or more (greedy)
* '+' Plus, match one or more (greedy)
* '?' Question, match zero or one (non-greedy)
* '[abc]' Character class, match if one of {'a', 'b', 'c'}
* '[^abc]' Inverted class, match if NOT one of {'a', 'b', 'c'} -- NOTE: feature is currently broken!
* '[a-zA-Z]' Character ranges, the character set of the ranges { a-z | A-Z }
* '\s' Whitespace, \t \f \r \n \v and spaces
* '\S' Non-whitespace
* '\w' Alphanumeric, [a-zA-Z0-9_]
* '\W' Non-alphanumeric
* '\d' Digits, [0-9]
* '\D' Non-digits
*
*
*/
#include "nfa.h"
#include <stdio.h>
/* Definitions: */
#define MAX_REGEXP_OBJECTS 30 /* Max number of regex symbols in expression. */
#define MAX_CHAR_CLASS_LEN 40 /* Max length of character-class buffer in. */
enum { UNUSED, DOT, BEGIN, END, QUESTIONMARK, STAR, PLUS, CHAR, CHAR_CLASS, INV_CHAR_CLASS, DIGIT, NOT_DIGIT, ALPHA, NOT_ALPHA, WHITESPACE, NOT_WHITESPACE, /* BRANCH */ };
typedef struct regex_t
{
unsigned char type; /* CHAR, STAR, etc. */
union
{
unsigned char ch; /* the character itself */
unsigned char* ccl; /* OR a pointer to characters in class */
};
} regex_t;
/* Private function declarations: */
__device__ static int matchpattern(regex_t* pattern, const char* text);
__device__ static int matchcharclass(char c, const char* str);
__device__ static int matchstar(regex_t p, regex_t* pattern, const char* text);
__device__ static int matchplus(regex_t p, regex_t* pattern, const char* text);
__device__ static int matchone(regex_t p, char c);
__device__ static int matchdigit(char c);
__device__ static int matchalpha(char c);
__device__ static int matchwhitespace(char c);
__device__ static int matchmetachar(char c, const char* str);
__device__ static int matchrange(char c, const char* str);
__device__ static int ismetachar(char c);
extern __device__ int re_matchp(re_t pattern, const char* text) {
if (pattern != 0) {
if (pattern[0].type == BEGIN)
return ((matchpattern(&pattern[1], text)) ? 1 : 0);
else return (matchpattern(pattern, text));
/* return (text[0] != '\0' && text[0] != '\n'); */
}
return 0;
}
re_t re_compile(const char* pattern)
{
/* The sizes of the two static arrays below substantiates the static RAM usage of this module.
MAX_REGEXP_OBJECTS is the max number of symbols in the expression.
MAX_CHAR_CLASS_LEN determines the size of buffer for chars in all char-classes in the expression. */
regex_t* re_compiled;
cudaMallocManaged(&re_compiled, MAX_REGEXP_OBJECTS);
static unsigned char ccl_buf[MAX_CHAR_CLASS_LEN];
int ccl_bufidx = 1;
char c; /* current char in pattern */
int i = 0; /* index into pattern */
int j = 0; /* index into re_compiled */
while (pattern[i] != '\0' && (j+1 < MAX_REGEXP_OBJECTS))
{
c = pattern[i];
switch (c)
{
/* Meta-characters: */
case '^': { re_compiled[j].type = BEGIN; } break;
case '$': { re_compiled[j].type = END; } break;
case '.': { re_compiled[j].type = DOT; } break;
case '*': { re_compiled[j].type = STAR; } break;
case '+': { re_compiled[j].type = PLUS; } break;
case '?': { re_compiled[j].type = QUESTIONMARK; } break;
/* case '|': { re_compiled[j].type = BRANCH; } break; <-- not working properly */
/* Escaped character-classes (\s \w ...): */
case '\\':
{
if (pattern[i+1] != '\0')
{
/* Skip the escape-char '\\' */
i += 1;
/* ... and check the next */
switch (pattern[i])
{
/* Meta-character: */
case 'd': { re_compiled[j].type = DIGIT; } break;
case 'D': { re_compiled[j].type = NOT_DIGIT; } break;
case 'w': { re_compiled[j].type = ALPHA; } break;
case 'W': { re_compiled[j].type = NOT_ALPHA; } break;
case 's': { re_compiled[j].type = WHITESPACE; } break;
case 'S': { re_compiled[j].type = NOT_WHITESPACE; } break;
/* Escaped character, e.g. '.' or '$' */
default:
{
re_compiled[j].type = CHAR;
re_compiled[j].ch = pattern[i];
} break;
}
}
/* '\\' as last char in pattern -> invalid regular expression. */
/*
else
{
re_compiled[j].type = CHAR;
re_compiled[j].ch = pattern[i];
}
*/
} break;
/* Character class: */
case '[':
{
/* Remember where the char-buffer starts. */
int buf_begin = ccl_bufidx;
/* Look-ahead to determine if negated */
if (pattern[i+1] == '^')
{
re_compiled[j].type = INV_CHAR_CLASS;
i += 1; /* Increment i to avoid including '^' in the char-buffer */
}
else
{
re_compiled[j].type = CHAR_CLASS;
}
/* Copy characters inside [..] to buffer */
while ( (pattern[++i] != ']')
&& (pattern[i] != '\0')) /* Missing ] */
{
if (pattern[i] == '\\')
{
if (ccl_bufidx >= MAX_CHAR_CLASS_LEN - 1)
{
//fputs("exceeded internal buffer!\n", stderr);
return 0;
}
ccl_buf[ccl_bufidx++] = pattern[i++];
}
else if (ccl_bufidx >= MAX_CHAR_CLASS_LEN)
{
//fputs("exceeded internal buffer!\n", stderr);
return 0;
}
ccl_buf[ccl_bufidx++] = pattern[i];
}
if (ccl_bufidx >= MAX_CHAR_CLASS_LEN)
{
/* Catches cases such as [00000000000000000000000000000000000000][ */
//fputs("exceeded internal buffer!\n", stderr);
return 0;
}
/* Null-terminate string end */
ccl_buf[ccl_bufidx++] = 0;
re_compiled[j].ccl = &ccl_buf[buf_begin];
} break;
/* Other characters: */
default:
{
re_compiled[j].type = CHAR;
re_compiled[j].ch = c;
} break;
}
i += 1;
j += 1;
}
/* 'UNUSED' is a sentinel used to indicate end-of-pattern */
re_compiled[j].type = UNUSED;
return re_compiled;
}
void re_print(regex_t* pattern)
{
const char* types[] = { "UNUSED", "DOT", "BEGIN", "END", "QUESTIONMARK", "STAR", "PLUS", "CHAR", "CHAR_CLASS", "INV_CHAR_CLASS", "DIGIT", "NOT_DIGIT", "ALPHA", "NOT_ALPHA", "WHITESPACE", "NOT_WHITESPACE", "BRANCH" };
int i;
for (i = 0; i < MAX_REGEXP_OBJECTS; ++i)
{
if (pattern[i].type == UNUSED)
{
break;
}
printf("type: %s", types[pattern[i].type]);
if (pattern[i].type == CHAR_CLASS || pattern[i].type == INV_CHAR_CLASS)
{
printf(" [");
int j;
char c;
for (j = 0; j < MAX_CHAR_CLASS_LEN; ++j)
{
c = pattern[i].ccl[j];
if ((c == '\0') || (c == ']'))
{
break;
}
printf("%c", c);
}
printf("]");
}
else if (pattern[i].type == CHAR)
{
printf(" '%c'", pattern[i].ch);
}
printf("\n");
}
}
/* Private functions: */
__device__ static int matchdigit(char c)
{
return ((c >= '0') && (c <= '9'));
}
__device__ static int matchalpha(char c)
{
return ((c >= 'a') && (c <= 'z')) || ((c >= 'A') && (c <= 'Z'));
}
__device__ static int matchwhitespace(char c)
{
return ((c == ' ') || (c == '\t') || (c == '\n') || (c == '\r') || (c == '\f') || (c == '\v'));
}
__device__ static int matchalphanum(char c)
{
return ((c == '_') || matchalpha(c) || matchdigit(c));
}
__device__ static int matchrange(char c, const char* str)
{
return ((c != '-') && (str[0] != '\0') && (str[0] != '-') &&
(str[1] == '-') && (str[1] != '\0') &&
(str[2] != '\0') && ((c >= str[0]) && (c <= str[2])));
}
__device__ static int ismetachar(char c)
{
return ((c == 's') || (c == 'S') || (c == 'w') || (c == 'W') || (c == 'd') || (c == 'D'));
}
__device__ static int matchmetachar(char c, const char* str)
{
switch (str[0])
{
case 'd': return matchdigit(c);
case 'D': return !matchdigit(c);
case 'w': return matchalphanum(c);
case 'W': return !matchalphanum(c);
case 's': return matchwhitespace(c);
case 'S': return !matchwhitespace(c);
default: return (c == str[0]);
}
}
__device__ static int matchcharclass(char c, const char* str)
{
do
{
if (matchrange(c, str))
{
return 1;
}
else if (str[0] == '\\')
{
/* Escape-char: increment str-ptr and match on next char */
str += 1;
if (matchmetachar(c, str))
{
return 1;
}
else if ((c == str[0]) && !ismetachar(c))
{
return 1;
}
}
else if (c == str[0])
{
if (c == '-')
{
return ((str[-1] == '\0') || (str[1] == '\0'));
}
else
{
return 1;
}
}
}
while (*str++ != '\0');
return 0;
}
__device__ static int matchone(regex_t p, char c)
{
switch (p.type)
{
case DOT: return 1;
case CHAR_CLASS: return matchcharclass(c, (const char*)p.ccl);
case INV_CHAR_CLASS: return !matchcharclass(c, (const char*)p.ccl);
case DIGIT: return matchdigit(c);
case NOT_DIGIT: return !matchdigit(c);
case ALPHA: return matchalphanum(c);
case NOT_ALPHA: return !matchalphanum(c);
case WHITESPACE: return matchwhitespace(c);
case NOT_WHITESPACE: return !matchwhitespace(c);
default: return (p.ch == c);
}
}
__device__ static int matchstar(regex_t p, regex_t* pattern, const char* text)
{
do
{
if (matchpattern(pattern, text))
return 1;
}
while ((text[0] != '\0' && text[0] != '\n') && matchone(p, *text++));
return 0;
}
__device__ static int matchplus(regex_t p, regex_t* pattern, const char* text)
{
while ((text[0] != '\0' && text[0] != '\n') && matchone(p, *text++))
{
if (matchpattern(pattern, text))
return 1;
}
return 0;
}
__device__ static int matchquestion(regex_t p, regex_t* pattern, const char* text)
{
if (p.type == UNUSED)
return 1;
if (matchpattern(pattern, text))
return 1;
if (*text && matchone(p, *text++))
return matchpattern(pattern, text);
return 0;
}
#if 0
/* Recursive matching */
static int matchpattern(regex_t* pattern, const char* text)
{
if ((pattern[0].type == UNUSED) || (pattern[1].type == QUESTIONMARK))
{
return matchquestion(pattern[1], &pattern[2], text);
}
else if (pattern[1].type == STAR)
{
return matchstar(pattern[0], &pattern[2], text);
}
else if (pattern[1].type == PLUS)
{
return matchplus(pattern[0], &pattern[2], text);
}
else if ((pattern[0].type == END) && pattern[1].type == UNUSED)
{
return text[0] == '\0';
}
else if ((text[0] != '\0') && matchone(pattern[0], text[0]))
{
return matchpattern(&pattern[1], text+1);
}
else
{
return 0;
}
}
#else
/* Iterative matching */
__device__ static int matchpattern(regex_t* pattern, const char* text)
{
do
{
if ((pattern[0].type == UNUSED) || (pattern[1].type == QUESTIONMARK))
{
return matchquestion(pattern[0], &pattern[2], text);
}
else if (pattern[1].type == STAR)
{
return matchstar(pattern[0], &pattern[2], text);
}
else if (pattern[1].type == PLUS)
{
return matchplus(pattern[0], &pattern[2], text);
}
else if ((pattern[0].type == END) && pattern[1].type == UNUSED)
{
return (text[0] == '\0' || text[0] == '\n');
}
/* Branching is not working properly
else if (pattern[1].type == BRANCH)
{
return (matchpattern(pattern, text) || matchpattern(&pattern[2], text));
}
*/
}
while ((text[0] != '\0' && text[0] != '\n') && matchone(*pattern++, *text++));
return 0;
}
#endif