Hardik Soni

Thierry Turletti

Walid Dabbous

P4Bricks: Enabling multiprocessing using Linker-based network data plane architecture

Packet-level programming languages such as P4 usually require to describe all packet processing functionalities for a given programmable network device within a single program. However, this approach monopolizes the device by a single large network application program, which prevents possible addition of new functionalities by other independently written network applications.

We propose P4Bricks, a system which aims to deploy and execute multiple independently developed and compiled P4 programs on the same reconfigurable hardware device. P4Bricks is based on a Linker component that merges the programmable parsers/deparsers and restructures the logical pipeline of P4 programs by refactoring, decomposing and scheduling the pipelines' tables. It merges P4 programs according to packet processing semantics (parallel or sequential) specified by the network operator and runs the programs on the stages of the same hardware pipeline, thereby enabling multiprocessing. This paper presents the initial design of our system with an ongoing implementation and studies P4 language's fundamental constructs facilitating merging of independently written programs. P4Bricks 2.1 P4 background

Introduction

P4 [START_REF]The P4 Language Consortium. P4 16 language specification version 1[END_REF] is a high-level language for programming protocol-independent packet processors. It allows reconfiguring packet processing behavior of already deployed data plane hardware devices, introducing new protocols and their processing to the devices, hence decoupling packet processing hardware from software. This provides high degree of flexibility for programming new network applications and packet processing functionalities using reconfigurable hardware like RMT [START_REF] Bosshart | Forwarding metamorphosis: Fast programmable match-action processing in hardware for sdn[END_REF] and Intel Flexpipe TM . With P4, a network operator can execute one program at a time on the target reconfigurable network device. However, as the number of features or network applications to be supported by the device grows, P4 programs increase in complexity and size. The development and maintenance of such monolithic P4 programs containing all the possible features is error prone and needs huge time and effort as program complexity grows. On the other hand, this approach does not allow to easily compose independently written modules in a single P4 program. Network devices can have different packet processing requirements according to their role and location in the network. Deploying one large program for a small subset of applications results in inefficient utilization of the reconfigurable device resources.

P5 [START_REF] Abhashkumar | P5: Policy-driven optimization of p4 pipeline[END_REF] optimizes resource utilization by leveraging policy intents specifying which features are required to remove excess applications and packet processing functionalities. However, with P5 there is still a large monolithic program to be configured based on P4Bricks 1. Introduction the policy intents. ClickP4 [START_REF] Zhou | Clickp4: Towards modular programming of p4[END_REF] proposes a smart modular approach in which separate modules can be developed within the ClickP4 programming framework and then manually integrated into ClickP4 configuration files to create a large program. However, programmers are required to know the code of the ClickP4 library modules to integrate a new module into the ClickP4 framework as source code modifications may be required for modules on the already developed code base. Basically, P5 allows removing extra modules and features from already composed small P4 programs, whereas ClickP4 gives choice to select from a list of modules. Most importantly, with both P5 and ClickP4, packet processing functionalities on a device can not be easily composed using independently developed and compiled P4 programs.

Hyper4 [START_REF] Hancock | Hyper4: Using p4 to virtualize the programmable data plane[END_REF], HyperV [START_REF] Zhang | Hyperv: A high performance hypervisor for virtualization of the programmable data plane[END_REF] and MPVisor [START_REF] Zhang | Mpvisor: A modular programmable data plane hypervisor[END_REF] propose virtualization of programmable data plane in order to deploy and run independently developed multiple P4 programs on the same network device at the same time. In these approaches, a general purpose P4 program working as a hypervisor for programmable data plane is developed, which can be configured to achieve functionally equivalent packet processing behavior of multiple P4 programs hosted by it. However, virtualization requires minimum 6-7× and 3-5× more match-action stages for every P4 program compared to its native execution for Hyper4 and HyperV, respectively. Also, such approaches show significant performance degradation for bandwidth and delay, thereby nullifying the benefit of high performance reconfigurable hardware.

Meanwhile, executing efficiently multiple P4 programs at a time on a same target device is highly desirable. We believe that a network operator should be able to easily add any features on its target device with programs potentially developed by different providers.

We present the design and architecture of P4Bricks, our under development system that aims to deploy and execute multiple independently developed and compiled P4 programs on the same reconfigurable device. P4Bricks comprises two components, called Linker and Runtime. The Linker component merges the programmable parsers and deparsers and restructures the logical pipeline of P4 programs by refactoring, decomposing and scheduling their match-action tables (MATs). The Runtime component translates the dynamic table updates from the control planes of different applications into the tables of the merged pipeline. P4Bricks merges and executes MATs of multiple compiled P4 programs on the stages of the same hardware pipeline, thereby enabling multiprocessing. The idea is to provide a seamless execution environment for P4 programs in a multiprogram environment without any changes required in its control interface and MATs definitions. With P4Bricks network operators can specify the packet processing policy on the target device in terms of compiled P4 programs and composition operators using a simple command line interface. This report presents the initial design of our system with an ongoing implementation and studies P4 language's fundamental constructs facilitating merging of independently written programs. It is organized as follows. Section 2 provides an overview of the P4Bricks system. Section 3 describes Linker, which composes compiled P4 programs using the only knowledge of MATs definitions. Then, Section 4 describes the Runtime module that interacts with the Linker component and the control plane of P4 programs to manage flow entries in the MATs of the programs.

System Overview

In this section, we provide a brief overview of the P4 language and we then introduce our system, describing merging of programmable blocks at link time and their management at runtime.

P4 background

P4 [START_REF]The P4 Language Consortium. P4 16 language specification version 1[END_REF] is a high-level language, based on programmable blocks, used to define protocolindependent packet processing by programming data plane of reconfigurable target devices. A P4-compatible target manufacturer provides P4 architecture defining the programmable blocks and describing several hardware related information. Essentially, the P4 architecture of a target provides programmable components and declares interface to program them and exposes specific and already implemented constructs(e.g., checksum units and algorithms) that can be used and manipulated through APIs. The programmable parser, deparser and logical match-action pipeline are the main blocks used to program the data plane of packet processing targets. In addition of providing data plane programmability, P4 generates the APIs for control plane of the target device to communicate with the data plane. The APIs allow to manage the state of data plane objects from the control plane.

P4 provides the following fundamental abstractions to program reconfigurable target devices.

• Header types -to describe the format (ordered sequence of fields and their size) of headers within a packet.

• Parser -to describe all the possible sequences of headers within received packets, mechanism to identify the header sequences, headers and the values of the fields.

• Actions -are already implemented in the target, hence their behaviors are fixed. They are used manipulate header fields and metadata, also may take data as input from control plane at runtime.

• Tables -allows P4 programmer to define match keys and associate them with actions. A Match key can have multiple header and metadata fields. A Match key provides complex match and decision capabilities.

• Match-action unit -implements a table for execution at runtime. It performs table lookup using the match key to search associated actions and data, and executes the actions.

• Control Flow -to describe packet processing program, which includes invoking sequence of Match-action units. Packet reassembly can also be defined using control flow.

• extern -are architecture and hardware specific objects and implementations, which can be manipulated but not programmed. Because, their processing behaviors are implemented in the target.

• Intrinsic metadata -are architecture and target specific entities associated with each packet(e.g., interfaces)

• User-defined metadata -are data structures and variables defined in a program to maintain the program specific per packet state during packet processing.

The P4 language is constituted of four different sub-languages used for different purposes.

1. The core language -to describe types, variables, scoping, declarations, statements etc.

2. A sub-language for describing parsers -having specific constructs to describe packet parsing.

3. A sub-language for describing processing -using match-action units and to define traditional imperative control flow for packet processing 4. A sub-language for describing architecture -to define and declare types of programmable blocks for the target, architecture specific data and functions.

Next, we briefly discuss sub-languages for parsers and packet processing along with some of their fundamental constructs used to program Parsers, Deparser and logical match-action pipeline.

Programmable Parser and Packet Parsing in P4

All the packet processing targets must identify packets' protocol headers and their fields to determine how packets can be processed. Identifying and extracting protocol headers and their fields within packets is called packet parsing. Packet parsing is a complex and challenging as the packet size, protocol header types and their payload vary across the networks and packets within the same network. Header type defines the format of a protocol header within a packet, programmers can describe it by specifying sequence of fields and their sizes. Every header type has a length field and an identifier field indicating length of the current header and encapsulated protocol type, respectively, to facilitate extraction of header fields and subsequent protocols headers. The first protocol header fields are extracted based on the network type from the start of packet bit-stream. Parser identifies the next protocol header to extract from the values of current protocol header's fields using a given parse graph, which captures all possible sequences of protocol headers using a DAG. Fixed parsers can parse packets according to the parser and protocol types defined at design time of the target. Where as programmable parsers allow to modify the parse graph and protocol headers types at any time according to new packet processing requirements.

P4 allows to define programmable parser blocks for P4- ) and method calls (to process parsed fields using extract and invoke other parser blocks defined within the same program). P4-compatible target implements the parser FSM in their programmable parser unit [START_REF] Gibb | Design principles for packet parsers[END_REF], where FSM is converted into state transition table and loaded into the unit's memory blocks. Packets parsed into header using defined parser blocks instances are further processed using control blocks.

Control Block

P4 provides a specific language constructs to program control block describing processing of the parsed headers and other metadata in the parser block. Actions and tables are among the fundamental packet processing abstractions declared within the body of control block. P4 uses the same control block but with specialized signature as the programming interface deparsing.

Programming match-action units

P4 language construct tables are used to describe packet processing in match-action units of targets. In this thesis, we use Match-

Deparser Control Block

Packets are reassembled using the processed header instances by programming the deparser control block. Deparsing is described using a control block having a mandatory specific parameter of type packet out. P4 provides a language construct, emit, to reassemble the packet. Emit takes header instance as an argument and if the header instance is valid, it is appended to the packet. If the header is not valid, no operation is performed.

