source: docs/Working/re/ppopp-re.tex @ 3498

Last change on this file since 3498 was 3493, checked in by cameron, 6 years ago

Basic GPU description.

File size: 37.3 KB
Line 
1%-----------------------------------------------------------------------------
2%
3%               Template for sigplanconf LaTeX Class
4%
5% Name:         sigplanconf-template.tex
6%
7% Purpose:      A template for sigplanconf.cls, which is a LaTeX 2e class
8%               file for SIGPLAN conference proceedings.
9%
10% Guide:        Refer to "Author's Guide to the ACM SIGPLAN Class,"
11%               sigplanconf-guide.pdf
12%
13% Author:       Paul C. Anagnostopoulos
14%               Windfall Software
15%               978 371-2316
16%               paul@windfall.com
17%
18% Created:      15 February 2005
19%
20%-----------------------------------------------------------------------------
21
22
23\documentclass[preprint]{sigplanconf}
24
25% The following \documentclass options may be useful:
26
27% preprint      Remove this option only once the paper is in final form.
28% 10pt          To set in 10-point type instead of 9-point.
29% 11pt          To set in 11-point type instead of 9-point.
30% authoryear    To obtain author/year citation style instead of numeric.
31
32\usepackage{amsmath}
33\usepackage{pgfplots}
34
35\begin{document}
36
37\special{papersize=8.5in,11in}
38\setlength{\pdfpageheight}{\paperheight}
39\setlength{\pdfpagewidth}{\paperwidth}
40
41\conferenceinfo{PPoPP 2014}{February 15-19, 2014, Orlando, Florida, United States} 
42\copyrightyear{2013} 
43\copyrightdata{978-1-nnnn-nnnn-n/yy/mm} 
44\doi{nnnnnnn.nnnnnnn}
45
46% Uncomment one of the following two, if you are not going for the
47% traditional copyright transfer agreement.
48
49%\exclusivelicense                % ACM gets exclusive license to publish,
50                                  % you retain copyright
51
52%\permissiontopublish             % ACM gets nonexclusive license to publish
53                                  % (paid open-access papers,
54                                  % short abstracts)
55
56\titlebanner{}        % These are ignored unless
57\preprintfooter{Bitwise Data Parallel Grep}   % 'preprint' option specified.
58
59\title{Bitwise Data Parallelism in Regular Expression Matching}
60%\subtitle{Subtitle Text, if any}
61
62\authorinfo{Anonymous Authors}{Institutions}{emails}
63%\authorinfo{Robert D. Cameron \and Kenneth S. Herdy \and Dan Lin \and Meng Lin \and Ben Hull \and Thomas S. Shermer \and Arrvindh Shriraman}
64%          {Simon Fraser University}
65%           {\{cameron,ksherdy,lindanl,linmengl,bhull,shermer,ashriram\}@cs.sfu.ca}
66
67\maketitle
68
69\begin{abstract}
70\input{abstract}
71\end{abstract}
72\category{Theory of computation}{Formal languages and automata theory}{Regular languages}
73\category{Computer systems organization}{Parallel architectures}{Single instruction, multiple data}
74
75% general terms are not compulsory anymore,
76% you may leave them out
77%\terms
78%term1, term2
79
80\keywords
81regular expression matching, grep, parallel bit stream technology
82
83\section{Introduction}
84
85The use of regular expressions to search texts for occurrences
86of string patterns has a long history and
87remains a pervasive technique throughout computing applications today.
88% {\em a brief history}
89The origins of regular expression matching date back to automata theory
90developed by Kleene in the 1950s \cite{kleene1951}.
91Thompson \cite{thompson1968} is credited with the first construction to convert regular expressions
92to nondeterministic finite automata (NFA).
93Following Thompson's approach, a regular expression of length $m$ is first converted
94to an NFA with $O(m)$ nodes. It is then possible to search a text of length $n$ using the
95NFA in worst case $O(mn)$ time. Often, a more efficient choice
96is to convert the NFA into a DFA. A DFA has only a single active state at any time
97in the matching process and
98hence it is possible to search a text at of length $n$ in $O(n)$ time.
99However, it is well known that the conversion of an NFA to an equivalent DFA may result
100in state explosion. That is, the number of resultant DFA states may increase exponentially.
101In \cite{Baeza-yates_anew} a new approach to text searching was proposed based on bit-parallelism \cite{baeza1992new}.
102This technique takes advantage of the intrinsic parallelism of bitwise operations
103within a computer word. Given a $w$-bit word, the Shift-Or algorithm \cite{Baeza-yates_anew} algorithm uses the
104bit-parallel approach to
105simulate an NFA in $O(nm/w)$ worst-case time.
106
107A disadvantage of the bit-parallel Shift-Or pattern matching approach
108in comparison to simple string matching algorithms is an inability to skip input characters.
109For example, the Boyer-Moore family of algorithms \cite{boyer1977fast} skip input characters
110to achieve sublinear times in the average case. Backward Dawg Matching
111(BDM) string matching algorithms \cite{crochemore1994text} based on suffix automata are able to skip characters.
112The Backward Nondeterministic Dawg Matching (BNDM) pattern matching algorithm \cite{wu1992fast} 
113combines the bit-parallel advantages of Shift-Or and with the character skipping advantages of the BDM algorithm.
114The nrgrep pattern matching tool is built over the BNDM algorithm,
115and hence the name nrgrep \cite{navarro2000}.
116
117There has been considerable interest in using parallelization techniques
118to improve the performance of regular expression matching on parallel hardware
119such as multi-core processors (CPUs), graphics processing units (GPUs),
120field-programmable gate arrays (FPGAs), and specialized architectures such as
121the Cell Broadband Engine (Cell BE). % FPGA results (synthesis of patterns into logic circuits) vs. memory based approaches (STTs in memory)
122%CPU
123Scarpazza and Braudaway \cite{scarpazza2008fast} demonstrated that
124text processing algorithms that exhibit irregular memory access patterns
125can be efficiently executed on multicore hardware.
126In related work, Pasetto et al. presented a flexible tool that
127performs small-ruleset regular expression matching at a rate of
1282.88 Gbps per chip on Intel Xeon E5472 hardware \cite{pasetto2010}.
129Naghmouchi et al. demonstrated that the Aho-Corasick (AC)
130string matching algorithm \cite{aho1975} is well suited for parallel
131implementation on multi-core CPUs, GPUs and the Cell BE \cite{scarpazza2011top, naghmouchi2010}.
132On each hardware, both thread-level parallelism (additional cores) and data-level parallelism
133(wide SIMD units) are leveraged for performance.
134Salapura et. al., advocated the use of vector-style processing for regular expressions
135in business analytics applications and leveraged the SIMD hardware available
136on multi-core processors to acheive a speedup of better than 1.8 over a
137range of data sizes of interest \cite{salapura2012accelerating}.
138%Cell
139In \cite{scarpazza2008}, Scarpazza and Russell presented a SIMD tokenizer
140that delivered 1.00–1.78 Gbps on a single
141Cell BE chip and extended this approach for emulation on the Intel Larrabee
142instruction set \cite{scarpazza2009larrabee}.
143On the Cell BE, Scarpazza \cite{scarpazza2009cell} described a pattern matching
144implementation that delivered a throughput of 40
145Gbps for a small dictionary of approximately 100 patterns, and a throughput of 3.3-3.4
146Gbps for a larger dictionary of thousands of patterns. Iorio and van Lunteren \cite{iorio2008} 
147presented a string matching implementation for automata that achieves
1484 Gbps on the Cell BE.
149% GPU
150In more recent work, Tumeo et al. \cite{tumeo2010efficient} presented a chunk-based
151implementation of the AC algorithm for
152accelerating string matching on GPUs. Lin et al., proposed
153the Parallel Failureless Aho-Corasick (PFAC)
154algorithm to accelerate pattern matching on GPU hardware and
155achieved 143 Gbps throughput, 14.74 times faster
156than the AC algorithm performed on a four core
157multi-core processor using OpenMP \cite{lin2013accelerating}.
158
159Whereas the existing approaches to parallelization have been
160based on adapting traditional sequential algorithms to emergent
161parallel architectures, we introduce both a new algorithmic
162approach and its implementation on SIMD and GPU architectures.
163This approach relies on a bitwise data parallel view of text
164streams as well as a surprising use of addition to match
165runs of characters in a single step.  The closest previous
166work is that underlying bit-parallel XML parsing using 128-bit SSE2 SIMD
167technology together with a parallel scanning primitive also
168based on addition \cite{cameron2011parallel}.   
169However, in contrast to the deterministic, longest-match
170scanning associated with the ScanThru primitive of that
171work, we introduce here a new primitive MatchStar
172that can be used in full generality for nondeterministic
173regular expression matching.   We also introduce a long-stream
174addition technique involving a further application of MatchStar
175that enables us to scale the technique to $n$-bit addition
176in $\lceil\log_{64}{n}\rceil$ steps.   We ultimately apply this technique,
177for example, to perform
178synchronized 4096-bit addition on GPU warps of 64 threads.
179
180There is also a strong keyword match between the bit-parallel
181data streams used in our approach and the bit-parallelism
182used for NFA state transitions in the classical algorithms of
183Wu and Manber \cite{wu1992agrep}, Baez-Yates and Gonnet \cite{baeza1992new}
184and Navarro and Raffinot \cite{navarro1998bit}.
185However those algorithms use bit-parallelism in a fundamentally
186different way: representing all possible current NFA states
187as a bit vector and performing parallel transitions to a new
188set of states using table lookups and bitwise logic.    Whereas
189our approach can match multiple characters per step, bit-parallel
190NFA algorithms proceed through the input one byte at a time.
191Nevertheless, the agrep \cite{wu1992agrep} and
192nrgrep \cite{navarro2000} programs implemented using these techniques remain
193among the strongest competitors in regular expression matching
194performance, so we include them in our comparative evaluation.
195
196The remainder of this paper is organized as follows.
197Section \ref{sec:grep} briefly describes regular expression
198notation and the grep problem.
199Section \ref{sec:bitwise} presents our basic algorithm and MatchStar
200using a model of arbitrary-length bit-parallel data streams.
201Section \ref{sec:blockwise} discusses the block-by-block
202implementation of our techniques including the long stream
203addition techniques for 256-bit addition with AVX2 and
2044096-bit additions with GPGPU SIMT.
205Section \ref{sec:analysis} 
206Section \ref{sec:SSE2} 
207Section \ref{sec:AVX2} 
208Section \ref{sec:GPU} 
209Section \ref{sec:Concl} concludes the paper with a discussion of areas for future work.
210
211\section{Regular Expression Notation and Grep}\label{sec:grep}
212
213We follow common Posix notation for regular expressions.
214A regular expression specifies a set of strings through
215a pattern notation.   Individual characters normally
216stand for themselves, unless they are one of the
217special characters \verb:*+?[{\(|^$.: that serve as metacharacters
218of the notation system.  Thus the regular expression \verb:cat:
219is a pattern for the set consisting of the single 3-character
220string ``\verb:cat:''.   The special characters must be escaped
221with a backslash to prevent interpretation as metacharacter, thus
222\verb:\$: represents the dollar-sign and \verb:\\\\: represent
223the string consisting of two backslash characters.
224Character class bracket expressions are pattern elements
225that allow any character in a given class to be used in a particular
226context.  For example, \verb:[@#%]: is a regular expression
227that stands for any of the three given symbols.  Contiguous
228ranges of characters may be specified using hyphens;
229for example \verb:[0-9]: for digits and \verb:[A-Za-z0-9_]:
230for any alphanumeric character or underscore.  If the
231caret character immediately follows the opening bracket,
232the class is negated, thus \verb:[^0-9]: stands for
233any character except a digit.  The period metacharacter
234\verb:.: stands for the class of all characters.
235
236Consecutive pattern elements stand for strings formed by
237concatenation, thus \verb:[cd][ao][tg]: stands for the
238set of 8 three-letter strings ``\verb:cat:'' through ``\verb:dog:''.
239The alternation operator \verb:|: allows a pattern to be
240defined to have to alternative forms, thus \verb:cat|dog:
241matches either ``\verb:cat:'' or ``\verb:dog:''.  Concatenation
242takes precedence over alternation, but parenthesis may be
243used to change this, thus \verb:(ab|cd)[0-9]: stands for any
244digit following one of the two prefixes  ``\verb:ab:'' or ``\verb:cd:''.
245
246Repetition operators may be appended to a pattern to specify
247a variable number of occurrences of that pattern. 
248The Kleene \verb:*: specifies zero-or-more occurrences
249matching the previous pattern, while \verb:+: specifies one-or
250more occurrences.  Thus \verb:[a-z][a-z]*: and \verb:[a-z]+:
251both specify the same set: strings of at least one lower-case
252letter.  The postfix operator \verb:?: specifies an optional
253component, i.e., zero-or-one occurrence of strings matching
254the element.  Specific bounds may be given within braces:
255\verb:(ab){3}: specifies the string ``\verb:ababab:'',
256\verb:[0-9A-Fa-f]{2,4}: specifies strings of two, three
257or four hexadecimal digits, and \verb:[A-Z]{4,}: specifies
258strings of at least 4 consecutive capital letters.
259
260The grep program searches a file for lines containing matches
261to a regular expression using any of the above notations.
262In addition, the pattern elements \verb:^: and \verb:$:
263may be used to match respectively the beginning or the
264end of a line.  In line-based tools such as grep, \verb:.:
265matches any character except newlines; matches cannot extend
266over lines.
267Normally, grep prints all matching
268lines to its output.  However, grep programs typically
269allow a command line flag such as \verb:-c: to specify
270that only a count of matching lines be produced; we use
271this option in our experimental evaluation to focus
272our comparisons on the performance of the underlying matching
273algorithms.
274
275\section{Matching with Bit-Parallel Data Streams}\label{sec:bitwise}
276
277Whereas the traditional approaches to regular expression matching
278using NFAs, DFAs or backtracking all rely on a byte-at-a-time
279processing model, the approach  we introduce in this paper is based
280on quite a different concept:  a data-parallel approach to simultaneous
281processing of data stream elements.  Indeed, our most abstract model
282is that of unbounded data parallelism: processing all elements of
283the input data stream simultaneously.   In essence, data streams are viewed
284as (very large) integers.   The fundamental operations are bitwise
285logic, stream shifting and long-stream addition.
286
287Depending on the available parallel processing resources, an actual
288implementation may divide an input stream into blocks  and process
289the blocks sequentially.   Within each block  all elements of the
290input stream are processed together, relying the availability of
291bitwise logic and addition scaled to the block size.   On commodity
292Intel and AMD processors with 128-bit SIMD capabilities (SSE2),
293we typically process input streams 128 bytes at a time.   In this
294case, we rely on the Parabix tool chain \cite{lin2012parabix}
295to handle the details of compilation to block-by-block processing.
296For our GPGPU implementation, we have developed a long-stream
297addition technique that allows us to perform 4096-bit additions
298using 64 threads working in lock-step SIMT fashion.  Using scripts
299to modify the output of the Parabix tools, we effectively divide
300the input into blocks of 4K bytes processed in a fully data-parallel
301manner.
302
303\begin{figure}[tbh]
304\begin{center}
305\begin{tabular}{cr}\\
306input data  & \verb`a4534q--b29z---az---a4q--bca22z--`\\
307$B_7$ & \verb`.................................`\\
308$B_6$ & \verb`1....1..1..1...11...1.1..111..1..`\\
309$B_5$ & \verb`111111111111111111111111111111111`\\
310$B_4$ & \verb`.11111...111....1....11.....111..`\\
311$B_3$ & \verb`......11..11111.1111...11.....111`\\
312$B_2$ & \verb`.11.1.11....111..111.1.11......11`\\
313$B_1$ & \verb`...1....11.1....1........11.111..`\\
314$B_0$ & \verb`1.11.111..1.1111.1111.111.11...11`\\
315\verb:[a]: & \verb`1..............1....1......1.....`\\
316\verb:[z]: & \verb`...........1....1.............1..`\\
317\verb:[0-9]: & \verb`.1111....11..........1......11...`
318\end{tabular}
319
320\end{center}
321\caption{Basis and Character Class Streams}
322\label{fig:streams}
323\end{figure}
324
325A key concept in this streaming approach is the derivation of bit streams
326that are parallel to the input data stream, i.e., in one-to-one
327correspondence with the data element positions of the input
328streams.   Typically, the input stream is a byte stream comprising
329the 8-bit character code units of a particular encoding such
330as extended ASCII, ISO-8859-1 or UTF-8.   However, the method may also
331easily be used with wider code units such as the 16-bit code units of
332UTF-16.   In the case of a byte stream, the first step is to transpose
333the byte stream into eight parallel bit streams, such that bit stream
334$i$ comprises the $i^\text{th}$ bit of each byte.   These streams form
335a set of basis bit streams from which many other parallel bit
336streams can be calculated, such as character class bit
337streams such that each bit $j$ of the stream specifies
338whether character $j$ of the input stream is in the class
339or not.  Figure \ref{fig:streams} shows an example of an
340input byte stream in ASCII, the eight basis bit streams of the
341transposed representation, and the character class bit streams
342\verb:[a]:,
343\verb:[z]:, and
344\verb:[0-9]:
345that may be computed from the basis bit streams using bitwise logic.
346Zero bits are marked with periods ({\tt .}) so that the one bits stand out.
347Transposition and character class construction are straightforward
348using the Parabix tool chain \cite{lin2012parabix}.
349
350\begin{figure}[tbh]
351\begin{center}
352\begin{tabular}{cr}\\
353input data  & \verb`a4534q--b29z---az---a4q--bca22z--`\\
354$M_1$ & \verb`.1..............1....1......1....`\\
355$M_2$ & \verb`.11111..........1....11.....111..`\\
356$M_3$ & \verb`.................1.............1.`
357\end{tabular}
358
359\end{center}
360\caption{Marker Streams in Matching {\tt a[0-9]*z}}
361\label{fig:streams2}
362\end{figure}
363
364\paragraph*{Marker Streams.}  Now consider how bit-parallel data
365streams can be used in regular expression matching.   Consider
366the problem of searching the input stream of Figure \ref{fig:streams}
367to finding occurrence of strings matching
368the regular expression \verb:a[0-9]*z:.
369The matching process involves the concept of {\em marker streams}, that
370is streams that mark the positions of current matches during the
371overall process.  In this case there are three marker streams computed
372during the match process, namely,
373$M_1$ representing match positions after an initial \verb:a:
374character has been found, $M_2$ representing positions
375reachable from positions marked by $M_1$ by further matching zero or
376more digits (\verb:[0-9]*:) and finally $M_3$ the stream
377marking positions after a final \verb:z: has been found.
378Without describing the details of how these streams are computed
379for the time being, Figure \ref{fig:streams2} shows what each
380of these streams should be for our example matching problem.
381Note our convention that a marker stream contains a 1 bit
382at the next character position to be matched, that is,
383immediately past the last position that was matched.
384
385
386\paragraph*{MatchStar.}
387MatchStar takes a marker bitstream and a character class bitstream as input.  It returns all positions that can be reached by advancing the marker bitstream zero or more times through the character class bitstream.
388
389\begin{figure}[tbh]
390\begin{center}
391\begin{tabular}{cr}\\
392input data  & \verb`a4534q--b29z---az---a4q--bca22z--`\\
393$M_1$ & \verb`.1..............1....1......1....`\\
394$D = \text{\tt [0-9]}$ & \verb`.1111....11..........1......11...`\\
395$T_0 = M_1 \wedge D$ & \verb`.1...................1......1....`\\
396$T_1 = T_0 + D$ & \verb`.....1...11...........1.......1..`\\
397$T_2 = T_1 \oplus D$ & \verb`.11111...............11.....111..`\\
398$M_2 = T_2 \, | \, M_1$ & \verb`.11111..........1....11.....111..`
399\end{tabular}
400
401\end{center}
402\caption{$M_2 = \text{MatchStar}(M_1, D)$}
403\label{fig:matchstar}
404\end{figure}
405
406
407Figure \ref{fig:matchstar} illustrates the MatchStar method. In this figure,
408it is important to note that our bitstreams are shown in natural left-to-right order reflecting the
409conventional presentation of our character data input.   However, this reverses the normal
410order of presentation when considering bitstreams as numeric values.  The key point here is
411that when we perform bitstream addition, we will show bit movement from left-to-right.
412For example, $\verb:111.: + \verb:1...: = \verb:...1:$.
413
414The first row of the figure is the input data,
415the second and third rows are the input bitstreams: the initial marker position bitstream and the
416character class bitstream for digits derived from input data. 
417
418In the first operation ($T_0$), marker positions that cannot be advanced are temporarily removed from consideration by masking off marker positions that aren't character class positions using bitwise logic.  Next, the temporary marker bitstream is added to the character class bitstream. 
419The addition produces 1s in three types of positions.  There will be a 1 immediately following a block of character class positions that spanned one or more marker positions, at any character class positions that weren't affected by the addition (and are not part of the desired output), and at any marker position that wasn't the first in its block of character class positions.  Any character class positions that have a 0 in $T_1$ were affected by the addition and are part of the desired output.  These positions are obtained and the undesired 1 bits are removed by XORing with the character class stream. $T_2$ is now only missing marker positions that were removed in the first step as well as marker positions that were 1s in $T_1$.  The
420output marker stream is obtained by ORing $T_2$ with the initial marker stream.
421
422In general, given a marker stream $M$ and a character class stream $C$,
423the operation of MatchStar is defined by the following equation. 
424\[\text{MatchStar}(M, C) = (((M \wedge C) + C)  \oplus C) | M\]
425Given a set of initial marker positions, the result stream marks
426all possible positions that can be reached by 0 or more occurrences
427of characters in class $C$ from each position in $M$
428
429\paragraph*{Compilation Algorithm.}
430
431
432\section{Block-at-a-Time Processing}\label{sec:blockwise}
433
434The unbounded stream model of the previous section must of course
435be translated an implementation that proceeds block-at-a-time for
436realistic application.  In this, we primarily rely on the Pablo
437compiler of the Parabix toolchain \cite{lin2012parabix}.  Given input
438statements expressed as arbitrary-length bitstream equations, Pablo
439produces block-at-a-time C++ code that initializes and maintains all the necessary
440carry bits for each of the additions and shifts involved in the
441bitstream calculations.   
442
443In the present work, our principal contribution to the block-at-a-time
444model is the technique of long-stream addition described below.
445Otherwise, we were able to use Pablo directly in compiling our
446SSE2 and AVX2 implementations.   Our GPU implementation required
447some scripting to modify the output of the Pablo compiler for our
448purpose.
449
450\paragraph*{Long-Stream Addition.}  The maximum word size for
451addition on commodity processors is typically 64 bits.  In order
452to implement long-stream addition for block sizes of 256 or larger,
453a method for propagating carries through the individual stages of
45464-bit addition is required.  However, the normal technique of
455sequential addition using add-with-carry instructions, for example,
456is far from ideal.
457
458We have developed a technique using SIMD or SIMT methods for constant-time
459long-stream addition up to 4096 bits.   
460We assume the availability of the following SIMD/SIMT operations
461operating on vectors of $f$ 64-bit fields.
462\begin{itemize}
463\item \verb#simd<64>::add(X, Y)#: vertical SIMD addition of corresponding 64-bit fields
464in two vectors to produce a result vector of $f$ 64-bit fields.
465\item  \verb#simd<64>::eq(X, -1)# :  comparison of the 64-bit fields
466of \verb:x: each with the constant value -1 (all bits 1), producing
467an $f$-bit mask value,
468\item  \verb#hsimd<64>::mask(X)# : gathering the high bit of each 64-bit
469field into a single compressed $f$-bit mask value, and
470\item normal bitwise logic operations on $f$-bit masks, and
471\item  \verb#simd<64>::spread(x)# : distributing the bits of
472an $f$ bit mask, one bit each to the $f$ 64-bit fields of a vector.
473\end{itemize}
474
475Given these operations, our method for long stream addition of
476two $f \times 64$ bit values \verb:X: and \verb:Y: is the following.
477\begin{enumerate}
478\item Form the vector of 64-bit sums of \verb:x: and \verb:y:.
479\[\text{\tt R} = \text{\tt simd<64>::add(X, Y)} \]
480
481\item Extract the $f$-bit masks of \verb:X:, \verb:Y: and \verb:R:.
482\[\text{\tt x} = \text{\tt hsimd<64>::mask(X)} \]
483\[\text{\tt y} = \text{\tt hsimd<64>::mask(Y)} \]
484\[\text{\tt r} = \text{\tt hsimd<64>::mask(R)} \]
485
486\item Compute an $f$-bit mask of carries generated for each of the
48764-bit additions of \verb:X: and \verb:Y:.
488\[\text{\tt c} = (\text{\tt x} \wedge \text{\tt y}) \vee ((\text{\tt x} \vee \text{\tt y}) \wedge \neg \text{\tt r})\]
489
490\item Compute an $f$-bit mask of all fields of {\tt R} that will overflow with
491an incoming carry bit.  This is the {\em bubble mask}.
492\[\text{\tt b} = \text{\tt simd<64>::eq(R, -1)}\]
493
494\item Determine an $f$-bit mask identifying the fields of {\tt R} that need to be
495incremented to produce the final sum.  Here we find a new application of
496MatchStar!
497\[\text{\tt i} = \text{\tt MatchStar(c*2, b)}\]
498
499This is the key step.  The mask {\tt c} of outgoing carries must be
500shifted one position ({\tt c*2}) so that each outgoing carry bit becomes associated
501with the next digit.  At the incoming position, the carry will
502increment the 64-bit digit.   However, if this digit is all ones (as
503signalled by the corresponding bit of bubble mask {\tt b}, then the addition
504will generate another carry.  In fact, if there is a sequence of
505digits that are all ones, then the carry must bubble through
506each of them.   This is just MatchStar!
507
508\item Compute the final result {\tt Z}.
509\[\text{\tt Z} = \text{\tt simd<64>::add(R, simd<64>::spread(i))}\]
510
511\end{enumerate}
512\begin{figure}
513\begin{center}
514\begin{tabular}{c||r|r|r|r|r|r|r|r||}\cline{2-9}
515{\tt X} & {\tt 19} & {\tt 31} & {\tt BA} & {\tt 4C} & {\tt 3D} & {\tt 45} & {\tt 21} & {\tt F1} \\ \cline{2-9}
516{\tt Y} & {\tt 22} & {\tt 12} & {\tt 45} & {\tt B3} & {\tt E2} & {\tt 16} & {\tt 17} & {\tt 36} \\ \cline{2-9}
517{\tt R} & {\tt 3B} & {\tt 43} & {\tt FF} & {\tt FF} & {\tt 1F} & {\tt 5B} & {\tt 38} & {\tt 27} \\ \cline{2-9}
518{\tt x} & {\tt 0} & {\tt 0} & {\tt 1} & {\tt 0} & {\tt 0} & {\tt 0} & {\tt 0} & {\tt 1} \\ \cline{2-9}
519{\tt y} & {\tt 0} & {\tt 0} & {\tt 0} & {\tt 1} & {\tt 1} & {\tt 0} & {\tt 0} & {\tt 0} \\ \cline{2-9}
520{\tt r} & {\tt 0} & {\tt 0} & {\tt 1} & {\tt 1} & {\tt 0} & {\tt 0} & {\tt 0} & {\tt 0} \\ \cline{2-9}
521{\tt c} & {\tt 0} & {\tt 0} & {\tt 0} & {\tt 0} & {\tt 1} & {\tt 0} & {\tt 0} & {\tt 1} \\ \cline{2-9}
522{\tt c*2} & {\tt 0} & {\tt 0} & {\tt 0} & {\tt 1} & {\tt 0} & {\tt 0} & {\tt 1} & {\tt 0} \\ \cline{2-9}
523{\tt b} & {\tt 0} & {\tt 0} & {\tt 1} & {\tt 1} & {\tt 0} & {\tt 0} & {\tt 0} & {\tt 0} \\ \cline{2-9}
524{\tt i} & {\tt 0} & {\tt 1} & {\tt 1} & {\tt 1} & {\tt 0} & {\tt 0} & {\tt 1} & {\tt 0} \\ \cline{2-9}
525{\tt Z} & {\tt 3B} & {\tt 44} & {\tt 0} & {\tt 0} & {\tt 1F} & {\tt 5B} & {\tt 39} & {\tt 27} \\ \cline{2-9}
526\end{tabular}
527\end{center}
528\caption{Long Stream Addition}\label{fig:longadd}
529\end{figure}
530
531Figure \ref{fig:longadd} illustrates the process.  In the figure,
532we illustrate the process with 8-bit fields rather than 64-bit fields
533and show all field values in hexadecimal notation.  Note that
534two of the individual 8-bit additions produce carries, while two
535others produce {\tt FF} values that generate bubble bits.  The
536net result is that four of the original 8-bit sums must be
537incremented to produce the long stream result.
538
539A slight extension to the process produces a long-stream full adder
540that can be used in chained addition.  In this case, the
541adder must take an additional carry-in bit
542{\tt p} and produce a carry-out bit {\tt q}.
543This may be accomplished by incorporating {\tt p}
544in calculating the increment mask in the low bit position,
545and then extracting the carry-out {\tt q} from the high bit position.
546\[\text{\tt i} = \text{\tt MatchStar(c*2+p, b)}\]
547\[\text{\tt q} = \text{\tt i >> f}\]
548
549As described subsequently, we use a two-level long-stream addition technique
550in both our AVX2 and GPU implementations.  In principle, one can extend
551the technique to additional levels.  Using 64-bit adders throughout,
552$\lceil\log_{64}{n}\rceil$ steps are needed for $n$-bit addition.
553A three-level scheme could coordinate
55464 groups each performing 4096-bit long additions in a two-level structure.
555However, whether there are reasonable architectures that can support fine-grained
556SIMT style coordinate at this level is an open question.
557
558Using the methods outlined, it is quite conceivable that instruction
559set extensions to support long-stream addition could be added for
560future SIMD and GPU processors.   Given the fundamental nature
561of addition as a primitive and its novel application to regular
562expression matching as shown herein, it seems reasonable to expect
563such instructions to become available.
564\raggedbottom
565\section{Analytical Comparison with DFA and NFA Implementations}\label{sec:analysis}
566
567\begin{enumerate}
568\item Operations
569\item Memory behaviour per input byte: note tables of DFA/NFA.
570
571Bille and Throup \em{Faster regular expression matching}\cite{bille2009faster}
572
573\end{enumerate}
574
575
576
577\section{Commodity SIMD Implementation and Experimental Evaluation}\label{sec:SSE2}
578
579
580\subsection{Implementation Notes}
581\subsection{Evaluation Methodology}
582\subsection{Comparison}
583\begin{figure}
584\begin{center}
585\begin{tikzpicture}
586\begin{axis}[
587xtick=data,
588ylabel=Cycles per Byte,
589xticklabels={@,Date,Email,URIorEmail,xquote},
590tick label style={font=\tiny},
591enlargelimits=0.15,
592legend style={at={(0.5,-0.15)},
593anchor=north,legend columns=-1},
594ymax=8,
595ybar,
596bar width=7pt,
597]
598\addplot
599file {data/cycles1.dat};
600\addplot
601file {data/cycles2.dat};
602\addplot
603file {data/cycles3.dat};
604 
605\legend{Bitstreams,NRGrep,Grep,Annot}
606\end{axis}
607\end{tikzpicture}
608\end{center}
609\caption{Cycles per Byte}
610\end{figure}
611 
612\begin{figure}
613\begin{center}
614\begin{tikzpicture}
615\begin{axis}[
616xtick=data,
617ylabel=Instructions per Byte,
618xticklabels={@,Date,Email,URIorEmail,xquote},
619tick label style={font=\tiny},
620enlargelimits=0.15,
621legend style={at={(0.5,-0.15)},
622anchor=north,legend columns=-1},
623ymax=16,
624ybar,
625bar width=7pt,
626]
627\addplot
628file {data/instructions1.dat};
629\addplot
630file {data/instructions2.dat};
631\addplot
632file {data/instructions3.dat};
633 
634\legend{Bitstreams,NRGrep,Grep,Annot}
635\end{axis}
636\end{tikzpicture}
637\end{center}
638\caption{Instructions per Byte}
639\end{figure}
640
641\begin{figure}
642\begin{center}
643\begin{tikzpicture}
644\begin{axis}[
645xtick=data,
646ylabel=Instructions per Cycle,
647xticklabels={@,Date,Email,URIorEmail,xquote},
648tick label style={font=\tiny},
649enlargelimits=0.15,
650legend style={at={(0.5,-0.15)},
651anchor=north,legend columns=-1},
652ybar,
653bar width=7pt,
654]
655\addplot
656file {data/ipc1.dat};
657\addplot
658file {data/ipc2.dat};
659\addplot
660file {data/ipc3.dat};
661
662\legend{Bitstreams,NRGrep,Grep,Annot}
663\end{axis}
664\end{tikzpicture}
665\end{center}
666\caption{Instructions per Cycle}
667\end{figure}
668
669\begin{figure}
670\begin{center}
671\begin{tikzpicture}
672\begin{axis}[
673xtick=data,
674ylabel=Branch Misses per Byte,
675xticklabels={@,Date,Email,URIorEmail,xquote},
676tick label style={font=\tiny},
677enlargelimits=0.15,
678legend style={at={(0.5,-0.15)},
679anchor=north,legend columns=-1},
680ymax=0.03,
681ybar,
682bar width=7pt,
683]
684\addplot
685file {data/branch-misses1.dat};
686\addplot
687file {data/branch-misses2.dat};
688\addplot
689file {data/branch-misses3.dat};
690
691\legend{Bitstreams,NRGrep,Grep,Annot}
692\end{axis}
693\end{tikzpicture}
694\end{center}
695\caption{Branch Misses per Byte}
696\end{figure}
697
698
699
700\section{SIMD Scalability}\label{sec:AVX2}
701
702
703
704
705\begin{figure}
706\begin{center}
707\begin{tikzpicture}
708\begin{axis}[
709xtick=data,
710ylabel=Cycles per Byte,
711xticklabels={@,Date,Email,URIorEmail,xquote},
712tick label style={font=\tiny},
713enlargelimits=0.15,
714legend style={at={(0.5,-0.15)},
715anchor=north,legend columns=-1},
716ybar,
717bar width=7pt,
718]
719\addplot
720file {data/ssecycles.dat};
721\addplot
722file {data/avxcycles.dat};
723
724\legend{SSE2,AVX2,Annot}
725\end{axis}
726\end{tikzpicture}
727\end{center}
728\caption{Cycles per Byte}
729\end{figure}
730
731\begin{figure}
732\begin{center}
733\begin{tikzpicture}
734\begin{axis}[
735xtick=data,
736ylabel=Instructions per Byte,
737xticklabels={@,Date,Email,URIorEmail,xquote},
738tick label style={font=\tiny},
739enlargelimits=0.15,
740legend style={at={(0.5,-0.15)},
741anchor=north,legend columns=-1},
742ybar,
743bar width=7pt,
744]
745\addplot
746file {data/sseinstructions.dat};
747\addplot
748file {data/avxinstructions.dat};
749
750\legend{SSE2,AVX2,Annot}
751\end{axis}
752\end{tikzpicture}
753\end{center}
754\caption{Instructions per Byte}
755\end{figure}
756
757\begin{figure}
758\begin{center}
759\begin{tikzpicture}
760\begin{axis}[
761xtick=data,
762ylabel=Instructions per Cycle,
763xticklabels={@,Date,Email,URIorEmail,xquote},
764tick label style={font=\tiny},
765enlargelimits=0.15,
766legend style={at={(0.5,-0.15)},
767anchor=north,legend columns=-1},
768ybar,
769bar width=7pt,
770]
771\addplot
772file {data/sseipc.dat};
773\addplot
774file {data/avxipc.dat};
775
776
777\legend{SSE2,AVX2,Annot}
778\end{axis}
779\end{tikzpicture}
780\end{center}
781\caption{Instructions per Cycle}
782\end{figure}
783
784\begin{figure}
785\begin{center}
786\begin{tikzpicture}
787\begin{axis}[
788xtick=data,
789ylabel=Branch Misses per Byte,
790xticklabels={@,Date,Email,URIorEmail,xquote},
791tick label style={font=\tiny},
792enlargelimits=0.15,
793legend style={at={(0.5,-0.15)},
794anchor=north,legend columns=-1},
795ybar,
796bar width=7pt,
797]
798\addplot
799file {data/ssebranch-misses.dat};
800\addplot
801file {data/avxbranch-misses.dat};
802
803\legend{SSE2,AVX2,Annot}
804\end{axis}
805\end{tikzpicture}
806\end{center}
807\caption{Branch Misses per Byte}
808\end{figure}
809
810
811
812
813\subsection{AVX Stream Addition}
814 \begin{figure*}[tbh]
815\begin{center}
816\begin{verbatim}
817void add_ci_co(bitblock_t x, bitblock_t y, carry_t carry_in, carry_t & carry_out, bitblock_t & sum) {
818  bitblock_t all_ones = simd256<1>::constant<1>();
819  bitblock_t gen = simd_and(x, y);
820  bitblock_t prop = simd_xor(x, y);
821  bitblock_t partial_sum = simd256<64>::add(x, y);
822  bitblock_t carry = simd_or(gen, simd_andc(prop, partial_sum));
823  bitblock_t bubble = simd256<64>::eq(partial_sum, all_ones);
824  uint64_t carry_mask = hsimd256<64>::signmask(carry) * 2 + convert(carry_in);
825  uint64_t bubble_mask = hsimd256<64>::signmask(bubble);
826  uint64_t carry_scan_thru_bubbles = (carry_mask + bubble_mask) &~ bubble_mask;
827  uint64_t increments = carry_scan_thru_bubbles | (carry_scan_thru_bubbles - carry_mask);
828  carry_out = convert(increments >> 4);
829  uint64_t spread = 0x0000200040008001 * increments & 0x0001000100010001;
830  sum = simd256<64>::add(partial_sum, _mm256_cvtepu16_epi64(avx_select_lo128(convert(spread))));
831}
832
833\end{verbatim}
834
835\end{center}
836\caption{AVX2 256-bit Addition}
837\label{fig:AVX2add}
838\end{figure*}
839
840
841\section{GPU Implementation}\label{sec:GPU}
842
843To further assess the scalability of our regular expression matching
844using bit-parallel data streams, we implemented a GPGPU version
845in OpenCL.   We arranged for 64 work groups each having 64
846threads.  Input files are divided in data parallel fashion among
847the 64 work groups.  Each work group carries out the regular
848expression matching operations 4096 bytes at a time using SIMT
849processing.  Figure \ref{fig:GPUadd} shows our implementation
850of long-stream addition on the GPU.  Each thread maintains
851its own carry and bubble values and performs synchronized
852updates with the other threads using a six-step parallel-prefix
853style process.
854
855
856
857\begin{figure*}[tbh]
858\begin{center}
859\begin{verbatim}
860inline BitBlock adc(int idx, BitBlock a, BitBlock b, __local BitBlock *carry, _
861                    _local BitBlock *bubble, BitBlock *group_carry, const int carryno){
862        BitBlock carry_mask;
863        BitBlock bubble_mask;
864
865        BitBlock partial_sum = a+b;
866        BitBlock gen = a&b;
867        BitBlock prop = a^b;
868        carry[idx] = ((gen | (prop & ~partial_sum))&CARRY_BIT_MASK)>>(WORK_GROUP_SIZE-1-idx);
869        bubble[idx] = (partial_sum + 1)? 0:(((BitBlock)1)<<idx);
870       
871        barrier(CLK_LOCAL_MEM_FENCE);
872        for(int offset=WORK_GROUP_SIZE/2; offset>0; offset=offset>>1){
873                carry[idx] = carry[idx]|carry[idx^offset];
874                bubble[idx] = bubble[idx]|bubble[idx^offset];
875                barrier(CLK_LOCAL_MEM_FENCE);
876        }
877       
878        carry_mask = (carry[0]<<1)|group_carry[carryno];
879        bubble_mask = bubble[0];
880       
881        BitBlock s = (carry_mask + bubble_mask) & ~bubble_mask;
882        BitBlock inc = s | (s-carry_mask);
883        BitBlock rslt = partial_sum + ((inc>>idx)&0x1);
884        group_carry[carryno] = (carry[0]|(bubble_mask & inc))>>63;
885        return rslt;
886}
887\end{verbatim}
888
889Figure \ref{fig:SSE-AVX-GPU} compares the performance of
890our SSE2, AVX and GPU implementations.   
891
892\end{center}
893\caption{OpenCL 4096-bit Addition}
894\label{fig:GPUadd}
895\end{figure*}
896
897
898
899\begin{figure}
900\begin{center}
901\begin{tikzpicture}
902\begin{axis}[
903xtick=data,
904ylabel=Running Time (ms per byte),
905xticklabels={@,Date,Email,URIorEmail,xquote},
906tick label style={font=\tiny},
907enlargelimits=0.15,
908legend style={at={(0.5,-0.15)},
909anchor=north,legend columns=-1},
910ybar,
911bar width=7pt,
912]
913\addplot
914file {data/ssetime.dat};
915\addplot
916file {data/avxtime.dat};
917\addplot
918file {data/gputime.dat};
919
920\legend{SSE2,AVX2,GPU,Annot}
921\end{axis}
922\end{tikzpicture}
923\end{center}
924\caption{Running Time}\label{fig:SSE-AVX-GPU}
925\end{figure}
926
927
928
929
930\section{Miscellaneous}
931\subsection{Skipping}
932\input{re-Unicode}
933
934\input{conclusion}
935
936
937
938%\appendix
939%\section{Appendix Title}
940
941%This is the text of the appendix, if you need one.
942
943\acks
944
945This research was supported by grants from the Natural Sciences and Engineering Research Council of Canada and
946MITACS, Inc.
947
948% We recommend abbrvnat bibliography style.
949
950\bibliographystyle{abbrvnat}
951
952% The bibliography should be embedded for final submission.
953 
954\bibliography{reference}
955
956%\begin{thebibliography}{}
957%\softraggedright
958
959%\bibitem[Smith et~al.(2009)Smith, Jones]{smith02}
960%P. Q. Smith, and X. Y. Jones. ...reference text...
961%
962%\end{thebibliography}
963
964
965\end{document}
966
967
Note: See TracBrowser for help on using the repository browser.