The P4Bricks System

P4Bricks enables network operators to deploy and execute multiple independently developed and compiled P4 programs on the same reconfigurable target device. P4Bricks allows network operator to define packet processing policy of the device using P4 programs and composition operators. It merges P4 programs according to packet processing semantics (parallel or sequential) specified by the network operator and runs the programs on the stages of the same hardware pipeline, thereby enabling multiprocessing. P4Bricks comprises two components, called Linker and Runtime. Linker takes as input the data plane configuration files generated by P4 compiler from source files of independently written P4 programs along with composition operators to apply on the programs, see Figure 1. It merges the programmable parser, pipeline and deparser blocks defined in the compiled configuration files of P4 programs according to packet processing policy described using parallel and sequential composition operators. It restructures the logical pipeline of P4 programs by refactoring, decomposing and scheduling their match-action tables (MATs). Linker does not assume knowledge of flow entries in MATs of P4 programs while merging the pipelines at link time. P4Bricks considers the packet header instances (extracted or emitted) along with Userdefined and Intrinsic metadata associated with each packet as data plane resources.

We consider the packet header instances and intrinsic metadata as shared data plane resources accessed for packet processing by the logical pipelines of the different P4 programs. However, we do not share user-defined metadata of one program with the other P4 programs because they are used by programmers to store the program specific state during packet processing. Linker replaces program specific IDs of header types and instances given by the compiler with the mapped ones in UIDTs before merging pipelines according to the composition operators. This unifies different references (IDs and names) for equivalent header types and instances used in P4 programs, thereby identifying sharing of data plane resources. As Intrinsic metadata represents data structures and variable related to architecture and target, they are uniformly referenced across independently developed P4 programs. User-defined metadata of every program is always given a unique ID using program's name to restrict other programs from accessing it.

P4 allows to program packet deparsing using a control block having at least one parameter of a specialized type, packet out. A deparser control block allows to generate the emit sequence of header instances to reassemble the packet, as defined by the programmer. If the instance is valid, emitting a header instance will append the instance to the outgoing packet. Similar to parser, two deparser blocks can be merged if the P4 compiler generates a DAG of emit instances, which encodes all possible emits functions after appending a header instance. However, as the P4 compiler generates a list to append header instances, the topological order (providing relative location of appending the header instance) can not be identified. The topological order is essential requirement to merge instances of different header types at the same level of the network stack. Let us consider that deparser of program P A emits header instances in the order of Ethernet, IPv4, ICMP and TCP, whereas deparser of program P B provides emit sequence of Ethernet, VLAN, IPv6 and UDP. Because of the semantics of emit calls that append the header instance only if it is valid, and the strict order imposed by the sequences of these calls, it is not possible to identify the correct merged emit order of header instances.

In particular, it is not possible to deduct that the Ethernet header can be followed by either VLAN, IPv6 or IPv4. In our current P4Bricks implementation, we use the P4Bricks 3. Linker merged parse graph to identify the topological order among the header instances to be compatible with current P4 specifications. We note that if a future version of P4 makes use of DAGs to represent deparser control block, this could allow merging of deparsers without dependence to the parser block.

In P4, MATs are defined using 1) match keys composed of header fields and runtime state data of the program, and 2) actions to be taken based on matching on the keys. The packet processing control flow in the compiled configuration file of a P4 program is commonly represented as a DAG, called control flow graph (CFG), with each node representing packet processing using a MAT. The edges in CFG represent control dependency among MATs capturing packet processing order in the program. For each pipeline of a program, Linker decomposes the MATs CFG by adding resources as nodes, splitting each MAT node into match and action control nodes, and adding dependencies between control nodes and resource nodes according to resources accessed. Linker generates read-write operation schedule graph (OSG) for each resource from the decomposed CFG to capture all possible access orders and types (read or write) of operations executed on the resource due to packet processing control flow in the pipeline. Linker merges packet processing CFGs of all the P4 programs and the OSGs generated from them for each resource according to composition operators. Then, Linker refactors the MATs, regenerates the CFG and maps the refactored MATs to physical pipeline stages while respecting the merged read-write OSG for each resource, the MAT control flow of all the P4 programs and available physical match memory type and capacity in the stages. We introduce two concepts to facilitate this restructuring : 1) Vertically decomposing the MATs into sub MATs and 2) Performing out-of-order write operations in OSG of any resource. These techniques allow mapping of a sub MAT on available physical match memory type (e.g., exact, ternary or Longest Prefix Match) in the physical pipeline stage, even if the complete MAT can not be scheduled to the stage due to control dependency. Apart from creating new MATs and CFGs, Linker produces the mappings between new MATs mapped to physical pipeline stages and the MATs of all the P4 programs. It also prepares the MATs mappings and update schedules of the decomposed MATs that will be used by the Runtime component, as shown in Figure 1.

Runtime executes in the control plane of the target device and acts as a proxy to the control plane of P4 programs in order to manage their MATs defined in configuration files. It uses UIDTs and MAT mappings generated during linking to translate MATs update from the control plane of programs to the tables mapped to physical pipeline stages by Linker. Runtime is responsible for maintaining referential integrity across the sub MATs of a decomposed MAT to provide consistent MAT update. For every decomposed MAT with its sub MATs mapped to different stages of physical pipeline, Runtime updates the entries of the sub MATs according to the schedule generated by Linker. Moreover, it regulates the flow updates from the control plane of all the P4 programs to enforce flow space isolation dictated by the network operator.

Linker

In this section, we describe the static linking process of compiled configuration files of multiple independently written P4 programs. Essentially, merging of two parsers requires creating a union of their Directed Acyclic Graphs (DAGs) encoded as FSMs. However, as programs are independently implemented, they may not use the same identifiers and names. So, it is necessary to identify the equivalent header types, instances and parse states defined in the parser blocks. Also, explicit state transitions to accept and reject states and implicit transitions to these states resulting from error handling function call verify are required to merge for semantically correct merging of two parser FSMs. In the following we define the notion of equivalence of header types, parse states and header instances between two programs and explain how to find the equivalence relationship. Then, we describe our method to merge the states and create union of DAGs and thereby FSMs.

Equivalence of Header Types

A header type is defined as an ordered sequence of fields with bit-widths. Two fixedlength header types are equivalent if their ordered sequences of bit-widths are the same. Regarding variable length types, the length of the variable bit-width fields must depend on the same fixed width field in the header types and their maximum lengths must be P4Bricks 3.1 Merging Parsers and Deparsers identical. The length indicator field is uniquely identified by its bit-width and its offset from the start of the header. Using these definitions, we create a UID for each header type as an ordered sequence of bit-widths corresponding to its fields. In case of a variable length type, a UID is created considering its maximum possible length and the identifiers of the length indicator field. Header types UIDT maintains the mapping between UIDs and program specific identifiers for all the header types in all the programs.

Equivalence of Parse States and Header Instances

Parse states extract the bit streams into instances of header types defined in the program. So, the equivalence of parse states and header instances are correlated. The instances extracted from equivalent parse states are mapped to each other in the instance UIDT.

A parse state of a P4 program's parser is equivalent to a parse state of another program's parser if they satisfy the following conditions:

C1: both states extract bits from the same location in the bit stream of the packet. So, the current bit index points to the same location in the bit stream of the packet when the parsers visit the states.

C2: both states advance the current bit index by the same number of bits and extract the equivalent header types.

C3: if both states have select expression, then the lookup fields used in the expressions should be equivalent 1 .

These conditions cover the scenario of lookahead construct also, where the states do not advance the current bit index but read the same set of bits to identify next transition states or store the value. Let us take an example of merging two parse graphs of P4 programs to process Data center and Enterprise 2 network traffic, shown in Figures 3a and3b, respectively. The parse graph of Data center has two VLAN states extracting double tagged headers in two VLAN instances a and b, whereas the parse graph of the Enterprise network has a single VLAN state extracting header into the single instance named x. As shown in Figures 3c and3d, when merging the two parse graphs, V LAN x should be equivalent to V LAN a but not to V LAN b . Indeed, only V LAN x and V LAN a states extract the bits from the same current bit in the topological order to satisfy condition C1. Algorithm 1 describes the steps to identify equivalent parse states and header instances. First, we select the set of nodes (i.e., states) having 0 in-degree at each iteration for each parser graph in order to traverse the parse graphs in topological order (lines 2-3). Then, we find equivalence states by verifying C2 and C3 conditions for all possible pairs of states from the two sets(line 6). After that, we remove the 0 in-degree nodes with their outgoing edges and continue iterating till all the nodes are visited.

Traffic Isolation and Error Handling

Parsers of two programs may transit to different final states (accept or reject) while parsing the same packet's bit stream. In case of error detected by the verify statement 1 We assume all keyset expressions in select to be known values at compile time. 2 The parse graphs are inspired from Figure 3 in [START_REF] Gibb | Design principles for packet parsers[END_REF]. We add an intrinsic metadata field, called indicator, with a bit width equal to the number of P4 programs to be merged. Each bit of this field indicates validity of the packet header instances for a program. In every state of the merged parse graph, we add an assignment statement masking the indicator field with a binary code. The binary code specifies per program validity of the packet if the parser transits to the state. Moreover, we appropriately set the bits of binary code, while merging parsers' accepttransition and reject-transition states to indicate the validity of packet header instances for the programs according to the final states in their parse graphs.

We prepare a binary code for each merged state while merging equivalence parse states of parsers. First, we assign the binary code according to the equivalence relationship among the parsers states. For instance, in the example of Figure 3c, V LAN b and IP V 6 will mask the indicator field with 0b01 and 0b10 binary codes. Next, when an accepttransition state from one parser (U DP of Enterprise, Figure 3b) is required to merge with equivalent intermediate state (U DP of Data Center, Figure 3a) from another, we set the bit in the binary code corresponding to the former parser (Enterprise) in all successive states in the merged parser. Hence, V XLAN will mask the indicator field with binary code 0b11, even if it does not have any equivalent state in Enterprise parser.

The merged parse graph transits to final reject state only when all the parsers transit to their logical reject state, otherwise merged parser transits to final accept state.

Linker replaces the verify statements in compiled P4 configuration file to P4Bricks data plane architecture specific implementation with the following signature. We emphasize here that P4 programmers use the verify statement with the signature provided by the P4 language. The Linker component in P4Bricks is responsible for translating the verify statements of each program by invoking the calls with appropriate values of the indicator field and program id arguments.

Merging of Parse Graphs

We begin with merging the select transitions (i.e., edges in parse graphs) of two equivalent states by taking union of their keysets to create a single state. Regarding transition state of each select case, we find their equivalent state from the mappings and recursively merge them as described in Algorithm 2 . If the two states have the same keyset, the corresponding transition states must be equivalent and we merge them too. We note that for a same keyset, two parsers can not transit to two different states. For instance, the value 0x0800 of EtherType can not be used to transit to IPv4 state by one parser and to IPv6 state by the other. Allowing such ambiguous transitions creates non-deterministic FSM, resulting in a scenario where a packet can be parsed in different ways, which creates an ambiguity during packet processing.

In the case of sequential processing, merging the parsers of two programs is not sufficient to find the equivalent header instances between them. Let us consider an example of chaining of encapsulation-decapsulation network functions, where the first program pushes new header instances and the second parses the pushed header instances to process traffic in the network. Executing them on the same target with sequential composition requires identifying the equivalence relationship between header instances. For this purpose, we map the topological order of instances parsed from the merged parse graph to the sequence of emitted header instances in the deparser control block of the first program. Algorithm 3 describes detailed steps for it. In algorithm 3, we iteratively map 0 in-degree instances from the merged parser to the emitted instance from the deparser of the first program in sequence and removing instances from both. At each iteration, we search for an unmapped instance of merged parser having equivalent header type to the emitted unmapped instance. If there is any such instance pair, we create an equivalence mapping between them.

Apart from select, extract and verify, parse states may have assignment, variable and constant declarations statements. Variables and constants declared in any program are user-defined metadata and are not shared with other programs. We concatenate the lists of assignment, variable and constant declaration statements used in equivalent states along with translated verify statements, even though they may perform redundant Remove hi E from the list operations (e.g., verifying checksum or fields' value). Our current design for Linker fails to merge parsers of P4 programs and stops linking, if any of the equivalent parse states has assignment and method call statements modifying the shared data plane resources.

Deparser

P4 programmers can use emit function calls in a specialized control block defined for packet deparsing. If the header instance specified in a function argument is valid, it is appended to the packet else the function call does not perform any operation. To merge deparsers, we use merged parse graph to identify topological order of header instances. We note that the use of merged parse graph is not a semantically correct approach for finding the topological order among header instances in the deparsers of programs. Indeed, a parser may use a completely different network protocol stack to extract and process header instances than the one used in the deparser to reassembly the packet. More precisely, P4 programmer may (1) define parser to extract instances of some of the defined header types thereby decapsulating packet data from the header instances, (2) process the packet and finally (3) emit instances of disjoint header types to encapsulate the packet data before forwarding it to a interface. In such scenario, use of merged parse graph does not provide topological order among header instances to be emitted before sending out packets. Hence, we restrict merging to P4 programs, which use the same network protocol stack for parsing and deparsing of packets and do not perform encapsulations. This allows the merging of P4 programs to be compatible with semantics of deparser control block described in current P4-16 language specification version 1.0.0 and still find topological order of header instances emitted by deparsers.

In this work, we generate emit sequence of header instances by performing topological sort on the merged parse graph. We note that, we can perform sequential merging of parsers of an encapsulating P4 program followed by a decapsulating P4 program, because DAG in parser of decapsulating program is indirectly providing topological order of encapsulated header instance. However, in case of merging deparsers of multiple P4Bricks 3.2 Logical Pipelines Restructuring decapsulating programs, the topological order of header instances being emitted can not be identified with the current P4 specifications, thereby restricting the use of our proposal.

Logical Pipelines Restructuring

First, we replace all the occurrences of program specific identifiers of header types and instances with their UIDs using the mappings in the UIDTs generated while merging the parser and deparser blocks. Then, we decompose CFG stored in each P4 program's configuration file generated by P4 compiler. We create an OSG for each data plane resource from the decomposed CFG of each P4 program to capture all possible sequences of operations performed on the resource by the program. We describe decomposition of CFG and construction of OSG in Section 3.2.1. For each data plane resource, we apply composition operators on its OSGs generated from multiple P4 programs and create a single OSG for the resource, as explained in Section 3.2.2. This OSG captures all possible sequences of operations performed on the resource under given composition of P4 programs. Next, we merge decomposed CFGs of P4 programs according to composition operator by adding required control dependencies across the nodes of CFGs (Section 3.2.3). Also, we add dependencies from indicator metadata field to the control node without incoming edge in decomposed CFG of each P4 program. Such a control node specifies entry point of packet processing control flow of a given P4 program. Using the decomposed-and-merged CFGs and each data plane resource's OSG, we refactor original MATs of P4 programs and map them on physical pipeline stages. In Section 3.2.4, we describe the mechanism to map and refactor MATs on physical pipeline stages. To facilitate refactoring of MATs and mapping them on physical pipeline stages while respecting dependencies in merged OSGs and decomposed CFGs, we introduce 1) Vertically decomposing the MATs into sub MATs and 2) Performing out-of-order write operations in OSG of any resource in Sections 3.2.5 and 3.2.6, respectively.

Decomposing CFGs and Constructing Operation Schedules

First we decompose CFGs and capture schedules of the operations performed on each data plane resources. Each CFG node represents packet processing by a MAT involving match and action phases. We split each CFG node and reinterpret the packet processing phases as match and action control nodes operating on shared data plane resources, as shown in Figure 4a.

Match control nodes can only perform read operations on resources representing key fields. Action control nodes can perform read and write operations on resources, because all the actions declared within a table may take resources as input and modify them.

We add stage dependency from match to action control node of the same table. We decompose the CFG shown in 4b using this MAT representation. Figure 4c shows an example where several control nodes in the pipeline access a single resource R1. The decomposed CFG captures all control dependencies of the P4 program among control nodes and their operations on the data plane resources. CFGs generated by the P4 compiler are DAGs, but decomposed CFG may not be DAG. Because, we added every data plane resource as a graph node and a resource can be accessed by different control We create an operation schedule graph (called OSG) for each resource from the decomposed CFG to capture all possible sequences of operations that could be performed on the resource at runtime. As the CFG generated by P4 compiler for a P4 program is a DAG, so is the OSG of each resource derived from the CFG. Because, we merely split each CFG node into match and action control nodes, add stage dependency between them but do not add any packet processing control edge in CFG for decomposition. Within an OSG, nodes can be of two types (node-R, node-W ), depending on the type of operation (i.e., read or write) performed on the resource by either a match or an action control node. Figure 4d shows the OSG for resource R1 created from the graph of Figure 4c. Even if the decomposed CFG of a program has cycles, all the access to resource nodes in a cycle follows a directed path in OSGs of the resources. For each data plane resource, we create one OSG per P4 program (derived from its CFG) and we merge all the OSGs by applying the composition operators specified by the network operator.

Applying Composition Operators on Resources

Multiple P4 programs can process traffic in parallel provided that they are processing disjoint data plane resources (or traffic flows). When processing disjoint data plane resources, there can only be one program with write operations on a given data plane resource and all the read operations on this resource from other programs must complete before any possible write operation.

Figure 5a and5b show an example of OSGs for a resource derived from CFGs of two programs X and Y. All the read operations of the graph from Y must be scheduled before any write operations of the graph from X. To apply parallel composition for P4Bricks

3.2 Logical Pipelines Restructuring M1 X -R A2 X -W A5 X -W A4 X -R A3 X -W (a) OSG from X.p4 M1 Y -R A2 Y -R M3 Y -R (b) OSG from Y.p4 M1 Z -R A2 Z -W A4 Z -R A5 Z -W A3 Z -W (c) OSG from Z.p4 M1 X -R A2 X -W A5 X -W A4 X -R A3 X -W M1 Y -R A2 Y -R M3 Y -R (d) Parallel (X Y) M1 Z -R A2 Z -W A4 Z -R A5 Z -W A3 Z -W M1 X -R A2 X -W A5 X -W A4 X -R A3 X -W (e) Sequential (X Z)
Figure 5: Two Different Merging of Resource's OSGs shared traffic flows, we add dependencies from all possible last reads of schedule graph from Y to all possible first writes of schedule graph from X, as shown in Figure 5d. In case of programs operating on disjoint traffic flows, we do not add any node dependency across the OSGs and consider merged graph as disconnected acyclic graph. In this case, Runtime enforces isolation of flows of programs in their MATs.

In case of sequential composition, once the first program completes all its operations on a data plane resource then and only then the next program is allowed to access the resource. We take an example of merging resource's OSGs from two programs X and Z, shown in Figures 5a and5c, according to X after Z sequential composition. To apply the sequential composition's constraints, we add dependencies from all possible last operations of X to all possible first operations of Z as shown in Figure 5e. Similarly, we can merge resource's OSGs using any number of P4 programs by adding appropriate node dependencies among them, according to the specified operators and order. We use the merged OSGs of all the data plane resources, while restructuring the logical pipelines and mapping the MATs to physical pipeline stages.

Merging Decomposed CFGs

We merge decomposed CFGs by creating union of their nodes (i.e., data plane resource and control nodes) and edges. Every decomposed CFG has unique control nodes and control edges compared to other programs' decomposed CFGs. Because P4 programs may access shared data plane resources and perform operation on them, graph union connects decomposed CFGs of the programs through common resource nodes. If programs are accessing disjoint set of data plane resources, union of decomposed CFGs creates disconnected graph. The disconnected graph implies that merged program can have multiple control flows and MATs of each CFG can be scheduled and mapped on physical pipeline stages without any dependency from other CFGs' MATs. We add dependencies between indicator metadata field, created while merging parsers, and the match control node representing entry point of packet processing control flow in decomposed CFG of each P4 program. It allows to enforce required packet processing isolation In case of parallel composition, we do not need to add any dependencies across the control nodes, as all the programs should process the assigned traffic flows. In case of sequential composition, some actions from the first program must be executed before the next program can start processing, even if there is no dependencies between operation schedules of any shared data plane resource from the composition constraint. For example, drop actions from a firewall program must be scheduled before any possible action for the next program. Therefore we add control node dependencies from such specific action control nodes of the first program to all the control nodes without incoming control edge of next program.

The merging of decomposed CFGs captures all possible control flow paths of each P4 program's CFG under specified composition. Similar to decomposed CFG, decomposedand-merged CFGs can also have cycles. Every access sequence on every resource node resulting from a cycle should follow a corresponding directed path in the resource's merged OSG. However, under parallel composition, there may not exist a directed path for access sequences of all the resources in their respective merged OSGs. Such a scenario emerges when program P 1 writes on a set of resources read by program P 2 and P 1 reads a disjoint set of resources written by P 2. Figures 6a and6b shows decomposed CFGs of single table programs P1 and P2 both accessing resource R1 and R2. The cycle, R1-M1-A1-R2-M2-A2-R1, results in a sequence of operations (M1-R, A2-W) on resource R1 and (A1-W, M2-R) on resource R2 in the decomposed-and-merged CFG(Figure 6c). In the merged OSG of resource R1, a corresponding path can be found for the sequence of operations on it. However, for resource R2 no corresponding path can be found in its merged OSG (Figure 6e). If we consider the other direction R2-M2-A2-R1-M1-A1-R2 in the cycle, no corresponding path can be found in merged OSG of resource R1 instead of R2. So, in any case either of the resources is operated out of its scheduled order represented in OSG of the resource. To perform required operation sequence resulting from a cycle on a resource, we perform out-of-order write operation with respect to the resource's merged OSG not having directed path according to the operation sequence. We describe the procedure to schedule out-of-order write operation in OSGs in Section 3.2.6. Next, we discuss usage of merged OSGs and decomposed CFGs of the P4 programs to refactor MATs defined in the programs and map them on physical pipeline stages. Physical pipeline of a hardware target can be modeled as a DAG of stages, as shown in Figure 7, with each stage having different memory types with finite capacities and match capabilities, as described in [START_REF] Jose | Compiling packet programs to reconfigurable switches[END_REF]. The DAG provides the execution order of stages and possible parallelism among them. Moreover, decomposed-and-merged CFGs and merged resources' OSGs provide the scheduling order constraints on 1) match and action control nodes and 2) operations performed by these nodes on the resources. We refactor MATs and orderly map them to the physical pipeline stages by recomposing match and actions control nodes of the decomposed-and-merged CFG while respecting the scheduling order and the match stage constraints, as outlined in Algorithm 4.

In OSG of a resource, a node-W (specifying write operation on the resource) without any incoming edge is considered as schedule-ready. Schedule-ready nodes satisfy all scheduling constraints and are considered as available to map to pipeline stages. Also, a node-R (read operation) without any incoming edge along with its exclusive descendants with read operation type are considered schedule-ready; e.g., 5c and5d, respectively, are schedule-ready nodes. In a decomposed CFG, a control node without any incoming control edge is considered as control-ready. If a control node has all of its read operations on resources in ready state, it is considered as complete read-ready. A control node without any incoming edge is considered as schedule-ready having no control dependency as well as all of its operations on resources are schedule-ready. Essentially, a complete read-ready and control-ready node is schedule-ready. We select the physical pipeline stages according to their execution order (topological order) to map the refactored MATs while respecting the constraints of stages (line 2). To efficiently utilize the match stage's memory, we relax memory capacity constraints originating from the width of MATs by vertically decomposing them. This allows allocating the match stage memory at the granularity of individual key field instead of all the key fields required to match the control node.

A2 Z -W , M 1 Z -R and M 1 Y -R, M 1 X -R, A2 Y -R, M 3 Y -R in Figures
Memory allocation may map either a subset or all of the key fields of a match control node, thereby creating sub matches of the logical MAT. We allocate memory of the current physical match stages by selecting schedule-ready node-R pertaining to match control nodes from OSGs for all resources while respecting memory match type and capacity constraints of physical stages (line 5). This approach may map the logical match of a MAT defined in a program to multiple nonparallel match stages. We focus on realizing match stage memory allocation at granularity of match operation on individual data plane resources (i.e., individual key field) rather than presenting any specific algorithm to select match operations to map on memory stages. Thereby, we omit the description of SelectionOperations function used in line 5 in Algorithm 4. Once read-ready operations of resources are mapped to the match stage, we remove the nodes and edges from the schedule graphs by making all their predecessors point to their respective successors (lines 9-11). For all the sub matches and match control node not having its action node in complete read-ready state, we create sub MATs as described in section 3.2.5 and allocate the actions to the action stage of the physical pipeline (lines 15-18). If any control-ready and complete read-ready match node is mapped to the stage, we schedule its dependent action node provided that it is at least complete read-ready. Also, if any write operation from the action node is not schedule-ready (line 20), we perform out-of-order write to map the MAT by recomposing match and action control nodes as detailed in 3.2.6.

Decomposing logical MATs into sub MATs

For every match control node, we recompose a MAT using schedule-ready read operations allocated by a selection algorithm to map on multiple nonparallel stages. If all the read operations from a match control node are not mapped to the same stage by used selection algorithm, recomposition creates a sub MAT and vertically decomposes the logical MAT corresponding to the match control node. We create a new data plane resource of user-defined metadata type for every decomposition of a logical MAT. We term this resource as the foreign key field associated with the recomposed sub MAT and its logical MAT. In this sub MAT, we add a new action setting the foreign key field to the given row ID on successful match at runtime. We consider sub MAT as secondary table having an action setting the foreign key field. The bit width of the foreign key field depends on the length of the logical MAT defined in its program. Next, we remove dependency edges to the match control node from the selected key field resources. We add the resource dependency of the foreign key field to the match control node and add corresponding read-ready operations in the OSG of the field. In successive iterations of mapping physical stages, a sub MAT recomposed with the foreign key field resource in its match key fields is considered as the primary table being referenced by foreign key of the secondary table. A sub MAT can be primary table associated to multiple secondary tables referenced by their respected foreign keys. This creates multiple sub MATs using subsets of the match key set, hence vertically decomposing the MAT into multiple sub MATs. We store mappings between the decomposed MAT and its sub MATs (line 18).

To realize flow updates in the decomposed MAT, all of its sub MATs are required to be updated with consistency at runtime. As the logical MAT is decomposed in multiple sub MATs with foreign key references, we create update schedule for sub MATs of every vertically decomposed logical MAT. The schedule provides a sequence to update sub MATs such that primary sub MAT is followed by its secondary, thereby maintaining referential integrity and consistency across the sub MATs for match and actions on foreign key fields. We arrange sub MATs of a vertically decomposed MAT in reverse topological order of their mapped physical pipeline stages. We push new sub MAT in front of the update sequence of sub MATs for the vertically decomposed MAT (line 18). The Runtime system uses this schedule to update the sub MATs mapped to the stages of the pipeline.

4. Runtime

Out-of-order Write Operation

To explain scheduling of out-of-order write operations in an OSG, we start with the simplest case of an operation schedule (OS) list, followed by an OS tree and we finally address the directed acyclic graph case for OSG, as shown in Figure 8. Let us consider the first case shown in Figure 8a, where a resource's operation schedule is a list describing the sequence of operations on it. A3-W is scheduled out-of-order by performing all the ancestor operations of the node on the dummy resource R1', hence we split the schedule into two and move the one with all the ancestors to the dummy resource R1'. In case of a tree (shown in Figure 8b with A4-W as out-of-order write node), we additionally move all the read operation nodes (e.g., M 2-R) before any write operation, exclusively reachable from the ancestors (e.g., M 1-R) of the out-of-order write node. Finally, a tree can be transformed to a DAG by adding edges or paths from a node's siblings or ancestors to its descendants or itself, as shown in Figure 8c. In this case, the descendants of an out-of-order write node reachable using alternative (not involving the write node, e.g., M 5-R) paths from its ancestors can be scheduled only once all of its other predecessor nodes (e.g., M 2-R) are scheduled, including the predecessor of the out-of-order write node (i.e., M 1-R). Also, the value to be used for such nodes (M 5-R) depends on two resources (R1 and R1') and during execution either of the resource will hold valid value depending on control flow of the execution at runtime. Hence, we add a flag as a data plane resource, a match (M X) and an action node (AX) nodes to copy the value of R1' to R1 by matching on the flag as shown in Figure 8d. This adds one MAT in the merged pipeline in order to recompose and map the MAT to the current pipeline stage using actions node not having all the write operations in ready state. Linker may decompose a logical MAT by creating sub MATs using subsets of its match key fields. Flow entries of a decomposed MAT may have the same values for match key fields of some of its sub MATs. As sub MATs match flows using subsets of key fields of decomposed MATs, multiple flow entries of a decomposed MAT may share the same flow entry in sub MATs. At runtime, every secondary sub MAT assigns a row ID value, unique across the entries in the sub MAT, to its foreign key field in its action on successful match its key fields. Primary sub MATs match flows using the foreign key fields and other key fields of the decomposed MAT to further specialize the match on the flow entry and eventually execute the actions specified for the flow entry in the decomposed MAT.

Algorithm 5 realizes addition of new flow entry in a decomposed MAT. Before performing insertion, we match key fields of each sub MAT in reverse order of update schedule P4Bricks 4. Runtime (packet processing order) to identify every sub MAT already having flow entry for its subset key fields of new flow of the decomposed MAT (line 4). If a sub MAT does not have matching flow with its subset key fields of new flow (line 5), we create new row IDs for the foreign key field of the sub MAT. We create a match and an action for the foreign key field (lines 9-10). Next, we insert flow entries in all the sub MATs in order of update schedule generated by the Linker. If a sub MAT does not have flow entry with its sub set of key fields and foreign key fields, we add corresponding new flow entryin the sub MAT (lines 12-14). To delete a flow entry of a decomposed MAT (see Algorithm 6), we need to identify the flow entry stored across the sub MATs with intermediate foreign key fields. We match sub MATs in packet processing oder to find values of all the foreign key fields set by actions in sub MATs and store the sub MAT and foreign key field mapping.(lines 4-7). We match key field values of sub MATs in order of update schedule to delete flow entry of a decomposed MAT (line 8). We delete match entry from a secondary sub MAT (line 9) only if the foreign key fields values, set by the corresponding action, are not used for matching in primary sub MATs along with key fields of the decomposed MAT. If a primary sub MAT has multiple match entries having the same values for foreign key fields and different values for flow key fields, we do not further delete entries from secondary sub MATs of the decomposed MAT (lines 11-14).

To update a flow entry of a decomposed MAT (Algorithm 7), we identify the flow entry stored across the sub MATs along with intermediate foreign key fields. We match key fields of flow pertaining to each sub MAT in reverse order of the update schedule (lines 4-8). The last primary sub MAT matching rest of the key fields of flow and foreign key fields uniquley identifies match of entry of decomposed MAT (line 9). Finally, we update

Conclusion

In this report, we outline our first design to compose independently written P4 programs. We describe methods used to merge parser blocks, use ID mappings generated by them in merging of the pipelines. We show that it is possible to merge logical pipelines by considering packet headers and fields as shared resources and by applying composition operations to these shared resources. Deparser is represented as a control block in the P4 language, but we hope that our work provide a good motivation to rethink design of deparser to facilitate conceptually correct merging.

P4Bricks is implemented as an extension of the BMv2 switch simulator. The current implementation includes merging of parsers (3.1.1 to 3.1.4) and restructuring of pipelines (3.2.1 to 3.2.3 & 3.2.5) with out-of-order write operations. An implementation of the selection algorithm to map match key fields to physical stage is ongoing.

Figure 2 :

 2 Figure 2: Parser FSM Structure

P4Bricks 3 . 1 1 : 3 States 6 ifFigure 3 :

 311363 Figure 3: Merging parse graphs of two parsers

Algorithm 2 :

 2 Merging of parse graphs input : Parse graphsP G A = (S A , T A ) & P G B = (S B , T B ) of programs A & B EquivalentStatesMap -States mapping

Figure 4 :

 4 Figure 4: An example of CFG decomposition and OSG construction for a resource

Figure 6 :

 6 Figure 6: Parallel Composition of decomposed CFGs and OSGs

Figure 7 :

 7 Figure 7: DAG representation of physical pipelines in reconfigurable devices[START_REF] Jose | Compiling packet programs to reconfigurable switches[END_REF]

P4Bricks 3 . 2 4 : 2 S 6 foreach r in R do 7 foreach 9 foreach

 3242679 Logical Pipelines Restructuring Algorithm Mapping MATs to Physical Pipeline Stages input : P G physical = (S, E) -DAG of physical pipeline stages S & edges E OSG m (r) -merged OSG of all the resource r ∈ R dCF G m -merged all the program's decomposed CFGs output: MATMappings -Maps every MAT to Ordered list of its Sub MATs 1 while S is not empty do // Topological-order traversal of P G physical match ← GetNodesWithoutIncomingEdge(S) 3 foreach r in R do 4 ReadyMatchOps[r] ← GetSchedReadyReadMatchOpNodes (OSG m (r)) 5 AllocatedOps[r] ← SelectionOperations (ReadyMatchOps[r], S match ) Op in AllocatedOps[r] do 8 Remove edge representing Op in dCF G m PrevOp in Predecessor (Op) do 10 foreach SuccOp in Successor (Op) do 11 Add P revOp to SuccOp directed edge in the OSG m (r) 12 MappedMatchCN ← GetMatchControlNodesFromOps (AllocatedOps) 13 foreach MatchCN in MappedMatchCN do 14 ActionCN ← GetNextActionControlNode (M atchCN ) 15 if MatchCN or ActionCN is not complete read-ready or MatchCN is not control-ready then 16 Create a sub MAT using schedule-ready Ops by MatchCN 17 Add Foreign key field resource node in S 18 Push sub MAT in mapped ordered list for the logical MAT in MATMappings 19 else // MatchCN is schedule-ready and ActionCN is complete read-ready 20 if ActionCN has not schedule-ready write op on a resource r then 21 Perform out-of-order write in OSG m (r) 22Remove S match from P G physical with their incident edges

Figure 8 : 2 UpdateSchedList 8 ForeignKeyField 9

 8289 Figure 8: Out-of-order Write Operations (OWOs) and Types of Operation Schedules

Algorithm 6 : 4 foreach 6 8 foreach 9 if

 64689 Delete Flow at Runtime in a Decomposed MAT input : Flow F = {M kf v , A} with match M kf v of key field-value & actions set A mat -UUID of the table to update flow F MATMappings -Maps every MAT to ordered list of its Sub MATs 1 Function DeleteFlow(F , mat, UpdateSchedList) 2 UpdateSchedList ← MATMappings.get(mat)3 MatchKFSet ← F (M kf v ), SubMAT in ReverseOrder ( UpdateSchedList) do 5 ForeignKeyField, rowID ← GetActionEntry (SubMAT, MatchKFSet) MatchKFSet.Add(MatchKeyValue (ForeignKeyField, rowID)) 7 KeyMATMap.Add(ForeignKeyField, SubMAT) SubMAT in UpdateSchedList do SubMAT is in KeyMATMap then 10 Delete Flow entry f sub having subset of MatchKFSet 11 foreach ( ForeignKeyField, rowID) in f sub do 12 if HasMatch ( SubMAT, ( ForeignKeyField, rowID)) then 13 SecSubMAT ← KeyMATMap.GetSubMAT(ForeignKeyField) 14 Remove SubMAT match keys from MatchKFSet

Algorithm 7 : 2 UpdateSchedList 4 foreach 5 f 6 ForeignKeyField 7 Remove 9 if

 7245679 Modify Flow at Runtime in a Decomposed MAT input : Flow F = {M kf v , A} with match M kf v of key field-value & actions set A mat -UUID of the table to update flow F MATMappings -Maps every MAT to ordered list of its Sub MATs 1 Function ModifyFlow(F , mat, UpdateSchedList) ← MATMappings.get(mat) 3 MatchKFSet ← F (M kf v ), SubMAT in ReverseOrder ( UpdateSchedList) do sub (M kf v ) ← GetKeyFieldValues (MatchKFSet, SubMAT (M kf )) , rowID ← GetActionEntry (SubMAT, MatchKFSet) SubMAT match keys from MatchKFSet 8 MatchKFSet.Add(MatchKeyValue (ForeignKeyField, rowID)) MatchKFSet is empty then 10 Update Flow entry with match f sub (M kf v ) with actions set A 11 return action entry in the primary sub MAT (line 10).

  compatible reconfigurable target. Parser blocks are programmed by defining packet header types and fields, declaring instances of the types and defining a Finite State Machine (FSM) to extract the header instances from the packets' bit streams. The parse graph is encoded as FSM in P4. The FSM structure is described by defining states and transitions among the states. The programmer must define one start state for FSM of a P4 parser. P4 provides two logical final states named accept and reject, which are not part of the FSM defined by the programmer. The programmer can define next transition to accept state or reject state

from any state of the parser to notify successful completion or failure packet in parsing, respectively. Each state has a name and a body comprising of a sequence of statements. The statements within a state's body describe the processing to perform when the parser P4Bricks 2.1 P4 background transits to the state. They can be local variable declarations, assignments, function invocation(e.g., verify -to validate already parsed data

  1 Merging Parsers and Deparsers in a parse state of one parser, the parser transits to reject state. However, the equivalent state from another parser may not have verify statement with the same condition and can transit to different states instead of the final reject state. Moreover, we illustrate the need of traffic isolation emerging from different possible transitions from equivalent states of two parsers with the example shown in Figure3, where one parser transits to a final state and another to an intermediate state from the equivalent state. In case of double tagged VLAN packet, the Enterprise program transits to reject state after extracting first VLAN header (V LAN x ). However, the Data center program continues extracting next VLAN header. Hence, parsed packet headers of double tagged VLAN traffic should be exclusively processed by the MATs of the Data center program and MATs of the Enterprise program should process VLAN traffic only if the packets have a single tag. Another example, if the match keys of MATs defined in both the programs use only TCP header fields, the IPv6 based TCP traffic should not be processed by the MATs of the Data center program. In order to provide traffic isolation emerging from parsers resulting in different final states for the same packet and seamless error handling to all the parsers, we devise an error handling mechanism constituting an intrinsic metadata field and P4Bricks target specific implementation of the verify call.

  For each P4 program, Linker allocates the parser error variable to hold error code generated by the program's parser, if the error base type is defined in configuration files of the program. Here, program id is used to identify the parser error variable and set it to err, if condition evaluates to false. Also, the bit associated with the caller P4 program in indicator is set to 0 to indicate invalidity of the packet header instances for the program. The semantics of the verify statement for P4Bricks enabled data plane

	P4Bricks	3.1 Merging Parsers and Deparsers
	architecture is given below.
	P 4B r i ck s P a rs e r Mo d e l . verify ( bool condition , error err , int
	program_id , bit <W > indicator ) {
	if ( condition == false ) {
	P 4 Br i c ks P a rs e r Mo d e l . parseError [ program_id ] = err ;
	indicator = indicator & (~(1 << n ) ) ;
	}
	if ( indicator == 0)
	goto reject ;
	}
	extern void verify ( in bool condition , in error err , in int
	program_id , inout bit <W > indicator ) ;

  Keyset B , N S B ) pairs in S B .KeysetToNextStateMap do Identifying equivalent header instances in sequential processing input : P G M erged = (S M , T M ) Merged Parse graph EmitHIList -Emit order sequence of header instances

	P4Bricks		3.1 Merging Parsers and Deparsers
	Algorithm 3:
	15	else	// Unique keyset across the case-lists
				// Recursive call
	20	S M ab .AddSelectCase(Keyset A , N S M ab )	// Add Merged State
		transitions
	21 for remaining (22 Repeat lines 16-20 using Keyset B , N S B , S B
	23	return S M ab

1 Function MergeParseGraph(P G A , P G B ) 2 InS A ← InitState(S A ), InS B ← InitState(S B ) 3 MergeStates (InS A , InS B , 0b0) // Merge init states of both parse graphs 4 Function MergeStates(S A , S B , IndF lag) // Recursively merge parse states 5 if S A and S B are already merged then 6 return merged state S M ab 7 Set bits for program A and B in IndM ask M ab 8 IndM ask M ab ← IndM ask M ab IndF lag // Merge state transitions using select case-lists of S A and S B 9 S B .KeysetToNextStateMap ← CreateMap (S B .SelectCaseList) 10 foreach (Keyset A , N S A ) pair in S A .SelectCaseList do 11 N S B ← S B .KeysetToNextStateMap.Find(Keyset A ) 12 if N S B then // S A & S B have transitions for the same keyset // Next States must be Equivalent 13 Assert (EquivalentStatesMap (N S B , N S A ) == 1) 14 Remove Keyset A from S B .KeysetToNextStateMap 16 N S B ← EquivalentStatesMap.Find(N S A )) 17 if S A is accept-transition state then 18 set bit for program A in IndF lag // packet valid for program A 19 N S M ab ← MergeStates (N S A , N S B , IndF lag) 1 while EmitHIList not empty ∨ S M = ∅ do 2 States B ← GetNodesWithoutIncomingEdge(S M ) 3 foreach hi E in EmitHIList do // Mapping instances using header type 4 foreach s M in States B do 5 hi M ← GetHeaderInstanceExtractedFromState (s M ) 6 if HdrTypeof(hi E ) ≡ HdrTypeof(hi M ) and hi M is unmapped then 7 Map (hi E , hi M ) in U IDT HeaderInstance 8 Remove s M and their outgoing edges 9 break 10

Acknowledgements

This work has been partly supported by the French ANR under the "Investments for the Future" Program reference #ANR-11-LABX-0031-01.


Safa Ouerghi

Rémi Boutteau

Xavier Savatier
email: xavier.savatier@esigelec.fr

Fethi Tlili
email: fethi.tlili@supcom.tn

Remi Boutteau

CUDA Accelerated Visual Egomotion Estimation for Robotic Navigation

Keywords: Egomotion, structure from motion, robotics, CUDA, GPU

published or not. The documents may come   L'archive ouverte pluridisciplinaire

INTRODUCTION

Accurate localization is a fundamental issue in autonomous navigation that has been extensively studied by the Robotics community. During the last years, cameras have become very popoular in Robotics allowing the developement of several vision-based methods for a real time localization. These methods primarily follow two main paradigms namely, SLAM (Simultaneous localization and Mapping) [START_REF] Durrant-Whyte | Simultaneous localisation and mapping (slam): Part i the essential algorithms[END_REF][START_REF] Dissanayake | A solution to the simultaneous localization and map building (slam) problem[END_REF] and real-time SFM (Structure from Motion) or Visual Odometry [START_REF] Nister | Visual odometry for ground vehicle applications[END_REF][START_REF] Maimone | Two years of visual odometry on the mars exploration rovers[END_REF]I. Comport et al., 2010). While SLAM methods tackle the issue of concurrent localization and mapping of a vehicle in an unknown environment, visual odometry calculates the egomotion by incrementally estimating the rotation and translation undergone by the vehicle using only the input of a single or multiple cameras. The visual odometry pipeline for the stereo scheme consists mainly in finding corresponding features between adjacent images in the video sequence and using the scene's epipolar geometry to calculate the position and orientation changes between the two images. A common way of determining the relative pose using two images taken by a calibrated cam-era is based on the estimation of the essential matrix that has been studied for decades. The first efficient implementation of the essential matrix estimation is proposed by Nister in [START_REF] Nister | An efficient solution to the five-point relative pose problem[END_REF] and uses only five point correspondances. The work of Stewenius built upon the work of Nister uses the Gröbner Basis to enhance the estimation accuracy [START_REF] Stewenius | Recent developments on direct relative orientation[END_REF]. However, in a real application, wrong matches can lead to severe errors in the measurements, which are called outliers and that occurs during the descriptors matching step. The typical way of dealing with outliers consists of first finding approximate model parameters by iteratively applying a minimal solution in a hypothesize-and-test scheme. This procedure allows us to identify the inlier subset, and then, a least-squares result is obtained by minimizing the reprojection error of all inliers via a linear solution or a non-linear optimization scheme, depending on the complexity of the problem. This scheme is called RANSAC and has been first proposed by Fischler and Bolles [START_REF] Fischler | Random sample consensus: A paradigm for model fitting with applications to image analysis and automated cartography[END_REF]. RANSAC can often find the correct solution even for high levels of contamination. However, the number of samples required to do so increases exponentially, and the associated computational cost is substantial. Especially for robotics systems, the challenges are more acute, due to their stringent time-response requirements. To solve these problems with a certain amount of speedup, the usage of GPU computation is a popular topic in the community. The Compute Unified Device Architecture (CUDA) has enabled graphics processors to be explicitly programmed as generalpurpose shared-memory multi-core processors with a high level of parallelism [START_REF] Lindholm | Nvidia tesla: A unified graphics and computing architecture[END_REF]. In fact, recently, many problems are being solved using programmable graphics hardware including feature matching and triangulation [START_REF] Li | Pitch angle estimation using a vehicle-mounted monocular camera for range measurement[END_REF], feature detectors [START_REF] Yonglong | Parallelization and optimization of sift on gpu using cuda[END_REF], large non-linear optimization problems such as bundle adjustment [START_REF] Wu | Multicore bundle adjustment[END_REF] and learning algorithms [START_REF] Chang | Libsvm: A library for support vector machines[END_REF].

In this paper, we focus on an efficient implementation of a state-of-the-art relative pose estimation based on the computation of the Essential matrix from five correspondances. We consider single GPU implementation and we describe the strategies to map the problem to CUDA architecture. Futhermore, new Kepler and Maxwell architecture features are used and analyzed, such as CUDA Dynamic Parallelism and new CuBLAS batched interfaces.

The outline of this paper is as follows: we briefly present the theory underlying the essential matrix estimation in section 2. Section 3 details the CUDA based implementation of the fivepoint essential matrix estimation algorithm within RANSAC. Afterwards, section 4 shows several experiments as examples of the speedup results obtained with our implementation. Finally section 6 gives the conclusion of the paper.

BACKGROUNDS

In this section, we provide an overview of the important backgrounds underlying the robust essential matrix estimation from 5 correspondances.

Essential matrix

The essential matrix E is a 3 × 3 matrix expressing the epipolar geometry between two calibrated camera systems [START_REF] Hartley | Multiple View Geometry in Computer Vision[END_REF]. That is, if a 3D point X is viewed in two images at locations u and v (where u and v are calibrated homogeneous image coordinates), then u and v are related such that v T E u = 0.

(1) Furthermore, if the two views have relative pose

[R|t] then E = [t] × R, (2)
where [t] × is the skew-symmetric matrix with the property that [t] × x = t × x.

However, from two images alone, the length of t cannot be determined. Therefore, E is only determined up to a scale.

Expanding Equation 1 gives a single linear constraint in the nine elements of E for every correspondence. From N correspondences, these equations can be stacked to form a 9×N matrix which null space obtained by singular value decomposition (SVD) gives a basis for the space in which E lies. The points within this vector space which are essential matrices are those which can be decomposed into a rotation and a translation. E can be decomposed in this way using an SVD decomposition

E = U   s 0 0 0 s 0 0 0 0   V T ,
which is equivalent to the following constraint providing an efficient test whether a matrix is approximately an essential matrix

EE T E - 1 2 trace(EE T )E = 0.
(3)

Computing E from five correspondances

Several algorithms have been developed to estimate the essential matrix, including, the seven-and eightpoint algorithms that are relatively fast [START_REF] Hartley | Multiple View Geometry in Computer Vision[END_REF]. However, for their use within RANSAC, essential matrix computations have relied on minimal subsets, which for essential matrix is five correspondences. Furthermore, Essential matrix estimation from five correspondances have shown a better accuracy than other faster algorithms with more correspondances. In essential matrix estimation, given five correspondences, four basis vectors satisfying Equation 1 can be computed by SVD. All linear combinations of these basis vectors satisfying Equation 3 are essential matrices that provide nine cubic constraints in the elements of E. The methods of Nister [START_REF] Nister | An efficient solution to the five-point relative pose problem[END_REF], and Stewenius et al. [START_REF] Stewenius | Recent developments on direct relative orientation[END_REF] both work by solving these nine equations. Stewenius et al. first showed that the equations can be written as MX = 0, where M is a 10 × 20 matrix. After gauss-jordan elimination, the system can be written [I B]X = 0, where I is a 10×10 identity matrix and B a 10×10 matrix.

Stewenius et al. used, subsequently, the action matrix concept to solve the systems in which a Gröbner basis is found. The 10 × 10 action matrix real eigenvalues and eigenvectors contain, hence, the solutions of polynomial equations.

Relative pose computation from Essential matrices solutions

Once the essential matrices solutions are computed, they have to be decomposed into rotation and translation. In fact, the decomposition follows the normal procedure for the general case [START_REF] Nister | An efficient solution to the five-point relative pose problem[END_REF], giving two possible solutions for the rotation, R a and R b , and two solutions for the translation as well, t a and t b , which have the same direction t determined up to a scale. Thus, if E ∼ USV T is the SVD of E, a matrix D is defined as

D =   0 1 0 -1 0 0 0 0 1   . (4) Then, R a = UDV T and R b = UD T V T . The solu- tion for the translation direction is t = [U 13 U 23 U 33 ] T .
Four pose configurations are, therefore, obtained for each essential matrix namely, (R a , t a ), (R b , t a ), (R a , t b ) and (R b , t b ). Consequently, a disambiguation has to be performed to output the correct movement undergone by the camera.

Robust estimation of the Essential matrix

Even if the underlying dataset is contaminated with outliers, RANSAC estimator can be used to robustly estimate the model parameters. RANSAC uses a randomly chosen subset of the entire dataset to compute a hypothesis. The remaining datapoints are used for validation. Repeating the hypothesis computation and validation with different subsets, the probability of finding a hypothesis that fits the data well increases. For a data set with a given proportion ε of outliers, the number of trials N required to give sufficiently high probability p to pick an outlier-free subset consisting of k point correspondences is

N = log(1 -p) log(1 -(1 -ε) k ) (5)
Since the confidence p is generally chosen to be p ≥ 0.99, the number of required RANSAC iterations N only depends on the number of parameters k and the assumed ratio of outliers ε. Usually, the ratio of outliers ε is unknown. Hence, we resort to an adaptive version of RANSAC, where, after each iteration, the number of inliers γ is counted and the outlier ratio is updated according to

ε = 1 - γ n , (6)
with n equal to the size of the dataset. The number of iterations N is therefore updated based on ε.

GPU IMPLEMENTATION OF RELATIVE POSE FROM FIVE 2D-2D CORRESPONDENCES

In this section, we present the implementation details of the essential matrix estimation from five correspondences within RANSAC using the CUDA programming model. As stated before, the eigenvalues of the action matrix contain the essential matrices solutions according to Stewenius's method [START_REF] Stewenius | Recent developments on direct relative orientation[END_REF]. However, a device based eigenvalue computation on CUDA doesn't exist yet. Hence, we have relied on the Matlab code provided by Chris Engels, based on the reduction to a single polynomial [START_REF] Stewenius | Matlab code for solving the fivepoint problem[END_REF]. This is done through the computation of the action matrix characteristic polynomial roots, equivalent to the action matrix eigenvalues.

Our parallelization approach is based on performing the required RANSAC iterations in parallel on CUDA to achieve a certain amount of speedup. This level of parallelism suggests the consideration of RANSAC iterations as a batch of parallel computations, each processing a small subset of data. Furthermore, we have relied on the use of cuBLAS, a highperformance implementation of BLAS-3 routines, for linear algebra computations (NVIDIA, 2015). As the matrices sizes in our problem are below 32 × 32, we have particularly exploited the batched interface of the cuBLAS librairy where many small dense matrices factorizations, to be performed simultaneously, are provided.

Thus, in total, four kernels have been employed operating at different levels of parallelism. The first, exploits the CuBLAS library batched interface, manages algebraic computations. It employs, therefore, a thread level parallelism and a nested warp level parallelism as it uses dynamic parallelism to call cuBLAS functions from within device. The second employs a straightforward parallelization and works at a threadlevel parallelism where each thread manages the remainder computations after the completion of the first kernel, i.e. one thread per RANSAC iteration. The third kernel is used to rate the models outputted by the previous kernel and works at a block level parallelism where each block validates a model relative to one RANSAC iteration. Finally, an additional kernel is used to compute RANSAC's best model and it simply performs a reduction to find the model with maximum number of inliers which represents the best model.

Cublas based kernel

This kernel is launched with one block and a number of threads equal to the number of required RANSAC iterations. The high level interface exposed by all implementations in this kernel is cuBLAS batched interface for solving a batch of N different systems with double precision. Besides the batch size and the matrix dimensions, the functions expect pointers to array of matrices. All arrays are assumed to be stored contiguously with a column major layout and accessed to in global memory through the handle of an array of pointers that we statically allocate as follows:

__device__ double* PointersArray[MaxBatchSize]
Firstly, a 9 × 5 hypothesis A[i], i = 0...batchSize -1 is computed from each random five correspondances by each thread. The hypotheses are written to global memory and referenced by an array of pointers as indicated above.

Secondly, the null-space of each hypothesis have to be computed by SVD. However, due to the absence of a GPU-based implementation of SVD decomposition, we use instead a QR decomposition to derive the null space. In fact, standard methods for determining the null space of a matrix are to use a QR decomposition or an SVD. If accuracy is paramount, the SVD is preferred but QR is faster. Using a QR decomposition, if A T = QR, and the rank of A is r, then the last nr columns of Q make up the null space for A. This is performed through a call to the cublas built-in function cublasDqr f Batched performing a QR factorization of each A[i] for i = 0, ..., batchSise-1. The decomposition output is presented in a packed format where the matrix R is the upper triangular part of each A[i] and the vectors v on the lower part are needed to compute the elementary reflectors. the matrix Q is, hence, not formed explicitly, but is represented as a product of these elementary reflectors.

As cublas dosen't provide a built-in routine to retrieve Q as Lapack does, we designed a child kernel called from the main kernel to simulaneously calculate the different reflectors and compute their product to retrieve Q.

The number of Thread-blocks in the launch configuration of the child kernel is identical to the batchSize, i.e. iterations. Each Thread-block computes a single matrix Q and a block-level parallelism is hence applied. The Thread-blocks are designed to be three-dimensional, where the x-dimension refers to the number of rows of each reflector, the y-dimension to the number of columns and the z-dimension to the number of reflectors. This allows each thread to handle one element in shared memory and consequently ensure a parallel computation of the different reflectors. It is worth noting that this configuration is possible because the matrix sizes in our problem are small (5 refrectors, each of size 9 × 9) and consequently, all reflectors fit at once in shared memory. The computation consists in loading, first, the A[blockIdx.x], and the array of scalars Tauarray[blockIdx.x] exited by cublasDqr f Batched into shared memory where the matrix Q is also allocated. The vector v i relative to each reflector q i is then putted in the required form, where v i (1

: i -1) = 0 and v i (i) = 1 with v i (i + 1 : m) on exit in A[blockIdx.x][i + 1 : m, i].
Each reflector q i has the form q i = I -Tau[i].v.transpose(v), computed for all reflectors by the pseudocode explicited in Figure 1 and finally, the product of all reflectors is computed to retrieve Q. Once the null space determined, the second step is to compute a 10 × 20 matrix M that is accelerated in the provided openSource code, through a symbolic computation of the expanded constraints. The matrix columns are then rearranged according to a predefined order. To save execution time and memory usage, we use to rearrange the matrix columns beforehand and to write it in column major for subsequent use of cublas functions. We hence output a permuted 20×10 matrix M .

Subsequently, the Reduced Row Echelon Form (RREF) of M have to be computed through a gauss jordan elimination, i.e. M = [I B]. Instead of carriying out a gauss jordan elimination on M, a factorization method may be used to find directly the matrix B from the existant matrix M. In fact, Cublas provides several batched interfaces for linear systems factorizations. We exploit the batched interface of LU factorization performing four GPU kernel calls for solving systems in the form (MX = b) as follows: 1. LU decomposition of M (P M = LU). 2. Permutation of the array b with the array of pivots P (y = P b). 3. Solution of the triangular lower system (L c = y). 4. Solution of the upper system to obtain the final solution (U x = c) With putting b as an array of pointers to null vector, Cublas directly provides cublasDgetrfBatched for the first step and cublasDgetrsBatched for the three other steps. We finally obtain the matrix B in exit of cublas-DgetrsBatched, solution of the system MX = 0.

RANSAC models computation kernel

At this level, the kernel is launched with one CUDA block and iterations number of threads. We only use global memory where the computations of the previous kernel are stored and small per thread arrays using registers and local memory. Each thread computes a 10 th degree polynomial using local variables. This is done by extracting from the RREF in global memory the coefficients of two 3 rd degree polynomials and a 4 th degree polynomial represented by private local arrays for each thread. These polynomials are afterwards convoluted then substracted and added to generate a single 10 th degree polynomial for each thread as explicited in the original Matlab code and which refers to the computation of the determinant of the characteristic polynomial. The covolution is performed in our implementation through a special device function presented as a symbolic computation of three polynomials of 3 rd , 3 rd and 4 th degrees respectively.

The key implementation of this kernel is the resolution of a batch of 10 th degree polynomials. In fact, we used a batched version of the Durand-Kerner Method in which we assign to each polynomial a thread. We start by giving a brief overview of the Durand-Kerner method, followed by our implementation details.

Durand-Kerner Method

The Durand Kerner Method allows the extraction of all roots ω 1 , ..., ω n of a polynomial

p(z) = n ∑ i=0 a i z n-i , ( 7
)
where a n = 0, a 0 = 1, a i εC. This method constructs a sequence, H(z k ) = z k+1 in C N with Z (0) being any initial vector and H is the Weierstrass operator making Z (k) i tends to the root ω i of the polynomial, defined as:

H i (z) = z i - P(z i ) ∏ j =i (z i -z j ) i = 1, ..., n
The iterations repeat until

|Z k i -Z k+1 i | Z k i or |P(z k i )
| is smaller than the desired accuracy.

GPU version of batched Durand-Kerner

The implementation of the Durand-Kerner on GPU, is basically sequential where each thread computes the ten complex roots of the 10 th degree polynomial. We defined the type COMPLEX denoting structs of complex numbers. We started from an initial complex guess z randomly chosen, and the vector of complex roots R of size the number of roots (10 in our problem) where, R[i] = z i , i = 1..n -1. The function poly evaluates at z a polynomial of the form of Equation 7where the vector A = a1, a2, a3, ..., a(n -2), a(n -1), a(n) denotes the coeffecients of our polynom.

As we are dealing with complex numbers, complex arithmetic has been employed denoted by compsubtract for complex numbers substraction and compdiv for complex division. As explicited in the following piece of code, we iterate until obtaining the desired accuracy expressed as a relative error of estimated roots below a predefined value as depicted in Figure 2.

Pseudocode2: GPU Version of Durand-Kerner method double maxDiff = 0; int iter=0; int maxIters =30; for( iter = 0; iter < maxIters; iter++ ) { maxDiff = 0; for (int j = 0; j < n; j ++) { COMPLEX B = poly(A, n, R[j]); for As explicited in Subsection 2.3, an SVD decomposition of the directly obtained essential matrices which are up to 10 (real solutions of 10 th degree polynomial) is used to decompose each solution into rotation and translation. This operation can take a significant portion of the computation time and we use, therefore, a specifically tailored singular value decomposition for essential matrices according to Equation 3, that is proposed in [START_REF] Nister | An efficient solution to the five-point relative pose problem[END_REF]) (Appendix B). In our implementation, each thread computes up to 10 essential matrices, and for each, four movement configurations are obtained.

(int k = 0; k < n; k++) { if (k != j) B = compdiv(B,compsubtract(R[j] , R[k])); } R[j] = compsubtract(R[j],B); maxDiff = max(maxDiff, abs(B.x)); } if( maxDiff <= 1e-10) break; }
However, in order to deal with all central camera models including perspective, dioptric, omnidirectional and catadioptric imaging devices, image measurements are represented as 3D bearing vectors: a unit vector originating at the camera center and pointing toward the landmark. Each bearing vector has only two degrees of freedom, which are the azimuth and elevation inside the camera reference frame as formulated in [START_REF] Kneip | Opengv: A unified and generalized approach to real-time calibrated geometric vision[END_REF]. Because a bearing vector has only two degrees of freedom, we frequently refer to it as a 2D information and it is normally expressed in a camera reference frame.

The disambiguation step that has, finally, to be performed by each thread consists in calculating the sum of reprojection error of the trinagulated 3D points relative to the corresponding bearing vectors used to compute the model. Finally, a single 4×3 transformation into the world reference frame matrix is returned by each thread referring to the lowest score of reprojection error between all essential matrices and pose configurations (up to 40). The transformation matrix is directly obtained from the already calculated rotation and translation.

Indeed, the triangulation method used in our implementation follows the general scheme employed in [START_REF] Kneip | Opengv: A unified and generalized approach to real-time calibrated geometric vision[END_REF]. The reprojection error of 3D bearing vectors was proposed in [START_REF] Kneip | Opengv: A unified and generalized approach to real-time calibrated geometric vision[END_REF] as well, and is computed by considering the angle between the measured bearing vector f meas , and the reprojected one f repr . In fact, the scalar product of f mes and f repr directly gives the angle between them, which is equal to cos θ as illustrated in Figure 3. The reprojection error is, therefore, expressed as

ε = 1 -f T meas f repro = 1 -cos θ. (8)

RANSAC rating kernel

In order to validate each estimated model, we compute a loss value for each datapoint of the dataset. The loss value is used to verify the model by computing the reprojection error of all triangulated bearing vectors of the dataset. Outliers are subsequently found by thresholding the reprojection errors, and the best model refers to the one with the maximum number of inliers. As the entire operation is in 3D, we use the thresholding scheme adopted in the Opengv librairy [START_REF] Kneip | Opengv: A unified and generalized approach to real-time calibrated geometric vision[END_REF]. This latter uses a threshold angle θ threshold to constrain f repr to lie within a cone of axis f meas and of opening angle θ threshold as depicted in Figure 3. The threshold error is given by

ε threshold = 1 -cos θ threshold = 1 -cos(arctan ψ l ), (9)
where ψ refers to the classical reprojection error threshold expressed in pixels and l to the focal length.

The model validation process considers multiple accesses to global memory to evaluate whether each correspondance of the dataset is an inlier or an outlier which is a very execution-time consuming. The shared memory is, hence, used as a cache to accelerate computations. The RANSAC rating kernel employs a block level parallelism and is lauched with iterations blocks to make each block handles a RANSAC model and 8 × warpsize threads. Since warpsize = 32, a total of 256 threads is launched per block and each thread in the block evaluates a point. To load datapoints in shared memory, a buffer is allocated of size 256 × s where s refers to the size of each datapoint. In case of bearing vectors, s = 6. Each thread triangulates bearing vector correspondances into a 3D point and computes its reprojection error according to Equation 8. This latter is, thereafter, compared to the precalculated threshold according to Equation 9 to decide whether the correspondance refers to an inlier or to an outlier. In our implementation, the number of inliers for 256 values is automatically returned via:

inlier_count=__syncthreads_count( reproj_error[threadIdx.x]<threshold);
The process of loading data into buffer and evaluating 256 reprojection errors is repeated ceil(datasetCount/256) times.

RANSAC best model computation kernel

This kernel is launched with one block and itearations threads and performs a reduction in shared memory to derive the best model which refers to the one with the maximum number of inliers.

In this section we evaluate the speed and accuracy of our CUDA based essential matrix solver within RANSAC and compare it against the CPU based implementation for general relative camera motion provided in the OpenGV library. This latter is an open-Source librairy that operates directly in 3D and provides implementations to solve the problems of computing the absolute or relative pose of a generalized camera [START_REF] Kneip | Opengv: A unified and generalized approach to real-time calibrated geometric vision[END_REF].

Random Problem Generation

To make synthetic data for our tests, we used the automatic benchmark for relative pose included in the Matlab interface of the OpenGV librairy. We used the provided experiment to create a random relative pose problem, that is, correspondences of bearing vectors in two viewpoints using two cameras at the number of 1000 correspondences. In fact, the number of 1000 correspondences has been chosen based on an averaged number obtained from real images. The experiment returns the observations in both viewpoints (assumed to be a stereo camera system), plus the ground truth values for the relative transformation parameters.

Timing

We have measured the mean time while running on the GPU and CPU (using OpenGV librairy). To compute the mean time, each estimation is repeated 20 times. The repetition rate is required since a single estimation can be much slower or much faster than the mean due to the randomization. We will present results of computations for both single-precision and double-precision datatypes. The system on which the code has been evaluated is equipped with an i7 CPU running at up to 3.5 GHz, the intel i7 CORE. The CUDA device is an NVIDIA GeForce GTX 850M running at 876 MHz with 4096 MB of GDDR5 device memory. The evaluation has been performed with CUDA version 7.5 integrated with VisualStudio 2012. At the first execution of the estimation, memory allocations have to be performed. This is required only once and takes about 6ms. To evaluate our implementation, 12 outlier ratios from ε = 0.05 to ε = 0.6 in steps of ε = 0.05 are evaluated. In Figure 4, we show the performance results of estimating camera relative pose from sets of 2D bearing vectors correspondences. Firstly, in Figure 4a, we compare the mean computation time of CPU and GPU implementations, in single and double precision. We show a mean computation time even more important for CPU reaching 86ms for an outlier ratio ε = 0.5 against 20.2ms for GPU in double precision and 18ms in single precision. With an outlier ratio of ε = 0.5 which is common for the essential matrix estimation from automatically computed point correspondences, we show in figure 4b that the speedup is about 4× compared to the CPU implementation in single and double precision. The speedup becomes more important for higher outlier ratios reaching 7× in single precision for ε = 0.6 against 4× in double precision. Furthermore, it is useful to visualize the intersection between each CPU and GPU evaluation, i.e. the outlier ratio where the speedup is equal to one. Figure 4 shows that there is no speedup for lower outlier ratios ε ≤ 0.2. This is because the needed number of iterations for ε = 0.2 is only 12 iterations. However, the minimum number of iterations used in GPU based implementation is 32 iterations referring to the warp size.

Accuracy

In Figure 5, we present the Root-mean-square error (RMSE) between the ground truth rotation matrix and our Cuda based implementation rotation matrix of the best model for different outlier ratios, for both single and double precision datatypes. Overall, single precision datatype showed good performance while en-suring higher levels of speedup for higher outlier ratios. The accuracy loss is mostly due to the batched durand-kerner version for solving 10 th degree polynomials where the maximum number of iterations is fixed.

CONCLUSIONS

In this paper we have presented a 2D-2D robust motion estimation on CUDA which is applicable to a wide range of problems and especially to autonomous navigation. We presented our parallelization strategy, based mainly on performing the required RANSAC iterations in parallel. We described our implementation dealing with several levels of parallelism namely, warp level parallelism, block level parallelism and thread level parallelism. In addition, we adapted the five-point essential matrix using Gröbner basis to CUDA ressources and programming model. Furthermore, we described our RANSAC implementation and the rating measure used which is based on the computation of the reprojection error of triangulated 3D points from bearing vectors. An evaluation of our implementation has been presented and the mean computation time of RANSAC for different outlier ratios has been measured. Overall, the implementation showed good performance, and a speedup 4 times faster than the CPU was measured for an outlier ratio ε = 0.5, common for the essential matrix estimation from automatically computed point correspondences.

More speedup was shown when dealing with higher outlier ratios.

Figure 1 :

 1 Figure 1: Pseudocode of reflectors computation in shared memory.

Figure 2 :

 2 Figure 2: Pseudocode of batched Durand-Kerner method on CUDA.

Figure 3 :

 3 Figure 3: Reprojection error computation in Opengv (Source: (Kneip and Furgale, 2014)).

  Performance of essential matrix estimation with RANSAC.

Figure 5 :

 5 Figure 5: RMSE in rotation between ground truth and CUDA best model.