M Stéphane

Renée El Melhem

Blaise Pascal

Keywords: positional notation of 42, 75, 75, 1, 8 Binary-coded decimal positional notation of 42, 75, 1, 9 Basic floating-point representation of 42, 75 as 0, 333984375 × 2 7 multiplexer

Les techniques de synthèse de haut niveau permettent aux programmeurs non spécialistes de générer des descriptions de circuits numériques en utilisant des langages de programmation généralistes. Cependant, les outils existants ne supportent qu'un petit nombre de formats numériques et un petit nombre d'opérateurs standards. Cette thèse présente plusieurs techniques pour rajouter le support de nouveaux formats et de nouveaux opérateurs. Dans un premier temps, l'étude se focalise sur ce qui est réalisable en se restreignant aux fonctionnalités de métaprogrammation du standard C++ supporté par les outils HLS. Une bibliothèque d'opérateurs élémentaires pour les formats IEEE-754 et posit de taille arbitraire est proposée. Elle sert de base à une étude de cas comparant le coût matériel de l'implémentation de ces deux formats. L'implémentation d'évaluateurs de fonctions mathématiques arbitraires se heurte aux limites de la première approche. Dans un second temps, l'étude se porte sur les possibilités offertes par la modification du flot de compilation HLS, avec comme objectif de supporter cette fonctionnalité. Une bibliothèque permettant au développeur de spécifier des opérateurs pour approximer des fonctions arbitraires en précision arbitraire est présentée. Deux approches pour l'interfaçage de cette bibliothèque avec les outils de HLS sont proposées, selon que l'on a ou pas accès aux sources des compilateurs HLS.

Acknowledgement

It has been a real pleasure to work on this thesis. While there have been many ups and downs along the way 1 , it has been an interesting journey, and many people have contributed making the interesting part of it more interesting, and the least enjoyable parts, well, less less enjoyable.

Among these people I would like to thank my supervisor, Florent de Dinechin, for the many insightful discussions and the support I could get from him, starting from my first internship on FloPoCo up to the end of the thesis defense.

I am also grateful to Ronan and Gauthier, who I enjoyed working with at Xilinx/AMD. In addition to providing interesting points of view on software development, they have been of great help for the development of the different backends for MArTo arbitrary function support.

I am not sure if I need to be grateful to SARS-CoV-2 without which I would probably not have met them.

I am also thankful to 小泉さん, my former supervisor at NTT Multimedia Intelligence Laboratory, for his advice about perseverance and failure relativization, which proved very valuable these last few years.

Working on this thesis wouldn't have been half as nice without the good mood of the office. So thank you to Agathe, Benoît, Diane, Lélio, Marie, Matthias, Orégane, and Pierre for being such nice colleagues! While I am of course grateful to all my family, a special mention goes to my brother and sister for their continuous support. Another one for the logistic support I received from my parents and my family in Lyon.

He (probably) did not do it on purpose, but Mochi has also been of a great help. Indeed, real-life cats are clearly underrated compared to rubber ducks for debugging purposes.

Finally, despite their constant efforts to give me more time to work on the manuscript, I do not thank the SNCF.

Introduction

Transistor density on integrated circuits has stopped to follow Moore's law predictions around year 2010. This is around the same time that clock rate ceased to be improved by new generations of central processing units (CPU). As a result, the main selling point of new generations of desktop and server CPUs has shifted towards more parallelism and more hardwaresupported operations. Compute-intensive applications benefit from these two improvements. Indeed, hardware-accelerated mathematical functions and fine-grained parallelism (such as Single Instruction Multiple Data (SIMD) operations or dynamic reordering of instructions for superscalar execution) reduce the runtime of one kernel iteration, while coarse-grain parallelism (through multiple cores or simultaneous multithreading) allows running multiple independent kernel invocations simultaneously, hence reducing the total execution time.

However, the set of numerical formats and operations supported by CPUs does not necessarily fit tightly the application needs. Moreover, computation kernels might expose some parallelism at a degree that cannot be handled properly via CPUs.

In order to further accelerate applications suffering from this kind of CPU limitations, application-specific integrated circuits (ASIC) can be developed. Two main barriers limit this development. The first one consists in the fixed costs associated with custom integrated circuit production. This barrier can be lowered by using Field-Programmable Gate Arrays (FPGAs) instead of ASICs, at the price of some performance loss. The second barrier, common to ASIC and FPGAs, is the required expertise needed to design those circuits.

Lowering the requested experience to develop this kind of hardware is one of the objectives of High-Level synthesis tools. These tools allow building a circuit from a high-level behavioral description written in a classical programming language, such as C++. This thesis focuses on these tools.

While HLS tools allow software developers to write circuits easily, their support for numerical formats and mathematical functions is limited to those provided by the source language. For instance, the only floating-point formats supported by C++-based HLS tools are the 32-bits and 64-bits IEEE-754 formats float and double. Yet, using application-specific numerical formats and operators is a way to improve the quality of results (QoR, which regroups design latency and surface) of circuit designs. For instance, as elementary operations on narrower formats operate on smaller bit vectors, their latency is reduced compared to the same operation on wider a format. Besides, they also occupy less chip area, which leaves more space for packing more operators on the chip. This eventually allows running more operations in parallel. Using for each computation the narrowest format that still allows to verify the overall application accuracy constraint is then a good way to optimize this application.

The work presented in this thesis aims at providing support for custom numerical formats and operators in an HLS context. This support takes the form of a portable library, MArTo, that a developer can exploit to define and use custom numerical types in HLS code. This library supports both IEEE-754 and posit numerical formats, which are presented along with other numerical formats, arithmetic operator formalization and FPGA architecture in chapter 1. Using the MArTo allows fast evaluation of the performance/accuracy tradeoff entailed by the usage of a given numerical for a given application. MArTo is built on top of a second library, hint, that abstract operations on arbitrary size integers. Both of these libraries use C++ meta-programming constructs that are supported by HLS tools to implement the basic arithmetic operations of the supported formats. This design and the usage of the two libraries is further detailed in chapter 2, which also present a case study comparing the implementation cost of the two supported floating-point formats, IEEE-754 and posit.

A second optimization opportunities with regard to computation is the evaluation of mathematical functions. When only a limited set of mathematical functions is provided to the user (as it is the case with the C libm library), evaluating a composed function such as ↦ → cos(sin( )) requires composing evaluation of library-supported functions (computing sin( ) first, before evaluating cos on the result for the example function). However, ad-hoc function evaluators can be implemented to evaluate directly arbitrary expressions, allowing reduced latency and better control on precision compared to the composition approach. While generating such evaluators is a well studied problem, especially when both input and output format of the evaluator is fixed-point, this functionality is not supported by HLS tools.

The architecture of such evaluators require pre-computed values, such as coefficients of a polynomial approximating the function to evaluate. Obtaining these values is non-trivial, which prevent computing them in a pure metaprogrammatic fashion. Instead, it is more desirable to use existing tools that can provide state-of-the-art evaluation plans for arbitrary functions. However, evaluation plan has to be determined before performing the HLS compilation step.

Two methods to modify the HLS compilation flow are described in chapter 3. The first approach does not require direct modification of the HLS compiler. This makes it quite portable, but it requires that extra compilation steps are performed by the user, and it suffers from some fundamental limitations. The second approach consists in adding built-in support for arbitrary expression evaluation directly in the compiler. This method solves the issues of the first approach, but is not portable anymore. This disadvantage is attenuated by using a compiler-independent expression compiler that can be used by multiple HLS compilers.

Evaluation of arbitrary function is only one of the many possible optimized high-level primitives that could be made available to HLS developers, while providing automatic optimized implementation generation. Some perspective on building such "Real High-Level Synthesis" systems are discussed in the conclusion of this thesis.

Context

The central goal of the work presented in this thesis is to enable the development and the integration with software tools of high quality application specific arithmetic hardware operators. Here, hardware refers to implementation as digital circuits. This work focuses on digital circuits at the logic level. At this abstraction level, digital circuits consist in collections of logic signals combined to produce something useful. How exactly these signals are combined, and how to produce these combinations is detailed in section 1.4. The "something useful" that results from this combination depends on the application. In the case of this work, it consists in computing the result of some computation. This requires to have some way to interpret the logic signals of a digital circuit as a numerical value. This is the aim of numerical formats, which are presented in section 1.2. Having a representation for numbers allows to perform computation on them. This is the purpose of hardware arithmetic operators, presented in section 1.3. Finally, in the context of this thesis, hardware operators are implemented on Field-Programmable Gate Arrays (FPGAs), the architecture and programming model of which are presented in section 1.4. Before delving into these details, section 1.1 briefly presents the core concepts of logic signal and logic vector, and introduces the notations that are used in the rest of this work.

Logic signals, logic vectors

At the logic level, a logic signal is the core element that digital circuits handle. A logic signal is at anytime in one of two states, called either high and low, true or false, set and unset, or '1' and '0'. In this work, logic signals are modeled as variables from B = {0, 1} and are named with lowercase letters.

It is sometime convenient to reason on a whole group of signal as one en-tity. These signals can be concatenated to form logic vectors. In this work, logic vector names start with a capital letter.

Vector component access is denoted using squared bracket notation. For instance, [2] is the third component of the vector . The width of a logic vector is its number of elements, and is expressed in bits.

In this work, vectors will often be represented as bit string. With this representation, the vector first element is the rightmost one. That is, if = 01, then [0] = 1 and [1] = 0.

Numerical formats

Performing computation with digital circuits involves to manipulate numbers with them. As these circuits only manipulates logic vectors, it is necessary to be able to use them to represent numbers. This is the role of numerical formats, which define decoding schemes to interpret a logic vector, the representation as a value. Table 1.1 gives an example of an arbitrary numerical format for 2-bits vectors. A value is associated to each representation. All encoded values of table 1.1 example are real numbers. This is the canonical case, but some formats might also encode non-real special values, such as infinities.

Formally, a format F is defined by

• its representation width F (the size of its representations) (2 in case of the example of table 1.1),

• its (possibly empty) special value set S F . This is the set that contains the nonreal values that should be representable by the format such as infinities. The format domain is then

D F = R ∪ S F
This set is empty in the example.

• Its decoding function F : B F → D F that maps a representation to the represented value. In the case of the example, it corresponds to reading the value associated to the representation from the table.

• The set of representable domain values F F (the image of the representation domain through the application of the decoding function). This corresponds to the set of values appearing in the second column of table 1.1.

Due to the finite number of representations, F cannot represent exactly all real values. Rounding functions are used to map a representable value to any real value. Two basic directed rounding functions can be defined:

• Round towards +∞ (or rounding up): The F -representable neighborhood of a real value is defined as:

↑ F :D F → F F ↦ →        min ∈ F F | ≥ , if ∈ R otherwise • Round towards -∞ (or rounding down): ↓ F :D F → F F ↦ →        max ∈ F F | ≤ , if ∈ R otherwise
rnb F ( ) = ↓ F ( ), ↑ F ( )
When has a representation in F , this set only contains . As figure 1.1 shows, for the example format, rnb(5) = {2, 12}.

Numerical formats

This neighborhood is used to define two rounding modes that this work focus on. Rounding to nearest, as its name suggests returns the element of the neighborhood with the smallest distance to the exact value.

N F ( ) = argmin ∈rnb F ( ) | -|
In case of the illustration example, N (5) = 2 as the distance from 5 to 2 is smaller than the distance from 5 to 12.

A tie-breaking rule determines which of the neighbor should be returned when the real value is a midpoint, i.e. when the value is equidistant to the two neighbors.

The second rounding mode is faithful rounding. It only requires that one of the neighbors is returned.

Numerical formats with arbitrary mapping such as the example of table are difficult to reason with and to use when the representation size grows. For larger formats, it is thus desirable to get some hardware-exploitable regularity in the representation-value mapping. Commonly used regular formats are fixed-point and floating-point. These specific formats are detailed in the following sections.

Fixed-point

Unsigned fixed-point formats

Unsigned fixed-point formats values are represented using binary positional notation. Binary positional notation is very similar to the usual decimal notation: the only difference is the radix. A number written in this system is expressed as a sum of powers of two. The fractional point in this system is used to determine which power of to corresponds to which bit. The bit immediately on its left is associated to 2 0 and the associated weight doubles each step leftward. Reciprocally, it is halved each step rightward. A bit's position is its distance with the 2 0 weighted bit, with a positive increase in the left direction. The weight associated to such a bit is simply 2 . the axis below the representation, and their weight is written above them. For instance, the greyed bit is at position -1 and its weight is 2 -1 . All non-written bits are implicit zeros, so for instance bits with position above 6 are all zeros. The represented value is the weighted sum of its bits so here

2 7 7 2 6 6 2 5 5 2 4 4 2 3 3 2 2 2 2 1 1 2 0 0 2 -1 -1 2 -2 -2 2 -3 -3 2 -4 -4 2 -5
= 2 5 + 2 3 + 2 1 + 2 -1 + 2 -2 = 42.75
The value of a number written in binary positional notation with bits at position is

= +∞ =-∞ • 2
However, having an infinity of representation bits is not physically possible. Practical unsigned fixed-point numeric formats are defined by their high and low bit positions, respectively and , such that ≥ . Such a format is denoted ufix( , ), and has a representation width of ufix( , ) = -+ 1. Its representation is a "chunk" of the infinite positional notation, which is illustrated by figure 1 The weight of the rightmost bit of the representation is called the unit in the last place (ulp) of the format. This is the smallest difference that exists between two distinct representable values. Formats of the type ufix( , 0) form a notable group : they allow to represent all positive integers up to 2 +1 -1. By extension, a value from ufix( , ) can be thought as a scaled integer = • 2 with an integer representable on ufix( -, 0). The plots of representable values of two 5-bits fixed-point formats of Alternative equivalent parametrizations exist for fixed-point formats. One frequently used parametrization, the . notation, describes these formats with two integers and which are respectively the number of integer and fraction bits. For instance ufix(4, -2) would be denoted 5.2. While this notation is good when the decimal point is included in the described formats, it becomes quite unintuitive when it is not. For instance, the format ufix(-3, -7) is named -2, 5 in this formalism, the -2 integer bits conveying the fact that in addition to not having integer bits, the format also "lacks" the two leftmost fractional bits. For this reason, this work uses the ufix( , ) notation.

Next section present extensions to this system that allow handling negative values.

Signed fixed-point formats

Two main systems coexist to add support of negative values to fixed-point formats. These two systems have the same parameters and as unsigned fixed-point. In these two systems, the leftmost bit is the sign bit . The remaining low bits consist in a representation in ufix( -1, ). In the following, denotes these bits and the value that they represent. The decomposition of a signed fixed-point value is detailed on figure 1.5.

2 4 2 3 2 2 2 1 2 0 2 -1 2 -2 2 -3 2 -4
, ufix(4, -4) representation of Sign bit The two systems differs in how the sign bit is interpreted in the decoding function. In the most natural one, the sign-magnitude encoding, it is used to determined if the encoded value is (when the sign bit is '0') or its opposite -(when the sign bit is '1'). Formally, the represented value sm is given by the relation

sm = (-1) •
It is straightforward to see that the opposite of each representable value is also a representable value.

While it is quite intuitive, this encoding system has two main disadvantages. First, the algorithm to get the representation of the successor is quite complicated, as it depends on the sign bit. Indeed, in the general case (when both the value and its successor have the same sign), should be replaced by the representation of its successor when is unset, and by the representation of its predecessor when is set. And a few additional corner cases have to be handled. One of them is directly related to the second disadvantage, which is that these formats have two signed representations for 0.

The second system, two's complement, keeps the weighted-sum principle of unsigned fixed-point format. It only differs with this format by the weight of the leftmost (sign) bit, which is -2 instead of 2 . The represented value 2 is then 2 = -2 + With two's complement, the successor computation is greatly simplified and there is no redundant representation for one given value. Due to that and to the fact zero is still representable in the format, opposite of representable values is not always representable itself. This asymmetry is visible on Figure 1.6b which shows the representable value for a two's complement signed fixed-point format defined with = 1 and = -3. Value -2 is representable, while 2 is not. By contrast, representable values of sign-magnitude fixed-point with same parameters (represented on figure 1.6a) cannot represent -2 either and replace this by a second representation for 0. In the following, sfix( , ) will denote 2's complement fixed-point format of parameters and .

Similarly to the unsigned case, sfix( , 0) defines a notable group of formats able to represent all integers such that Numerical formats Values of sfix( , ) can also be seen as scaled integer:

-3 -2 -1 1 2 3 ±0 R (a) Sign-magnitude -3 -2 -1 0 1 2 3 R (b) Two's complement
= • 2
with representable on sfix( -, 0).

Biased fixed-point formats

It is possible to add an "implicit" constant offset (i.e. independent of the representation) to the summation performed to decode a fixed-point representation. This shifts the representable values by the added amount. One example of such biased fixed-point representations is the Half-Unit Biased (HUB) encoding families [START_REF] Hormigo | New formats for computing with real-numbers under round-to-nearest[END_REF]. A HUB representation is defined similarly to fixed-point formats by its high and low bit positions and . The Unit in the format name refers to the scaled integer vision of the fixed-point the represented number: HUB decoding function adds a constant offset of half of a scaled unit in the decoding process. The bias value is then 2 -1 .

It has the advantage that rounding a value to the nearest HUB representation only consists in truncated all bits of position lower than .

Another application of biased fixed-point is presented along IEEE-754 encoding scheme.

Arbitrary radix fixed-point

Using binary positional notation on digital circuits that manipulate binary signal seems natural. It is however possible to represent number using different radixes.

Radix notation decomposes a value as a sum of weighted -digits, each weight being a power of depending on the position of the -digit .

=

Similarly to binary fixed-point, the representation in radix fixed-point stores a sequence of -digits. As the circuit can only handle logic signal, these -digits are themselves encoded in binary.

When = 2 , the representation does not differ from radix 2 representation, only the interpreted meaning of representation bits differs. Indeed, in this case the contribution of each -digit can be seen as an integer scaled by a power of two, and none of these contributions overlap. Figure 1.7 illustrates this by showing two possible interpretation of the binary-coded representation of 42.75. When is not a power of two, the binary encoding of the -digit is wasteful. Indeed, = log 2 ( ) bits are required to store one -digit, but only out of the 2 possible combination of those bits are effectively used. This is illustrated with figure 1.8 that gives the binary-coded radix-10 representation of the value 42.75. This representation requires 15 bits, where the binary representation only requires 8 bits. So in one case the representation does not change, and in the other the representation is wasteful, so why using such a system ? If = 0, there is definitely no interest to use a radix that is not a power of two, as integers are all representable using radix 2. With < 0 however, the set of exactly representable values changes. For instance 49.3 is not exactly representable in binary, but is representable in radix 10. So financial application for instance might be interested not to have representation rounding errors when they perform computation over amount of money which are expressed in decimal notation.

2 7 2 6 2 5 2 4 2 3 2 2 2 1 2 0 2 -1 2 -2 2 -3 2 -4 2 -5 8 -1 8 0 8 1 2 2 2 1 2 0 2 2 2 1 2 0 2 2 2 1 2 0 … … 1 0 1 0 1 0 1 1
10 -2 2 3 2 2 2 1 2 0 10 -1 2 3 2 2 2 1 2 0 10 0 2 3 2 2 2 1 2 0 10 1 2 3 2 2 2 1 2 0 0 1 0 0 0 0 1 0 0 1 1 1 0 1 0 1 4 2 7 5
One way to limit the waste of representation bits is to use a radix that is very near to the power of two immediately superior to it. For instance, it might be more interesting to use radix 1000 (10 3 ) than radix 10 [2]. Indeed, with radix 1000, 1000 out of the 1024 possible combination of the 10 bits required to store a 1000-digit effectively used. In the other hand, only 10 out of the 16 possible combination of the 4 bits required to store a radix-10 digit are used.

Floating point

Fixed-point formats lack dynamics: representing both high magnitude and very small numbers is not possible without using very wide (and impractical) formats. Floating-point formats solve this issue by using a representation similar to the so-called scientific notation. A value is represented as a significand scaled by some power of a given radix .

= •

The discussion about radix value is analog to the fixed-point case. In the present work, only binary ( = 2) floating-point formats are considered.

A basic floating-point encoding could be defined by a pair of signed fixedpoint formats (F , F ), respectively the exponent and significand format. The representation is the concatenation of a representation and from each of the formats. The represented value is obtained from the two fixed-point values:

= F ( 6 ) • 2 F ( )
The representation width of F specifies with how many "significant digits" the format can represent a number. This is called the precision of the floating-point format. In the other hand, the representation width of F is related to the dynamic range of the format. For two floating-point formats with identical precision, the one with higher dynamic range will be able to represent higher and lower magnitude values. format. However, this format has redundancy. Indeed, if the significand representation does not start and end with a 1, it is possible to shift it and reduce or increase the exponent to keep the same value. For instance, 42.75 can also be represented as 010000010101011 in the format from figure 1.9. Most of the time this redundancy is undesirable, as it is "wasting" codes that could be used to represent other values. Following sections describes well-defined floating-point encoding schemes that avoid this waste.

IEEE-754 encoding scheme

Compared to fixed-point where possibility of variation is quite limited, floating-point offers a lot of freedom in their conception. In the early days of computing, this lead to each CPU manufacturer having its own specific format. This is reflected in the first version of ANSI C (C89) standard. In this version the definition of floating-point types in C is very scarce. It is only specified that C has three floating point type (float, double and long double), verifying the property

F float ⊆ F double ⊆ F long double
Outside from this, "the representations of floating-point types are unspecified". In addition to have different representation, the semantic of operations was also not consistent between floating-point implementations. As a result, standard compliant code could produce completely different outputs when executed on different machines. IEEE-754 [START_REF]ISO/IEC/IEEE International Standard -Floating-point arithmetic[END_REF] is the standard that resulted from the industry effort to eliminate this heterogeneity. Its first edition was issued in 1985 and defines two fully specified binary floating-point formats (which C99 later adopted as representation for the C float and double types), and the 'extended" format family. It also defines the base operations that should be provided by compliant implementations. The 2008 revision introduced the concept of interchange formats, that are intended solely for the purpose of sharing a value, but not for performing computation. It also specified three decimal floating-point formats.

Numerical formats

An IEEE-754 binary format representation is composed of three fields:

• The sign bit on 1 bit,

• the biased exponent on bits,

• the fraction part on bits. Field sizes for fully specified formats are given on table 1.2. The special value set of IEEE-754 binary encoding schemes contains four values. The first two, positive and negative infinities are the overflow markers. These are the values that are returned when an operation has a result which is bigger (resp. smaller) than the biggest (resp. smallest) representable values. The third and fourth ones are quiet and signalling Not-a-Number (NaN). They arise as a result of invalid operation, such as taking the logarithm of a negative number. The difference between signaling and quiet NaNs concerns the environment surrounding the computation, so for the rest of this work they will be considered as a unique value.

The special values are all encoded with all the exponent field bits set to one. When the fraction bits are set to zero, the represented value is infinity (with signedness depending on the sign bit). For instance, in binary16 the representation for -∞ is 1111110000000000. As soon as one bit of the fraction field differs from zero, the represented value is NaN. Since the 2008 standard revision, the leftmost bit of NaN representation indicates whether it is a quiet NaN. In binary16, quiet NaN are then represented as s111111xxxxxxxxx The remaining bits (marked with x in the previous example) are called the NaN payload and can be used to store application-defined data. This usage is simplified with the new operations setpayload and getPayload introduced in the 2019 standard revision.

When at least one bit of is set to 0, the represented value is a numerical value. The decoding process is very similar to what has been introduced in the previous section. A first difference is that it is using sign-magnitude encoding. The second important difference is that the fraction is normalized in order to avoid redundant value representations. That is, an implicit bit at position 0 is prepended to the fraction, the complete explicit significand is = . . In the general case, the value of is 1.

Finally, the third main difference is that a biased exponent is stored in the representation. For a binary format of exponent width , the bias value is

= 2 -1 -1
The actual exponent is -. Having a biased exponent ensure that the representation ordering is the same as if they were interpreted as sign-magnitude representation of integers. Indeed, the biased exponent is an unsigned integer to which is concatenated the fraction that is also an unsigned integer. For the same exponent, the representation with the biggest fraction represent the biggest value, and when the exponent differs, the representation with the biggest exponent encodes the biggest value. When the sign differs only the sign is relevant. This property allows efficient comparison of IEEE-754 values by using integer comparison on the representation. This also helps to compute rounding, as a rounding neighborhood is constituted by values having successive representations.

The represented value is then Figure 1.10 gives the encoding in binary16 of the value 42.75. The value stored in the exponent is 20. As the exponent field has a width of 5, the bias value is 15, so the unbiased exponent is 5. Stored fraction is 0101011000, so the significand is 1.0101011000 (1.3359375 in decimal). The represented value is then 1.3359375 • 2 5 = 42.75.

= (-1) • 1. • 2 - 0 1 0 1 0 0 0 1 0 1 0 1 1 0 0 0
Having always the implicit bit set to one does not allow representing 0. This is solved by adding a special case for numeric values called the subnormal range of the format. When all bits of are set to 0, the implicit bit becomes 0. To avoid a gap between subnormal and normalized values, the exponent that corresponds to subnormal significand is the same as the minimal normal exponent. So the real value decoding process involves computing the actual Numerical formats significand = .

and the unbiased exponent

= -+ (1 -) with =        0 if = 0 1 otherwise
to get the encoded value = (-1) • 2 • .

In addition to allowing the representation of zero, the subnormal mechanism also brings interesting properties to the format. One of which is the following equivalence: Property 1 Given an IEEE-754 binary numerical format, and a valid IEEE-754 rounding mode ∀( , ) ∈ F 2 , ( -) = 0 ⇔ = Property 1 means that the difference between two representable values cannot round to zero unless the two values are equal. This property can be exploited to optimize code (for instance replacing left member of equivalence by right member, which replaces computing one subtraction and one comparison by only one comparison). Figure 1.11 gives an illustration of why subnormal mechanism is required to get this property to work. Difference between two very near number that have exponent near to the minimal exponent can result in a value that is smaller than 2 min , with min the minimal normal exponent. Without subnormal mechanism, these values are not exactly representable, so depending on the rounding mode, the result could be zero even if the two values are not equal.

Table 1.3 summarizes the different decoding ranges for IEEE-754 binary16 format.

Operations

The IEEE-754 standard defines five rounding modes for elementary operations. The directed rounding modes towards +∞ and -∞ corresponds to the one defined in section 1.3, with the notable exception that values above the maximal representable numerical value are rounded to +∞ (resp. -∞). A third rounding mode, towards 0, is defined as follows:

1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 0 0 0 0 0 0 0 0 1 min - = 1.1111111 • 2 min = 1.1111110 • 2 min = 1.0000000 • 2 min -7 = 0.0000001 • 2 min
0 ( ) =        ↑ ( ) if < 0 ↓ ( ) otherwise
Round-to-nearest is also supported with two tie-breaking rules:

• Ties-to-even rule specifies that the neighbor with representation ending with a 0 bit should be returned. This allows an unbiased rounding error, as there are as many midpoints above their "nearest even" than midpoints below.

• Ties-to-away returns the neighbor with higher magnitude. This rule has "polarized" rounding error: negative midpoints are always rounded to lower values, where positive midpoint are always rounded to higher values. It is however a bit less expensive to compute, as only one extra bit after the last fraction bit of the result has to be known. In comparison, computing correct ties-to-even also requires knowing if the value is exactly a midpoint or not.

Compared to the fixed-point case, the distance between two successive representable values is not constant. Reasoning in terms of absolute error with floating-point is then quite complicated, so it makes more sense to use relative error. For a real value , and a rounding function , the relative error is

= ( ) - | |
Inside a binade, the set of all representable normal values that share the same exponent, the distance between successive values is constant. For round-to-nearest, the biggest absolute error is obtained for midpoints. On a binade of exponent in the normalized range, distance between representable values and midpoints is

= 2 -1 • 2
The minimal representable value inside the binade is = 2 . The relative error committed on the binade is then bounded by

| max | = + - + = 2 --1) 2 + 2 -( +1) < 2 -( +1) 2 = 2 --1
This maximal relative error does not depend on the binade. For directed rounding, the maximum absolute error inside a binade is twice the maximum absolute error of round-to-nearest, so the bound on relative error is also doubled. The same reasoning applies to each subnormal binade, but considering the effective fraction width of the binade (which takes into account the leading zero bits in the denormalized fraction). Figure 1.12 shows the relative error bound per binade when rounding to nearest IEEE-754 binary16 representation. The relative error is important for very small values (reaching 1 for values that are rounded to zero), and decreases until the normal range is reached, on which the error relative error bound is binade-independent. IEEE-754 binary32 and binary64 fulfill the role of the all-purpose numeric formats that are used when dealing with non-integral values of unknown scales. Even the high number of redundant representation for NaN has found usage, such as compact type information and value packing in dynamically typed language interpreter, with the technique known as NaN-boxing [START_REF] Suzuki | Building modern javascript engine[END_REF]. However, this redundancy is more annoying with narrower formats, where having more representable values would be preferable. Machine-learning is an example of application class that does not need an important precision but still requires values with an important dynamic range. Multiple encoding scheme have been developed recently to meet the requirement of efficient narrow floating-point representation.

DLFloat [START_REF] Agrawal | DLFloat: A 16-b floating point format designed for deep learning training and inference[END_REF] is such a format. Its 16-bits representations regular decoding process is those of an IEEE-754 format with = 9 and = 6. To reduce the overhead of special values, NaN and infinities are fused so the special value set of DLFloat only contains one element, NaN-or-Infinity. This element has

2 -29 2 -22 2 -15 2 -8 2 -1 2 6 2 13 2 20 2 -12 2 -10 2 -8 2 -6 2 -4 2 -2 2 0
Upper-bound for -N ( )

Figure 1.12: Bound on relative-error magnitude for IEEE-754 binary16 representable range.

two redundant representations, that are when the fraction and exponent bits are all set to one (the value of the sign bit is disregarded). All the other representation for NaN and infinity in the regular IEEE-754 scheme are interpreted as normal values, which increase the exponent range of one. The choice of having two redundant representation of one special value instead of having two distinct representation for infinity and NaN is done to simplify the special case handling. For the same objective, the detailed rule of IEEE-754 specifying which zero should be returned is ignored, and the two representations of zero are perfectly equivalent. Finally, subnormal value support is dropped, and the normal range is extended by one instead, which makes zero an irregular value regarding the encoding scheme.

Simpler NaN overhead reduction consists in taking a bigger exponent width for the same total representation width. As special value always reserve one exponent value, increasing the number of available exponents decrease the proportion of redundant NaN representations. This is the approach that BFloat16 [START_REF]BFloat16 -hardware numerics definition[END_REF] (a 16-bit format having an exponent width of 8-bits, similarly to IEEE-754 binary32) or MSFP8 and MSFP9 [START_REF] Chung | Serving DNNs in real time at datacenter scale with project brainwave[END_REF] (8 and 9-bits formats having the same exponent width of 5 bits than IEEE-754 binary16). Table 1.4 illustrates the trade-off that the three 16-bits formats binary16, BFloat and DLFloat offers in terms of representable range, precision, and pro-

Posit encoding scheme

While previous schemes are all variation over the IEEE-754 binary formats, completely different schemes also exist. This is the case of Posit encoding [START_REF]Standard for Posit arithmetic[END_REF], [START_REF]Posit standard documentation[END_REF]. Posit encodes floating-point value using a mixed thermometer/positional encoding of the exponent. A Posit format is parametrized by two natural integers, its width and its exponent shift field width . Regular positive Posit values have their sign bit , the leftmost representation bit, set to 0. A sequence of identical bits constitute the regime of the representation, with When < -1, the bit following the range acts as an end-of-regime marker. Then up to bits constitutes exponent shift, and all the remaining bits constitutes the fraction bits . The represented value is decoded using the following relation:

= 1. • 2 2 • 2
With being computed out of the regime:

=        -1 if [0] = 1 - if [0] = 0 Figure 1.
14 presents the decomposition of two positive Posit<8,2> representations. In the case of figure 1.14a, the range is constituted of 2 bits with the value 1, so the encoded value is

[1.10] 2 • 2 1•2 2 • 2 1 = 1.5 • 2 5 = 48
The regime can extend up to a point where it "pushes out of representation" and bits. In this case, pushed out bits are considered having a value of 0. This is the case on figure 1.14b. Here all the fraction bits and the first bit are pushed out of the representation by the regime. The regime is constituted of 5 bits with value 0 so the encoded value is

1.0 • 2 -5•2 2 • 2 0 = 1 • 2 -20
Only exception to this scheme are 0, which is encoded with all bit set to 0, and for Not a Real, the result of invalid operation, that is encoded with the sign bit set to 1 and all other bits set to 0.

The opposite of each representable posit value is also a representable value. Except for 0, the opposite is obtained by taking the two's complement of the representation.

Posit arithmetic is saturating: operation results that are greater in magnitude that the greatest representable posit value are rounded to this greatest value. Symmetrically, posit operations having a non-zero result with magnitude lower than the lowest non-zero representable magnitude are rounded away from zero.

2 -9 2 -7 2 -5 2 -3 2 -1 2 1 2 3
Upper-bound for -N ( )

Figure 1.15: Bound on relative-error magnitude for posit<16,2> representable range.

In the most recent version of the posit standard, which was before defined as log 2 ( ) for standard posit formats has been fixed to 2 for all the formats. Due to the varying length of the fraction encoding, the relative error of posit value is more binade-dependent than for IEEE binary encoding. Figure 1.15 gives the bound on the relative error for representable posit values of the standard posit<16, 2> format. It can be seen that the format gives a good relative precision for values around 1 (for posit<16,2> the relative precision is strictly better than IEEE binary 16 between 2 -4 and 2 3 ), but degrades when going to high magnitude exponent, similarly to what happens in the IEEE subnormal range. All exponent in the representable binade range cannot be properly represented (which corresponds to having the regime pushing the ES bits outside the representation), so in extreme cases the relative error can even be larger than 1.

Logarithmic Number System

All formats introduced until now are only able to represent rational values. This is not necessarily the case, as can be seen with Logarithmic Number System [START_REF] Swartzlander | The sign/logarithm number system[END_REF], [START_REF] Kingsbury | Digital filtering using logarithmic arithmetic[END_REF] (LNS). LNS keeps the idea from floating-point of encoding a number as a value scaled by a power of a given radix . However, the significand is always 1, and the exponent is allowed to be non-integral. In radix LNS, the representation of a value is thus a fixed-point number , such that = With this restricted definition, it is not possible to represent negative values. A sign bit is prepended to the representation to allow a sign-magnitude encoding of negative values. This still does not allow the representation of 0, so a special encoding is required for it if having it in the representable value set is a necessity. For instance, the minimal exponent value can be dedicated to it. . LNS has however the drawback of having almost no representable successive integers. For instance, binary LNS can represent 0, 1 and 2, but next representable integers are power of 2.

2 -3 2 -2 2 -1 2 0 2 1 2 2 2 3
The rationale behind LNS is that it simplifies some elementary operations computations. Indeed, the representation of the product of two values is just the sum of their representations, and the power function evaluation is reduced to the computation of a product. However, this comes to the price of a complex addition.

Indeed, for 0 and 1 two values such that 0 > 1 , represented in LNS by their logarithm 0 and 1 , the computation of the representation of = 0 + 1 is evaluated as follow:

log 2 ( ) = log 2 ( 0 + 1 ) = log 2 2 0 + 2 1 = log 2 2 1 • 1 + 2 0 -1 = 1 + ( 0 -1 ) with ( ) = log 2 (1 + 2 ) .
Contrary to fixed-point addition and product, no evaluation algorithm taking only the representation bits as input exists to evaluate . Its evaluation requires dedicated function approximation operators, which increases the cost of computing it. The problem is similar for subtraction, and amplifies when trying to compute exact sum of products.

Hardware arithmetic operators

The motivation of developing arithmetic digital circuits is to be able to perform computation. For instance, one can be interested in computing the euclidean distance between two points ( 0 , 0 ) and ( 1 , 1 ). This consists in evaluating the function defined by

: ( 0 , 0 , 1 , 1 ) ∈ R 4 ↦ → ( 1 -0 ) 2 + ( 1 -0 ) 2
Having numerical formats allows to represent the function inputs using bit vector. The purpose of hardware arithmetic operators is to compute the result of a given function applied on their inputs. Such an operator is obviously parametrized by the function it approximates, . In addition, the format of its inputs is given by its operand format list of, a -tuple of numeric formats. The size of the operand list matches the arity of the objective function. In addition, the destination format df is the numerical format according to which the operator output is encoded.

As an example, an operator can be defined to approximate the distance function , with all its inputs represented as 4-bits unsigned integers (corresponding to ufix(3, 0)), and having its output encoded as an sfix(4, -2). This operator would return 00101.00 (5) when applied to (4,0,1,4). A rounding relation ⊆ D df × F df specifying which values the operator is allowed to return when the result is not exactly representable is the last operator parameter. For instance, using the nearest neighboring as rounding relation for the example operator specifies that for the inputs (0, 0, 1, 1) it can only return 1.25 or 1.5, the representable values that surround the (not representable) exact value √

2.

The formal semantics of the operator is that given an input tuple , it returns the representation of a value ˜ ( ) such that ( ( ), ˜ ( )) ∈ .

For a given input value tuple , the difference between the operator result and exact result is ( ), the approximation error of the operator for this input.

Field-Programmable Gate Arrays

( ) = ( ) -˜ ( )
This work mostly focuses on the worst case absolute error of the operator

= max | ( )|
Once an arithmetic operator is fully specified, it can be implemented on real hardware. Implementing an operator consists in finding a succession of elementary operations that will transform inputs representations to output representation. These elementary operations depend on the type of circuit chosen for the implementation. For instance on CPU it consists in what is available through the instruction set. This thesis is focused on arithmetic on FPGAs, the next section describes this kind of device.

Field-Programmable Gate Arrays

The process of developing integrated-circuits comports some very high fixed costs related to the manufacturing process. In order to avoid paying these costs for dysfunctional circuits, lots of efforts is made to ensure circuit design correctness before starting the manufacturing process. Functional correctness can be partially tested through simulation on CPUs. However, it can be very slow for big design depending on the granularity of the simulation. Besides, this does not always allow to test the circuit system integration. This is to allow the fast simulation of arbitrary IC designs that Field-Programmable Gate Arrays (FPGAs) have been originally developed. Nowadays, they are also used as computation accelerators or used as network stream filters.

FPGA architecture

At a coarse grain FPGAs consist (as their name suggests) in a collection of reconfigurable logic blocks. Configuring the FPGA "programs" the logic blocks of the device to make it emulate a specific circuit.

Look-Up Tables

At a finer grain, FPGAs can be seen as a grid of interconnected programmable boolean function evaluator. The logic element at the core of boolean function evaluation are Look-Up Tables (LUT). They consists in very small memory blocks. For instance, on Xilinx UltraScale architectures [START_REF]Ultrascale architecture Configurable Logic Block[END_REF], LUTs are 2-bits memory addressed by a 5-bits word. These LUTs can be used to evaluate any boolean function of five inputs, with at most two outputs. Extra logic at the LUT output enable to choose one of the two outputs based on a sixth input. This setting allows the LUT to compute any boolean function of 6 inputs.

The competing Intel Stratix 10 architecture [START_REF]Intel Stratix 10 Logic Array Blocks and Adaptive Logic Modules user guide[END_REF] is based on 4 inputs LUTs that are grouped by 4 in what the architecture call an Adaptative Logic Module (ALM). In its widest configuration, an ALM can also be used as a 6-bits LUT of one output.

Table 1.5 show the (compressed) content of a 6-bits LUT that correspond to a 4:1 multiplexer. The output of the function is the value of the input signal , with the selection signal.

Table 1.5: LUT content for evaluating a 4:1 multiplexer.

3 2 1 0 0 0 x x x 0 0 0 0 x x x 1 1 0 1 x x 0 x 0 0 1 x x 1 x 1 1 0 x 0 x x 0 1 0 x 1 x x 1 1 1 0 x x x 0 1 1 1 x x x 1
Each bit of a vector returned by a function can be seen as being the result of an independent function. As a result, in the general case, it requires times more LUTs to compute a function returning a -bits vector than a function returning a single bit.

Function of more input bits can be constructed by splitting their truth table in multiple LUTs, each one adressed by the same input bits. The LUT result to keep is selected using a multiplexer controlled by the remaining input bits. The multiplexing can be built using LUTs parametrized with the truth table of table 1.5, or using fabric built-in multiplexers.

Routing

The role of the routing infrastructure is to allow connecting resulting signals to LUT inputs, multiplexers or other built-in resources. Routing basically consists in a wire grid surrounding logic elements. Multiplexers programmed during FPGA configuration select the endpoints for each routing wire. That way, arbitrary connection pattern can be implemented.

Number of LUTs available in recent high-end FPGAs exceed 10 6 . Having a dense connection graph between each LUT input and outputs is then im-Field-Programmable Gate Arrays possible. Besides, physical distance between LUTs has also to be taken into account, as the time needed for a signal to propagate along a wire that cross the entire die is way more important than to go from one LUT to its immediate neighbor. That is why routing is hierarchical, with low latency and dense connection between neighbor LUTs, and slower and sparser connection between distant LUTs.

LUTs are grouped to form coarser blocks (called respectively Logic Array Block and Configurable Logic Block in stratix 10 and UltraScale architectures). These blocks contain resources to combine efficiently their LUT outputs in useful ways. For instance, an UltraScale logic CLB slice contains all the multiplexers required to group its 8 LUTs into 4 functions of 7 bits, 2 function of 8 bits or one function of 9 bits.

In the general case, connecting a LUT output to the input of another LUT requires passing through routing, which is at least one order of magnitude slower than the LUT output computation latency. This is quite detrimental to computations involving dependency chains. This issue arises for instance when building adders.

Specialized primitives

A natural way to implement an adder of two -bits integers is to use LUTs configured as (1-bit) full adders. The corresponding truth table is reported in table 1.6. The input bits are connected to 0 and 1 . The output carry of the -th full-adder is used as input for the ( + 1)-th adder. The general scheme is reported on figure 1.17. With this setting, the carry propagation requires -1 "routing" steps, which dominates the overall latency of the computation.

For this specific case, UltraScale and Stratix 10 both offer fast carry lines. These lines allow fast propagation of a carry to the neighbor LUT. Schematic for UltraScale architecture reported on figure 1.18 illustrates the functioning of fast carry chain. Each LUT is configured as an half-adder (performing the addition of two bits, which corresponds to the truth table of table 1.6 keeping only the lines where = 0). In the UltraScale case, the result goes to the LUT output 6 and the output goes to 5. The carry propagation is done after LUT output. The fast result of LUT A will be xor-ed with the result of LUTB, computing the correct 1 value. And the value from LUTB controls the MUXCY multiplexer to send the correct carry signal to COUT. If the LUTB value is '0', then the input carry cannot propagate, so the LUTB is sent to COUT. When the value is '1' then the LUTB has necessary the value '0' and a carry will be propagated only if value is '1', that is why in this case the signal is propagated to COUT. As the neighbor LUTs are near to each other Table 1.6: Full-adder truth table.

0 1 0 0 0 0 0 0 0 1 1 0 0 1 0 1 0 0 1 1 0 1 1 0 0 1 0 1 0 1 0 1 1 1 0 0 1 1 1 1 1 1 -1 … 1 0 + -1 … 1 0 + 1 0 + 1 0 + 1 0 + 1 0 … +1 -1
. . . Field-Programmable Gate Arrays and the logic is simpler than full routing logic, fast carry propagation is faster than routing (and on UltraScale devices at least, has even smaller latency than LUT read). Fast-carry is an example of primitive added in the FPGA fabric to allow fast implementation of frequent operations (wide addition in this case). DSPs are another example of specialized blocks that provides efficient implementation of basic operations. They provide for instance fast small product computation (27x18 bits products in the case of UltraScale DSPs [START_REF]Ultrascale architecture dsp slice[END_REF], 27x27 for Stratix 10). Other provided operations includes pre-adding or bit-pattern comparison. Stratix 10 also offers hard operators for IEEE-754 binary32 floating point arithmetic.

In addition to computation primitives, memory primitives are also provided to avoid wasting routing to build large memories out of LUTs. Specialized LUT blocks to build distributed RAM or Block RAM are example of such primitives, with different size/latency trade-off.

Computing with FPGAs

While their original purpose was emulating arbitrary ASIC design, FPGAs are used nowadays as computation accelerators [START_REF] Caulfield | A cloud-scale acceleration architecture[END_REF], [START_REF] Sarkar | Fpga implementation of dna sequence alignment with traceback[END_REF]. They are well suited for processing data streams. The stream inputs can be read from FPGA interfaces or created out of a block of data retrieved from global memory. On CPU, the intermediate results of a computation are stored on general purpose registers, and spilled back to memory if not enough registers are available. In opposition, FPGAs can be seen as computation graph processor. One FPGA configuration corresponding to one computation graph. As long as the input of a primitive (LUT or DSP for instance) does not change, the intermediate result computed by this primitive is preserved locally. By using the registers present at logic block outputs, it is possible to build arbitrary deep computation pipeline. The critical path between two registers determines the clock frequency at which the FPGA design can run. The latency to process one element is the product of the clock period and the pipeline depth. If the stream is dense enough to keep the pipeline full, one stream element can be processed at each clock cycle, ensuring good throughput.

In addition, to reduce the latency, spatial parallelism can be exploited. On a single computing graph, all independent computations can run in parallel. Besides, multiple computation graphs can also coexist on the same configuration. They can be efficiently synchronized in a dataflow fashion: FIFOs are inserted at the interface between two computation graphs, and a graph computation starts only when all its input queues have elements ready to be processed.

From computation graph to FPGA configuration

The process of getting an FPGA configuration from a computation graph involves multiple steps. The first step, pipelining, requires specifying how to split the computation graph in multiple pipeline stages. This split will determine which computation nodes should be registered. The result of this stage is a register transfer level description of the computation graph.

The second stage of this compilation process is synthesis, which performs logic optimization before mapping arbitrary computation to FPGA primitives (e.g. LUTs and DSPs). The resulting is a logical netlist describing each required FPGA primitive, the connection graph between them and their configuration.

The final stages, place and route, determine the physical location for each of the primitive of the synthesis netlist and the routing configuration required to get the specified connectivity. At the end of this stages a physical netlist fully specifying the FPGA configuration is produced. This netlist can be used to produce an FPGA bitstream: the binary file that contains the device configuration.

Resources and timing analysis are performed on the Synthesis and Routed netlists to ensure that the computation graph fits on the device and determine the clock frequency at which the configuration can run. putation graph computing = ( + ) × . Choices made at an early stage have impact on the next stages. For instance, the register insertion point specified at the pipelining stage forces to + as an intermediate result. This hinders the usage of the DSP pre-adder, which would have otherwise allowed to fuse all the computation in one DSP Primitive. The task of pipelining is then quite complicated, as it should be aware of how the computation will be mapped to hardware primitives.

Computation graph descriptions

Hardware Description Languages

On the traditional FPGA configuration workflow, the pipelining step is done manually. The pipelined computation graph is given to the tool chain in a Hardware Description Language (HDL) such as Verilog [START_REF]Verilog hardware description language[END_REF] or VHDL [START_REF]VHDL language reference manual[END_REF]. List-ing 1.1 gives an example of a Verilog description of a circuit corresponding to the pipelined computation graph of figure 1.19b. A register is manually inserted to create a pipelined stage. As it can be seen on this example, traditional HDL description is at such a low-level that the programmer has to specify that registers are assigned on clock signal positive edges. Besides, the distinction between blocking and non-blocking assignment can be confusing. In addition, despite these languages being standardized, their support in commercial tools is incomplete, each toolchain supporting a different subset of it. Finally, traditional HDL are well suited to give a hardware description, but offer only very limited possibility to manipulate a description. For instance, it is not easy to get a generic "protocol" behavior, that would specify a generic interface, and add this behavior easily to existing components. This limits drastically the possibility of developing standalone generic reusable hardware description in these languages. The reusability and genericity is then obtained by developing generators in other languages that produce HDL description on demand.

Hardware Construction Languages

All these constraints hinder FPGAs adoption. To relax them, new generation of languages, hardware construction languages (HCL) have been developed. Usually, they consist in domain specific languages embedded in a general purpose host programming language. Developing hardware with these tools consists in writing a software application that builds a hardware description with the embedded language at runtime. Access to the host language allows the development of integrated component libraries, and higher level libraries that combine these components. Functions are provided to convert these representations to traditional HDLs. The most significant example of such tools is Chisel [START_REF] Bachrach | Chisel: Constructing hardware in a scala embedded language[END_REF], that uses Scala as host language. Listing 1.2 gives a description of the same circuit using chisel. The clock and reset signals are completely hidden in this description. Inserting a register is done using a dedicated class (here a register with initial value), and the custom assignment operator (:= symbol) is overloaded in a way that makes sense for register (update on clock rising edge and reset on reset). It is possible to get a handle on the clock signal if required, but the goal is to have primitives with good default behavior that avoid having to manipulate it directly. Line 13 shows how to print the Verilog corresponding to a PreAddProd.

High-Level Synthesis

While improving the usability by reducing the amount of boilerplate to describe components, and offering genericity, these recent HCL stay at a very low abstraction level. In order to broaden the audience even more, High Level Synthesis (HLS) tools have been developed. The aim of these tools is allowing programmers to give behavioral description of circuits as programs written in high level programming languages. While Modern HDL uses modern languages to generate circuit description, in HLS the program is the circuit description. Both AMD and Intel offer an HLS solution to program their FPGA, respectively Vitis HLS [START_REF]Vitis High-Level Synthesis user guide[END_REF] and Quartus HLS. Non-vendor alternatives exist both in academia, such as Bambu [START_REF] Ferrandi | Invited: Bambu: An open-source research framework for the high-level synthesis of complex applications[END_REF] or in the industry, such as LegUp (which originated from an academic project but is now supported by Microchip). Most of the tool use C++ as their source language. As with HDL, different subsets of the language are supported by the different tools. A first point of divergence is the target C++ standard version, which defines which syntactic constructs are supported. Apart from this, the main divergences is the level of support of the standard library constructs, especially regarding multi-thread and memory allocation related library functions. The work presented in this thesis mostly uses Vitis HLS, which officially supports C++14, without dynamic allocations and no standard library containers. Describing the computation graph that should be implemented on FPGAs directly as a language function allows its integration in more sophisticated libraries. One example of this is SYCL [START_REF]SYCL 2020 specification[END_REF], a C++ library that allows to program heterogeneous devices in a single-source fashion. Listing 1.4 gives an example of a SYCL application. The code of the lambda parameter of the submit method call will be compiled for FPGA. Invocation of this FPGA code from the CPU application and data synchronization between host and device memory is taken care of by the SYCL runtime. While the buffer programming model presented here is not always suitable for efficient FPGAs implementation, Intel has proposed an extension for a pipe based alternative programming model that is more fit to FPGAs [START_REF] Kinsner | Sycl intel dataflow pipes[END_REF].

}

It is possible for the user to specify constraints on the implementation directly from the code, by using custom pragmas or attributes. For instance, the latency of the function, the pipeline initiation interval, global limit on primitive usage or specific mapping of intermediate computation to defined primitives is controllable by pragmas. Some tools also allow the manual insertion of register using a special library type or a built-in function. This is however not the case anymore in recent Vitis HLS versions.

While HCL targets hardware developers wanting an improved productivity over traditional HDL flows, HLS tools target software engineers. The approach is more iterative and holistic. Full algorithm describing the intent of the circuit is first written. After functional verification, tuning can be performed by using pragmas. A big advantage of using a compiled software language for describing the behavior of the circuit is the possibility of very fast functional verification (by compiling the circuit code as software and running it on CPUs, with access to regular software debugging tools). Besides, co-simulation can be performed to check the correctness of the generated hardware.

High-Level Synthesis arithmetic support

C++ having initially been designed with CPUs as sole target, current standard version offers only basic integer type of power of 2 multiple of 8 bits up to 64 bits. While the introduction of the arbitrary width BitInt type of C23 [START_REF] Ballman | Adding a fundamental type for n-bit integers[END_REF] will probably find its counterpart in next C++ standard, current system provides their own library types to manipulate arbitrary width integer. For instance, the equivalent 12-bit circuit in Vitis HLS would be defined as in listing 1.5.

Careful choice of numerical formats for computation is a way to reduce area and latency by avoiding computing overprecise results. To this extent, HLS tools offer library types to use custom numeric formats, in addition to languages built-in types. The first class of these types are arbitrary width integers. The corresponding types are ac_int and ap_int for Quartus HLS and Vitis HLS respectively. These types are parametrized by a width and a boolean indicating whether they are signed. Signed integers negative values are represented in 2's complement. Parametrized fixed-point types are also provided.

Listing 1.5: HLS C++ description of a circuit computing ( + ) • with 12-bits inputs.

#include <ap_int.h>

using uint12_t = ap_uint<12>; ap_uint<25> preAddProd(uint12_t A, uint12_t B, uint12_t C) { return (A + B) * C;

}

For all these types, only basic arithmetic operations (addition, subtraction, product, division and modulo) are supported. In particular there is no support for usual transcendental functions evaluation. To get this functionality, either the user write its own arithmetic operator (which is a bit tedious if required for each computation and breaks completely the accessibility purpose of HLS) or interface with an externally provided operator.

High-Level Synthesis arithmetic support

HDL arithmetic core generators

Externally provided operators can be provided in the form of HDL description. Most of the time, these descriptions are produced by arithmetic cores generators. These tools take a description of the operator they should produce (either in text form or by providing a GUI), the hardware target for which to produce it, and output an HDL description of an implementation of the operator optimized for this target. An example of such generator is FloPoCo [START_REF] De Dinechin | Reflections on 10 years of FloPoCo[END_REF]. Listing 1.6 gives an example of the command line to produce a Fused Multiply Add operator (computing = • + with a single rounding) using FloPoCo. flopoco target=kintex7 IEEEFMA WE=8 WF=23 Listing 1.6: Flopoco command line to produce an IEEE-754 binary32 Fused Multiply Add operator for Kintex 7 architecture.

HDL files produced by a generator can then be integrated into an HLS program. The way to implement this integration is not standardized and depends on the toolchain. In general, the integration of an exogenous HDL IP inside an HLS program impacts the synthesis QoR. The external IP is already pipelined and can require that its inputs are registered. This has the effect that the IP is seen as a logic and pipelining optimization barrier for the HLS compiler, potentially increasing the latency and surface compared to an integrated HLS IP.

Toward on-demand HLS arithmetic operator implementation ?

To summarize, it is possible to program FPGAs at different level of abstraction.

• HDL allow to have fine-grain control over how things are implemented, but they are very verbose even for frequent operations. Their specification supports only very limited genericity.

• HCL provide a neat interface to generate HDL, allowing better expressivity and less verbosity while allowing almost the same fine-grain control than HLS, and offering more genericity by providing access to a general purpose programming language.

• HLS provides construct of a higher level of abstraction, enhancing the productivity of component writing. Describing a component as a host language program allows integration in a broader software environment, however in the current state HLS lacks some possibility of fine-grain tuning of im-plementation. One of such limitations is the lack of support of applicationspecific arithmetic primitives.

Can we have HLS with custom arithmetic operator support and fine grain control over their implementation ? This is the question to which this thesis aims at giving a (partial) answer. Chapter 2 will present a first approach to solve this issue with pure C++ library code. To support more complex usecases, chapter 3 will present alternative solutions that require compiler support.

A portable HLS-enabled library for custom numerical formats

This chapter introduces MArTo [START_REF] Uguen | Evaluating the hardware cost of the posit number system[END_REF], a proof-of-concept HLS library for basic operation support on custom numeric formats. This library implements product, subtraction and division operation for arbitrary fixed-point, IEEE-754 and posit formats. The library is written in pure C++, in order to be portable on various HLS toolchains.

The support of the elementary operations is based on arbitrary width integers arithmetic. Main HLS tools provide specific libraries for using this arithmetic. In order to be portable across toolchains, MArTo operators are built using a zero-cost abstraction layer for these libraries. This layer comes into the form of another library: hint [START_REF] Forget | A type-safe arbitrary precision arithmetic portability layer for HLS tools[END_REF]. It is introduced in section 2.1.

Hint, a portable abstraction layer for arbitrary width integer arithmetic

Arbitrary-sized integers are extremely useful when designing custom operators: for instance, IEEE-754 operators have 11-bit exponents and 52-bit fractions at the inputs and outputs. The adder data-path comports a 53-bit explicit fraction and a 56-bit data after effective addition. For multipliers, the intermediate mantissa product has a width of 106 bits. Being able to manipulates such integers is then an important feature for writing HLS-enabled arithmetic operators.

Integers and HLS

The support of arbitrary-sized bit vectors is not standardized among HLS tools. The nearest to a common standard is the ac_int templatised C++ library developed at Mentor Graphics [START_REF] Takach | Algorithm c (AC) datatypes[END_REF]. It is supported by the industrial tools Intel HLS and CatapultC, and the academic tool GAUT. However, the HLS tool with the most traction in reconfigurable computing is probably Xilinx Vitis HLS, and it uses a proprietary library called ap_int.

While ap_int and ac_int provide almost functionally equivalent support for basic arbitrary-sized signed and unsigned integers, their interfaces differ, and they do not have equivalent support for operations such as leading zero count.

Other tools only support widths up to 64-bits: the academic tool Augh [START_REF] Prost-Boucle | Fast and standalone design space exploration for high-level synthesis under resource constraints[END_REF] defines 64 new non-standard base types int1 to int64. Two other tools, LegUP and Panda/Bambu only support the standard C integer types so code must be written with only 8-, 16-, 32-, and 64-bit integers. If only standard-width types are available, a 17-bit integer must be emulated in the code using 32-bits numbers and bit-masks. These masked operations will hopefully be transformed to 17-bits integer operations during some optimization step. However, if this happened late in the process, the HLS schedule can be produced considering 32-bits operations, giving a pessimistic scheduling. Besides, this adds a lot of clutter in user code, and reduce its readability.

The necessity for the support of arbitrary width integer arithmetic slowly percolates in C/C++ standard, as it can be seen with the recent inclusion of the BitInt(n) types in C23. However, the semantics for these new types is underspecified. For instance, the result of overflowing operations on signed integer is "undefined behavior" in C/C++. This means that each implementation is free to return what it wants when such case arises.

Strong semantics

C++ is strongly typed, which means that each value has a type. Being able to handle bit-precise integer types is a great way to control the bus/register size of the corresponding hardware. While it is straightforward to control the type of declared variables, it is more complicated to define the exact meaning of compound expressions involving implicit intermediate types, and implicit type conversions.

For instance, the type of intermediate result (a+b) in the expression (a+b)+c is well specified in C/C++, but this specification is sometime nonintuitive and error-prone. If both a and b are unsigned 8-bit integers (declared with type uint8_t), the intermediate result (a+b) is implicitly promoted to the int type, that is implementation-defined but should have a width of at least 16 bits. If not optimized out by the compiler, this implicit conversion entails a waste of at least 7 bits. These rules have been established to limit the risk of accidental overflow, but they do not actually catch all cases. An example is given in listing 2.1 with operands having types of different signedness. In this example, C/C++ conversion rules specifies that b should be converted to uint32_t before performing the computation. That is why no sign extension is performed when the result is affected to a wider type. Mainstream tools such as VivadoHLS or CatapultC tend to choose the implicit intermediate types in a way that ensures that no information is lost. For instance, if both a and b are 32-bit integers, the implicit type of (a+b) should be a 33-bit integer to hold the possible carry out, unless the result of (a+b)+c is itself finally stored in a 32-bit integer, in which case all the arithmetic may happen modulo 2 32 . But this should nevertheless be specified.

Things are a bit trickier with shifts: left shifts may or may not lose the shifted out bits. Right shifts always loose the shifted-out bits, but in the signed case they may perform a sign extension, where the size of the intermediate format will matter. An example is the program from listing 2.2. With gcc 11.2 for x86_64 architecture, execution of the code from listing 2.2 prints 255 when compiled with the -O0 flag and 0 when compiled with the -O2 flag. This can be explained by different intermediate types for the intermediate result (a 64-bit int in the -O0 case, a 32-bit type in the -02 case). This actually conforms to the C++ standard, according which shifting a value of a number of bits bigger than the format width is undefined behavior. Such corner cases (C++ standard is full of undefined behavior) advocate for a library with consistent and completely specified semantics, in order to avoid users that are not aware of them to fall in this kind of trap.

Core arithmetic primitives for floating-point operators

Most floating-point operators (be it IEEE-754, posit-like, or other) rely on the following basic components:

• Arbitrary-precision addition, subtraction, multiplication. Multiplication can be implemented out of addition, but on reconfigurable targets it can also be built by assembling DSP blocks in a clever way. Therefore, multiplication should be a primitive, and its implementation is best left to back-end tools that know the target. Division and square root can be implemented either out of addition and tabulation, or out of multiplications. Whether or not this algorithmic choice should be left to the back-end tools is out of the scope of this thesis.

• Arbitrary precision shifters. There are standard operators in C/C++ for shift operations: << and >>. As we have already observed, it doesn't mean that their behavior is completely defined by the standard. In a processor, we usually have shift instructions that input the shift value and an integer, and output an integer of the same size (with possible loss of information). The C shift operators expose these instructions. Now if we are generating hardware, it is interesting to generalize as depicted in Figure 2.1: a shifter may be defined by an input width , a maximum shift distance , and an output width . The shift input will be an integer on log 2 bits. The shifter can be errorless (no shifted-out bits) if ≥ + . Besides, the "fill" bit of the shifter can also be specified (for instance to use the sign bit on a right shifter). Figure 2.9 gives an in-situ illustration of shifter usage.

• Arbitrary precision leading-bit counters. These operators count the number of successive leftmost bits sharing the same value. Leading-zero counters (LZC) and Leading-One counters (LOC) return a non-zero count only if the bits value is what they count (respectively '0' and '1'). For instance, LZC shifter by max. To reduce the data-path width and improve the delay, it is often useful to merge these operators:

• Addition with carry-in is in principle no more expensive than plain addition. Writing + + 1, or + + where one of the three variables is a single bit, should translate to a single adder.

• A shifter-sticky is a shifter whose output size is the same as the input size. Therefore, bits may be shifted out. The logical OR of shifted-out bits, historically called sticky-bit, is computed in parallel with the shift result. When computing a floating-point sum, fraction alignment is performed, by rightshifting fraction of the operand with lower exponent. The sticky bit keeps trace of those shifted-out bits to allow correct rounding. Compared to the naive alternative of computing a shift on a right-padded input before computing an OR-reduction on the output, performing the steps in parallel allows to reduce the operator latency and avoid the cost of the MUXes to shift bits that have no other aim than to be OR-reduced.

An in-situ example of shifter-sticky usage is visible on figure 2.7, page 68.

• A normalizer is a combined leading-zero counter and shifter. It is a leading-zero counter that, at the same time, shifts the input so that the leading bit of the output is the first non-zero bit. It outputs the normalized result, along with the number of zero bits that were counted. Posits make use of a similar combined leading-bit-counter and shifter, as illustrated on figure 2.5, page 65.

All the tools support C++ primitives (addition, multiplication, shifts), although the actual behavior does not necessarily follow C++ semantics. Sometimes, they are implemented as highly optimized IP cores. Sometimes, they are implemented as libraries.

The most notable missing basic operation is the leading bit count. Among the fused operations, only the add with carry-in is sometimes supported.

Type safety for arbitrary-precision integers in HLS

Current HLS integer libraries perform very little compile-time sanity checks. This section describes the Hint library variable declaration and its elementary methods along with their semantics and checks. These basic methods are used to build more complex operators. Code from listing 2.3 is an illustration of code using the hint library.

Variable declaration

The Hint library defines templated wrappers around underlying backends. A wrapper is a template class parametrized by an unsigned int describing the width of the represented integer type and a boolean defining whether it is signed. The library offers wrappers for the main arbitrary precision integer library: ap_int and ac_int for HLS, and GMP for software libraries. These wrapper template classes can be used as a template template parameter to build backend agnostic functions. The templated computation function of listing 2.3 illustrates this. It can for instance be instantiated with parameter of type VivadoWrapper<12, false> (a wrapper around ap_uint<12>), or IntelWrapper<12, false> (that wraps an ac_int<12, false>). Using C++20 concepts could simplify the syntax, but at the moment of writing this version of the standard is not supported by HLS toolchains.

Arithmetic operation result type and value assignment

The resulting type of any arithmetic operation is always the smallest type that ensures that no overflows can occur. For instance, the addition of two 6-bit integers at line 8 of listing 2.3 results in an 7 bit integer. This is similar to the choice made by industrial toolchains libraries. In these libraries however, assignment to a variable of mismatched size is a cause of silent truncation or padding. This can result in counter-intuitive behavior. For instance, Vitis HLS gives access to a built-in method __builtin_clz that takes an int as parameter and returns its leading-zero count. This function can be used on an ap_uint<24> without any error signaled by the compiler. However, as it is only defined for an int, the value will be silently extended to a 32-bit integer and the resulting leading-zero count will be too high of 8 compared to the expected result.

The hint library diverges from this approach by only allowing assignment of matching expression types. Without explicit truncation or padding, a compilation error fires when attempting to assign non-matching expression to a variable. In listing 2.3 for instance, if the return type of the function has been declared as Wrapper<5, false>, the return statement of line 9 would have been invalid, as the LZC on a 7-bit vector is a 3-bit unsigned value.

When porting to hint some arithmetical operators initially developed with ap_int, this strict semantics allowed the discovery of some place where internal computations were silent padding occured. While this issues did not affect the operator functionality, resizing correctly the relevant intermediate signals allowed to get a better default pipelining.

Slicing

Slicing consists in extracting a contiguous sub-vector from a bit vector. For instance, the slice 1101 can be extracted from the bits (3 to 6) of the vector 01101101. This operation is pervasive in arithmetic operators. It is for instance used to separate the different fields from an IEEE-754 representation.

Slicing methods for ac_int and ap_int have different restrictions, and may therefore not be interchangeable. In ac_int, a bit slice of size S starting as bit weight lsb of the variable var is written var.slc<S>(lsb). Conversely, the ap_int slicing method is var.range(msb,lsb) where msb and lsb are the weights of the MSB (most significant bit) and LSB (least significant bit) from which to slice var.

The difference is that the value of S must be known at compile time (more specifically, it should be a constexpr value). Therefore, a slice whose size varies in a loop will compile using ap_int but not using ac_int. However, the synthesized hardware bus size is fixed for the whole loop block (it makes no sense to "add dynamically" wires during the loop iterations1 ). So the choice made by ac_int seems more judicious as it allows compile-time check that the size of the slice is not bigger than the input value.

The approach followed in hint is even more restrictive, as the position of both limits of the slice needs to be constexpr values, in order to allow bound checks. The syntax is var.slice<msb,lsb>() where msb and lsb are the positions of the MSB and LSB of the slice. The usage of this method is visible on line 6 and 7 of listing 2.3.

Others operations

The Hint API can be extended with any other methods with the same spirit that all types must be checked at compile time. The provided operations are:

• concatenation, which join the bits from two variable, resulting in an integer having for width the sum of the two operand widths.

• bitwise operations such as and, or, xor that from two identical width variables returns a bits variable containing the corresponding bitwise operation.

• and/or reductions that returns a single bit result,

• a signal inverter,

• an operator that reverse the bit order of the variable (the MSB of the input becomes the LSB of the output for instance),

• a padding operator,

• a bit replicator,

• a comparison operators that only compares identical width and sign Hint variables,

• a multiplexer operator that takes a control bit and two identical width variables,

• explicit modular arithmetic operations,

• add with carry operation.

A concrete example of hint library usage is presented in listing A.2, which present the code of an IEEE-754 adder.

Software design of backend common interface

There are several ways to implement a common interface to the many backend arbitrary width integer libraries. Polymorphism through virtual inheritance is not an option as the library aims at adding no overhead compared to native library usage, so static polymorphism should be used. This section presents two possible approaches and their limitations.

Shared implicit interface

A first approach is to have independent classes for each backend. This class should implement the same shared interface. While being the more direct way to implement the library, this approach has two major drawbacks. The first one is that the static checks have to be duplicated in each backend class, thus being more prone to semantic divergence. The second drawback is that the common interface is not enforced by the compiler.

Curiously recurring template pattern

An alternative solution that solves these issue is by using the Curiously Recurring Template Pattern [START_REF] Coplien | Curiously recurring template patterns[END_REF] C++ idiom. This idiom enables static polymorphism by providing a template class that defines the public library interface as non-virtual methods. Backend classes inherit from the base class specialized for themselves. Interface (non-virtual) method delegates the computation to (also non virtual) derived class private implementation methods via casting this to a pointer on derived class. This allows the enforcement of a shared internal interface among backends (the base CRTP class cannot be specialized for a given backend if it does not provide the internal interface methods). In addition, the static checks can be factorized in the CRTP public interface code. This method uses no runtime polymorphism, so dispatching to the right backend method is done at compile time without the need of virtual tables. This approach worked correctly in software simulation for both Vitis HLS and Intel HLS. However, neither of the tools managed to pipeline correctly the code produced by this approach. Investigation tends to point the responsibility of an incomplete code inlining by the HLS compilers. Due to this issue, the first approach as been chosen for hint.

Cleaner interface with concepts and functional library

Two opportunities to improve hint interface have been identified. First, pure functions wrapping hint objects template methods. An example of such a wrapper is given in listing 2.4 for the slice method. This could improve greatly the readability of hint code by removing the need of prefixing template method usage with the template keyword in generic operator code. The second improvement comes from the concept mechanism introduced in C++20. Using this mechanism would allow expressing concisely the type requirement, and dropping the non-intuitive template template parameter. For instance the slice functional wrapper of listing 2.4 could be rewritten in the form presented in listing 2.5.

Overhead evaluation

The overhead evaluation is performed on the implementation of a leading zero counter (LZC). The lzc algorithm chosen in this work has been implemented using ac_int, ap_int and Hint. The synthesis results are given in Tables 2.1 and 2.2 for Intel and Xilinx respectively. The width (N) of the inputs corresponds to real world examples. Indeed, 26 and 55 bits are the widths of the LZC needed in single and double precision floating-point adders, while a 256-bits LZC is needed when dealing with posit32 exact product accumulators.

The comparison between the native type implementation and the hint type implementation shows no overhead when using Hint.

Combined operators

A combined shifter + sticky has been implemented. This operator computes a right shift, and a 1-bit signal, the sticky bit, indicating if any of the shifted-out bits had the value 0. It is useful when computing the addition of two floating point numbers. The shifter is used to align the fraction, and the sticky bit is required to compute correctly some rounding modes.

While a naive operation would first right-pad the input, shift the padded input and finally OR-reduce the slice that corresponds to shifted out bits. This doubles the size of the vector to shift, and is then quite costly. Instead, the or-reduction can be performed inline at each shifter stage.

The synthesis results of the shifter+sticky are presented in Tables 2.3 and 2.4. The width (N) of the operators comes once again from what is required in an IEEE-754 adder for binary32 (27 bits) and binary64 (56 bits). As both tables show, this optimization saves a considerable amount of logic on both targets. Table 2.3 does not report the number of cycles required for said operators. Indeed, IntelHLS was giving untrustworthy latency results. However, the circuits were co-simulated to ensure that they produced the correct mathematical results using ModelSim.

A second useful primitive is the normalizer, that combines a LZC and a left shifter. The semantic of this operator is to shift the input by its leading zero count, in order to bring its first '1' bit to the leftmost position. This is for instance useful to normalize the result of IEEE-754 products involving one subnormal number. This can be done naively by first computing the LZC, and then shifting by this amount. However, the LZC implementation is not far from computing successive shifts with masking, so it is possible to build an operator that performs the two operations simultaneously to reduce the computation latency compared to the naive implementation.

To evaluate such an operator, a combined hint lzc+shift is compared to a hint LZC followed by a native shift (>>). Tables 2.5 and 2.6 provide synthesis results of these implementations for both Intel and Xilinx FPGAs. In addition to sizes previously presented for the LZC.

For both Intel and Xilinx, the latency of the normalizer is improved compared to a serial LZC followed by shifter implementation. In some cases, the tools are even able to reduce the area of the operator.

Custom floating-point format library

The hint portability layer allowed to build MArTo [START_REF] Uguen | Evaluating the hardware cost of the posit number system[END_REF], a portable HLS library for custom-size IEEE-754 and posit arithmetic types support. MArTo allows the programmer to define custom types by specializing the template classes IEEENumber and PositNumber that are parametrized by the format dimension ( and for IEEE-754 formats, and for posit formats) and the hint backend used for the underlying integer operations. Once a type is defined, it can be used like a native numerical type.

Elementary operation support

On its current state, MArTo supports for basic operation is limited to addition, subtraction and product. The five IEEE-754 rouding modes are supported for IEEENumber operations. An example of code using the library is shown on listing 2.7. It is easy to replace the numerical format used in this code by changing the definition of my_fp_type. For IEEENumber, the default rounding mode when using overloaded +, -and * is round to nearest, ties to even. Example from listing 2.8 shows how to use a non-default rounding mode. In this example, the input type can only be an instance of IEEENumber, as PositNumber does not support alternative rounding modes. In addition to elementary operations, MArTo supports exact accumulation of products in wide fixed-point accumulators. This operation is detailed in the next section.

Exact fixed-point accumulation of floating-points products

Sum of product (or vector dot-product) is a pervasive operation in scientific computation. This is for instance the operation at the core of matrix product. This sum of product can be computed by using elementary addition and product. However, a rounding error is made at each elementary operations, and those errors accumulate. When the number of components of the vector is high, the resulting error can be important. This is even worse if this sum of product is part of the loop of an iterative algorithm.

An obvious solution to reduce the overall error in this case is to reduce the elementary rounding error by using floating-point formats with higher precision. However, this increases the memory required to store the elements, and the cost of elementary operation.

An alternative solution is to offer primitives that allows to compute such sum of product with only one final rounding. It is possible to provide allin-one sum-of-product accumulation operations for fixed number of inputs, such as the Fused Multiply Add (FMA) of IEEE-754 that computes FMA( , , ) = • + with only one rounding.

For arbitrary number of input, a solution initially proposed by Kulisch [START_REF] Kulisch | Computer arithmetic and validity: theory, implementation, and applications[END_REF] is to have fixed-point number wide enough to store exactly any floating-point product of the source format playing the role of accumulator. Each product to sum is computed without rounding (resulting in a floating point number with a fraction width twice bigger than original format, and exponent field having one bit more). The resulting fraction is added in the accumulator at the corresponding position (determined by the product exponent). When all products have been added, the accumulator stores the exact sum of the input products. The final operation is to round it back to the original format.

2 4 2 5 2 6 2 7 2 8 2 0 2 1 2 2 2 3 2 4 2 -4 2 -3 2 -2 2 -1 2 0 … … 0 0 1 1 0 1 1 0 1 1 0 0 1 1 0 0 1 0 1 1 1 1 0 0 1 1 1 1 1 0 0 1 0 1 0 1 0 1 0 1 1 1 0 1 1 1 0
In a basic implementation where the accumulator is stored as one bit vector, addition of a product is an expensive operation due to the delay required for carry propagation from the lowest to the highest bit. In order to reduce this latency, the accumulator can be split in multiple banks in carry-save fashion. With this setting, there is an overlap bit between two adjacent banks. When a product is added to the accumulator, the carry bit from the lower bank is also added, and the resulting carry bit is stored in current bank carryout bit waiting next summation to propagate. Figure 2.3 gives an example of one accumulation step. On this figure, the carry saved for each bank is highlighted with a gray background. At each accumulation step, bank-wise additions can be performed in parallel, as the carry bit is stored locally. In this setting, the cost of complete carry propagation is only paid when rounding the accumulator back to the floating-point format.

Posit standard requires the support of accumulator operations. In the posit world, these accumulators are called quires. The 2019 version of the IEEE-754 standard recommend support for the dot operation, which computes an approximation of the dot product of two floating-point vectors. However, it leaves the accuracy implementation defined.

MArTo defines several types to represent these wide accumulators in plain integer or carry save formats, overload the operator to allow convenient syntax for accumulation of custom type products, and provides the rounding operator to convert these accumulators to the elementary format they accumu-Custom floating-point format library late products of.

Operator implementations in MArTo Implementation of IEEE-754 elementary operations

MArTo IEEE-754 elementary operation implementations are quite standard, so they will not be described in great details. The adder has a single-path architecture (which had better QoR on tested targets than dual path implementation at the time of development). The only "subtlety" is the use of fused operators. For instance, the adder uses a combined shifter + sticky operator for fraction alignment. Architectural choices have been inspired by the corresponding floating-point architectures produced by FloPoCo [START_REF] De Dinechin | Reflections on 10 years of FloPoCo[END_REF]. The interested reader is invited to check [START_REF] De Dinechin | Application-Specific Arithmetic[END_REF] to find detailed discussion on possible IEEE-754 operator implementation. To get an idea of the overall design, the generic adder code is reported in listing A.2. The fused shifter and or reduction is visible on line 95 of this listing. This code is also interesting to look at in order to see the relative amount of work of the different computation "steps". Indeed, while they occupy an important number of line of code, special case handling and rounding logic actually consists in a few or/and reduction and boolean combination of these reduction results. These operations are not very expensive in regard to the large shifts required to align the fraction and renormalize the result when required.

On the other hand, there is less literature about Posit operator. That is the reason why their MArTo implementation is further detailed in the following sections.

Posit smallest floating-point superset

The general idea behind posit operator implementation is to convert the posit encoding to a fixed-width fields floating-point format, and perform the operation on this representation.

The sizes of posit value exponent and fraction depend on the regime size. However, the hardware bit-widths cannot change dynamically, they have to be designed for the worst case widths. This worst case therefore corresponds to the smallest fixed-precision floating-point format that is a superset of posit numbers. To match posit encoding, this floating-point format only includes normal numbers with fractions expressed in two's complement.

The parameters for the smallest fixed precision floating-point superset are derived as follows. The size of a two's complement fraction is the maximum precision of a posit value, obtained for the minimum length = 2 of the The maximal exponent is obtained when the regime length is -1. In this case, all the and bits are pushed out by the regime. Hence, the maximum exponent value is

max = ( -2)2 (2.1)
As the opposite exponent can also be reached, the number of bits needed to store the exponent is

= 1 + log 2 ( -2)2 = 1 + + log 2 ( -2) (2.2)
The parameter allows trading between the range of the format and its maximal precision.

Table 2.7 gives, for each of the standard posit formats, the exponent and fraction sizes of the smallest floating-point superset.

The hardware-friendly posit Intermediate Format

With this smallest posit FP superset, it becomes possible to define a new intermediate encoding for numbers in this set that is more adapted to hardware arithmetic operations, in the sense that all its fields have a fixed width. In this work, this encoding scheme is denoted posit Intermediate Format or PIF.

Figure 2.4 highlights the role of this format in an end-to-end posit arithmetic operator. Posits are first decoded to PIF. As PIF can represent all posit values, this operation is exact. Then, the arithmetic operation is performed on PIF data. Finally, the result is encoded back to posit. Since rounding (to an exponent-dependent position) must be performed in this encoding step, the output format of the PIF operation must be an Unrounded PIF, which is a PIF extended with all the additional information needed for standard-compliant rounding, detailed below in Section 2.2.3. We believe the 3-step approach of Figure 2.4 is inevitable for stand-alone posit operators (except for very small formats where a simple tabulation may be used). It is followed (more or less explicitly) by leading hardware posit implementations [START_REF] Podobas | Hardware implementation of POSITs and their application in FPGAs[END_REF]- [START_REF] Jaiswal | PACoGen: A hardware posit arithmetic core generator[END_REF].

The PIF should be designed with two objectives in mind:

• Posit to PIF conversion should be as simple as possible,

• Arithmetic operations should be efficiently computed on this representation.

Because of the second objective, PIF is a simple normalized floating-point representation that uses the parameters of Table 2.7. To address the first objective, the proposed PIF encodes both the exponent and significand in two's complement (where IEEE-754 uses a biased encoding for exponent and a signmagnitude representation for the significand). This avoids two's complement to sign-magnitude conversions (which may cost a carry propagation). It also has the side effect of slightly simplifying PIF addition of values with opposite signs.

Two's complement encoding for a normalized significand consists of a sign bit and a fraction on bits. The two's complement significand value is then = -2 + ¯ . .

However, (2.3) does not allow the encoding of zero. The proposed PIF encoding scheme introduces an extra bit which is one if and only if the represented value is strictly positive. The significand value becomes

= -2 + . . (2.4)
Zero is the only posit value whose PIF representation has both and set to zero. This enables efficient zero detection in arithmetic operators. For non zero values, exactly one of or is set.

PIF also has an extra isNaR bit, set to one if and only if the represented value is NaR. An alternative option could have been to use a non-posit value, for instance with both and set to one. This would trade one bit of representation for a few gates of encoding/decoding logic.

To summarize, the PIF encoding scheme is composed of the following fields:

• a isNaR flag,

• the sign bit ,

• the exponent stored in two's complement on bits,

• the weight one bit ,

• the fraction bits on bits.

The encoded value is

=        NaR if isNaR is one (-2 + + 0. ) × 2 otherwise (2.5)

Unrounded PIF encoding of the result of basic operations

In the general case, the result of an operation on two PIF values is not exactly representable as a PIF, and must be rounded. As PIF is a floating-point format, we may use textbook techniques [START_REF] Ercegovac | Digital Arithmetic[END_REF], [START_REF] Muller | Handbook of Floating-Point Arithmetic[END_REF] for this. For the basic operations (addition, multiplication, division and square root) the exact result can always be represented on at most 2 pif bits, then for the purpose of rounding the extra pif bits can be condensed into only two bits:

• an extra fraction bit at the LSB, called the round bit;

• a sticky bit, set if and only if the exact value is strictly greater than what is represented by the fraction extended with the round bit (but still smaller than the next representable value). In other words, a sticky bit of zero means that the value represented by the extended fraction is exact.

We define the UPIF (Unrounded PIF) format as a PIF with these two extra bits.

The floating-point literature often uses a third additional bit (called the guard bit), useful in the case when a 1-bit normalization of the significand may be needed. In the big picture of Figures 2.4, the PIF operator is in charge of this normalization, so no guard bit is needed in UPIF.

This work only demonstrates the use of the UPIF format on addition/subtraction and multiplication, but it is equally suitable for division and square root. Digit recurrence algorithms [START_REF] Ercegovac | Digital Arithmetic[END_REF] compute a remainder along with the quotient or square root, out of which the round and sticky bits can be computed. Multiplication-based algorithms [START_REF] Ercegovac | Digital Arithmetic[END_REF], [START_REF] Muller | Handbook of Floating-Point Arithmetic[END_REF] also can output their result in UPIF format -for instance by computing the remainder.

Table 2.7 gives the width for PIF and UPIF associated with standard posit formats.

Saturation management

Posit arithmetic does not offer an overflow detection mechanism to the user. When the exact result of an operation is bigger than the biggest representable value, this biggest representable value is returned.

This saturation could in principle be handled in a generic way in the "Encode to posit and round" block of Figure 2.4. However, as each operation leads to different overflow situations, it is more efficient to manage saturation in each PIF operator. The UPIF specification exposed previously actually assumes that saturation management has been performed by the operator, otherwise extra bits would have been needed. Another advantage is that some saturation situations may be detected in parallel with computation, thus reducing latency.

Posit to PIF decoder

The proposed posit decoder is depicted in Figure 2.5.

The "LZOC + Shift" block (LZOC stands for "leading zero/one counter") counts the range bits while discarding them, resulting in a normalized fraction.

The most significant exponent bits ℎ are computed out of the range count. If the leading bit is equal to , then ℎ = -(= ¯ +1); else ℎ = -1. An optimization is to skip the first range bit when counting, effectively computing = -1. Indeed, if the first range bit is equal to , ℎ = + 1 + 1 = ¯ , or ℎ = otherwise. This high bit decoding method improves the state of the art by avoiding an adder to compute -. The exponent least significant bits are obtained by xoring with the first bits of the aligned fraction. The PIF exponent is the concatenation of ℎ and (or is equal to ℎ if the corresponding format has = 0). The decoder is slightly simplified with = 0 posit formats, as it saves the XOR gates labeled * on Figure 2.5. The PIF fraction consists of the least significant bits of the aligned fraction. An OR reduction over the -1 rightmost bits of the posit input is used to detect both zero and NaR values, in conjunction with .

The weight 0 significand bit is computed out of and the detection of zero value.

The most expensive parts of this architecture are the "OR reduce" over -1 bits to detect NaR numbers, and the combined leading zero/one counter and shifter.

positN

LZOC + Shift / -1 OR reduce / -1 / -2 ℎ * / -3 / / / isNaR

UPIF to posit and PIF to posit

The complete UPIF to posit encoder architecture is shown in Figure 2.6.

First, the rounding bit is appended to the fraction, and the range is computed then prepended. The Shifter+Sticky component then simultaneously right-shifts this word and ORs the shifted-out bits into a unique sticky bit, Custom floating-point format library which is then ORed to the PIF sticky bit to get the final sticky bit.

Finally a round-up bit is computed (the right AND of Figure 2.6 implementing round to nearest, and the left AND the "ties to even" rule), and added to the final encoding.

The case = 0 requires additional logic (not shown in Figure 2.6) to detect and forbid a special case of rounding up that would cause the output to round to NaR or 0.

PIF floating-point operations

The architectures of the PIF adder/subtracter (Figure 2. Although the adder is a single-path architecture [START_REF] Muller | Handbook of Floating-Point Arithmetic[END_REF], its datapath can be minimized thanks to the classical observation that large shifts in the two shifters are mutually exclusive. Indeed, the normalizing LZOC+Shift of Figure 2.7 will only perform a large shift in a cancellation situation, but such a situation may only occur when the absolute exponent difference is smaller than 1, which means that the first shift was a very small one. Conversely, when the first shifter performs a large shift, the rightmost part of the significand can be immediately compressed into a sticky bit, since we know that it will not be shifted back by the second LZOC+Shift. All this allows us to keep most intermediate signals on

+ 2 to + 6 bits, where previous work [START_REF] Podobas | Hardware implementation of POSITs and their application in FPGAs[END_REF], [START_REF] Chaurasiya | Parameterized posit arithmetic hardware generator[END_REF] seem to use datapaths that are twice as large.

The posit multiplier shown in Figure 2.8 is straightforward. It performs the addition of the exponent, the signed product of the significands, then if necessary normalises the result (one-bit fraction shift and exponent update). This architecture aims at minimal area and delay. If energy efficiency is the goal, alternative architecture have been proposed that exploit the relation between exponent magnitude and precision to disable the computation of unneeded product LSBs [START_REF] Zhang | Design of power efficient posit multiplier[END_REF].

The bottom part of Figures 2.7 and 2.8 normalize the exact result computed by the top parts to a PIF. For both operators, the exact significand must be realigned, correcting the exponent accordingly.

Hardware support for exact accumulation

Several existing machine learning accelerators [START_REF] Brunie | Modified Fused Multiply and Add for exact low precision product accumulation[END_REF], [START_REF] Johnson | Rethinking floating point for deep learning[END_REF] already use variations of the exact accumulator to compute on IEEE-754 16-bit floating-point. Other application-specific uses have been suggested [START_REF] De Dinechin | An FPGA-specific approach to floatingpoint accumulation and sum-of-products[END_REF], [START_REF] Uguen | Application-specific arithmetic in high-level synthesis tools[END_REF]. For larger sizes, this could be a useful instance of "dark silicon" [START_REF] Taylor | Is dark silicon useful? harnessing the four horsemen of the coming dark silicon apocalypse[END_REF]. The posit standard, [START_REF]Standard for Posit arithmetic[END_REF], [START_REF]Posit standard documentation[END_REF] defines fused operations as those expressible as sums and differences of the exact product of two posits; no other fused operations are allowed. It also specifies a binary interchange format, the quire, which consists in a fixed point number of size defined by

= max + -min + 1
With max and min the maximal and minimal possible exponent for a given posit format product, and the number of extra-bits used to catch overflows.

In the sequel, we discuss the cost of hardware support for accumulating number in a quire. For draft standard formats, is set such that

= 2 2
As the parameter is a storage requirement, it defines a lower bound of the area cost. Figure 2.9 shows a high-level functional description of a quire accumulation, and shows that there is a -bit addition on the critical path from the quire to itself, which would also entail a large cycle delay. A technique that relaxes this constraint is reviewed in Section 2.2.3.

The posit standards [START_REF]Standard for Posit arithmetic[END_REF], [START_REF]Posit standard documentation[END_REF] specifies NaR as a special quire value. This means that this special value must be tested at each new quire operation. Instead, we add to the internal quire an isNaR flag bit, set when a NaR is added to the quire, and sticky until the quire is cleared. This isNaR bit can be encoded and decoded only when transferring quire to/from memory, however we suggest that it could even replace one of the quire carry bits in the interchange format.

In the posit context, it is natural to use two's complement for managing signs in the quire. Some implementations of Kulisch's exact accumulator seem to use a sign-magnitude representation for the accumulator [START_REF] Kulisch | Advanced Arithmetic for the Digital Computer: Design of Arithmetic Units[END_REF], matching the sign-magnitude representation of IEEE floating-point. However, a two's complement representation of the accumulator is more efficient even in the IEEE-754 context [START_REF] Brunie | Modified Fused Multiply and Add for exact low precision product accumulation[END_REF], [START_REF] Uguen | Design-space exploration for the Kulisch accumulator[END_REF].

Addition of products to the quire

The posit quire is able to perform exact sums and sums of products. Therefore, the input format of the quire is defined as the output of the exact multiplier from the top half of figure 2.8 (before the rounding logic).

To add a simple posit to the quire, it is first converted to PIF, then the PIF value is trivially extended to the exact product format.

The simplest implementation of the quire addition/subtraction is depicted in Figure 2.9. An exact posit product significand is shifted to the correct place to the quire format according to its exponent. A large adder then performs the addition with the previous quire value. The two's complement subtraction is performed at the cost of a XOR on the input and a carry-in to the adder, as in the posit adder/subtracter. For this the shifter must be a sign-extending one.

The simple architecture of Figure 2.9 can be used directly for small sizes (up to posit16). For larger sizes, the long carry propagation delay of the addition in this architecture will restrict the maximum frequency achievable. To address this, a cost-effective solution [START_REF] De Dinechin | An FPGA-specific approach to floatingpoint accumulation and sum-of-products[END_REF], [START_REF] Uguen | Design-space exploration for the Kulisch accumulator[END_REF] is to segment the quire into smaller words (typically standard 32-bit or 64-bit words). Carry propagation is then limited to a segment, and the carries between segments are stored in registers and propagated to the next segment during the next cycle. Another point of view is that the quire is kept in a high-radix carry-save redundant format (radix is 2 32 or 2 64 ). If such a format is used, its conversion to a nonredundant format will incur additional overhead to complete carry propagation. A cheap way of achieving this propagation is simply to dedicate a few Custom floating-point format library cycles to the completion of the carry propagation, during which the summand input to the quire is kept at zero. The number of carry-propagation cycles is /32 for 32-bit segments. These extra cycles are amortized if the quire is used for summing large numbers of values. The conversion of the quire value to a posit is divided in two steps. The quire is first converted to a UPIF value before the latter is encoded to a posit.

Conversion from quire to posit

There are four distinct cases to take into account when converting the quire to the UPIF:

• If the quire holds a NaR value, the result is NaR;

• If the quire value is larger in magnitude than the maximum-magnitude posit (overflow), the latter should be returned (saturation);

• If the quire value belongs to the representable posit range, it should be converted;

• If the quire value is smaller in magnitude than the minimum-magnitude non-zero posit (underflow), the latter should be returned (saturation); Figure 2.10 illustrates the different interesting zones of a quire. Detection of overflow consists in comparing all the bit in the overflow zone with the sign bit. If at least one differs, the posit overflows. The width of the overflow zone is given by:

= msb -max = max +
For quire values inside the posit range, a normalization should be performed, which uses a LZOC + shifter of input width , with

= max -min + 1
Finally, to detect the difference between an underflow value and a real zero, a wide OR is performed on the underflow zone of width , with

= min -lsb = max
For draft standard posit sizes, the simplified formulas are the following:

= 2 8 + 3 4 -2 = 2 8 - 4 = 2 4 - 2 + 1
Associated numerical values reported in Table 2.8.

Verification of functionnal correctness of the operators

In order to verify that the proposed architectures are correct, the following functional tests were first run in sofware:

• Exhaustive test against softposit of posit8 and posit16 addition and multiplication.

• Some corner case tests of quire addition/subtraction and conversion back to posit for posit16.

• Exhaustive test for addition/product of IEEE16 against SoftFloat, for the five IEEE-754 rounding modes.

Then the VHDL file produced by the Vivado HLS compiler for a 16-bit posit adder was exhaustively tested using a VHDL simulator. Scripts and sources to reproduce this experiment are accessible from the MArTo repository.

Finally, the standard posit16 multiplier was synthesized, placed and routed for the Zynq FPGA of a Zybo board using the Vivado HLS toolchain, and the resulting FPGA circuit was exhaustively checked against SoftPosit executed on the ARM core of the Zynq. All these tests were successful.

Case study: comparing IEEE-754 and posit hardware implementation cost

Many works have compared the accuracies of posit and IEEE-754 floatingpoint formats [START_REF] Carmichael | Performance-efficiency trade-off of low-precision numerical formats in deep neural networks[END_REF]- [START_REF] Buoncristiani | Evaluating the numerical stability of posit arithmetic[END_REF]. However, these works didn't evaluate the performance implication of replacing IEEE-754 by posit. This is the aim of this case study, which will evaluate two possible architecture for posit arithmetic units. The first one is the one described on figure 2.4, with posit used both as memory and internal register format. As shown on the figure however, this mostly adds a decoding step overhead to a classical floating-point operation (the PIF have no subnormal to support that simplify a bit these operators compared to IEEE-754, but the posit encoding is actually very similar to subnormal handling). As an alternative, posit can be used as memory format only, decoded on the memory path, and PIF as register format. The block-diagram of this architecture is presented on figure 2.11.

In order to avoid double-rouding and related issues, PIF in register should always represent the rounded result of the operation. This is the reason of the "UPIF inplace round" block of Figure 2.11b shows an UPIF in place round block. This block is roughly equivalent to the UPIF to posit encoder of figure 2.6, with the difference that a mask is shifted instead of the representation to. PIF to posit conversion that occurs when writing the value back to memory (as shown on figure 2.11b) is a simplification of the UPIF to posit encoder. Indeed, PIF values stores exactly representable posit values, so all the rounding logic can be removed.

Comparison of operator area and latency

Qualitative comparison

A first "qualitative" comparison between posit and IEEE-754 operators can be performed at the architecture level. There is no 1-to-1 match to the decoder step in IEEE-754 arithmetic operator implementation, as the representation already has fixed-width fields. However, it can be compared with the detection of special cases. On the IEEE-754 side, one OR reduction on the exponent bits is needed to detect subnormals, another OR reduction on the significand bits is needed to detect zeros, and two similar AND reductions are needed to detect respectively NaN and infinities. The two OR reduction have a similar input width in total to the posit decoder OR reduction block. Normalization of subnormal values requires a LZC and a shifter, that operate on the fraction bits. On the posit side, the combined LZOC + shifter operates on a wider value. Besides, the fact that it is an LZOC and not an LZC inclurs internal AND reductions that matches the cost of the IEEE-754 ones. All in all, the posit decoding process cost is expected to be very similar to the handling of special cases and subnormal values in IEEE-754.

The PIF arithmetic operator implementations are very similar to IEEE-754 operators, so area and latency should be quite similar.

Handling rounding and special cases encoding in IEEE-754 requires one integer addition of the representation width, a few gates to determine if the rounding should occur according to the rounding mode, and a few bitwise masking operations to handle the generation of NaN/infinities if required. By contrast, posit rounding and encoding process which requires the same addition, but also a potentially large shift seems more costly.

As a whole, it is expected that posit operators require more resources and have a higher latency compared to IEEE-754 operators for identical widths formats. The expected overhead coming mostly from an expensive rounding and encoding process.

Comparison with state of the art

As we eventually observe that posits are larger and slower than IEEE floats, it is important to be convincing that we are not using a substandard posit implementation. For this purpose, Tables 2.9, 2.10, 2.11 and 2.12 gather the results of best-effort comparisons with the state of the art in posit hardware at time of writing. It shows that MArTo definitely improved this state of the art.

There is less pressure to show that the MArTo implementation of IEEE floats is efficient. A comparison with Xilinx implementation of IEEE floats is provided in Table 2.14. There, the line labeled Xilinx Float corresponds to IP used by Vivado HLS when using the float and double data types in the HLS C++ (hence the lack of 16-bit results). This hard IP is the industry standard when using Vivado, and can be considered a state-of-the-art implementation IEEE-754 vs posit hardware cost comparison Table 2.9: Comparison with [START_REF] Chaurasiya | Parameterized posit arithmetic hardware generator[END_REF] for standard posit addition and product Op Format LUT DSP Delay (ns) [ As no sources is provided, we report as-is the figures from [START_REF] Chaurasiya | Parameterized posit arithmetic hardware generator[END_REF], obtained for a Zynq-7000 (xc7z020clg484-1) with Vivado 2017.4. To limit the possible effect of tool improvement on the synthesis, MArTo synthesis results have been obtained for the same part with Vivado HLS/Vivado 2018.3, the oldest version available for download at time of experimentation. Synthesis reported in [START_REF] Podobas | Hardware implementation of POSITs and their application in FPGAs[END_REF] trarget Stratix V FPGA. Results are read from a graphic plot, hence the approximate values. As there is no version of the Intel HLS toolchain that supports both Stratix V and the C++ 11 standard used in MArTo, the C++ to HDL compilation is done using Vivado HLS. The obtained HDL is then synthesised and routed for Stratix V using Quartus.

+ <16, 1> ∼500 0 ∼49 ∼550 <32, 2> ∼1000 0 ∼51 ∼520 × <16, 1> ∼330 1 ∼35 ∼600 <32, 2> ∼600 1 ∼38 ∼550 MArTo + <16, 1> 274
Despite being baroque, this toolchain seems to give good results, except for the <32, 2> product where it lacks the knowledge of the target's DSP possible configurations. Indeed, the product is computed using a 36x36 configuration of the DSP block, where a 27x27 configuration would be faster. IEEE-754 vs posit hardware cost comparison of floating-point for Xilinx FPGAs. It supports some of the IEEE features, such as infinity and NaN encoding. However, it is not IEEE-compliant: although the memory format is that of IEEE floats, subnormals are flushed to zero to save resources.

Considering that the Xilinx Float adders use DSP blocks to implement some of the shifters, the hardware costs of Xilinx adders and IEEE adders (obtained with MArTo) are really comparable. This illustrates that the overhead of subnormal handling in floating-point adders is small. Conversely, there is in Table 2.14 a very large difference in the resources used in multipliers. This demonstrates the cost of hardware subnormal handling in floating-point multipliers.

The Xilinx IP pipelining also seems to be fixed, and do not benefit from a relaxed clock constraint to reduce the latency, hence their large latency.

Since Xilinx floats lack subnormal support, the following bases on MArTo only the posit versus IEEE comparisons.

Comparison between posit and IEEE-754 operators

Table 2.13 compares combinatorial implementations of posits and floats of the same size on addition and multiplication. In this table, the "posit→posit" lines present results for the classical posit operators of Figure 2.4. The "PIF→PIF" lines presents results for the posit-compatible PIF operators that use the architecture of Figure 2.11b, including the inplace round component.

A first observation is that posit arithmetic is indisputably both larger and slower than IEEE-754 arithmetic. This contradicts the comparison in [START_REF] Chaurasiya | Parameterized posit arithmetic hardware generator[END_REF], which seems to use a very poor IEEE implementation.

As expected, the PIF-to-PIF operators are lighter and faster than the positto-posit ones. They still pay the price in area of a wider significand datapath (see Table 2.7) compared to IEEE operators: for the adders, PIF-to-PIF consume more LUTs than IEEE operators; for multipliers, they consume more DSP blocks (there is a step effect due to the discrete nature of DSP blocks). Again we observe in the IEEE multipliers the logic cost of subnormal support, but we also observe a comparable cost in the PIF multiplier, essentially due to the inplace round logic. Still, the PIF to PIF operators achieve delays that are closer to those of IEEE operators than to those of posit operators, which was their main motivation.

Note that the area cost of PIF/posit conversions (altogether about half the size of a complete IEEE adder) must still be paid in a posit arithmetic unit that uses the PIF-to-PIF approach. Only its delay (altogether about half the delay of a complete IEEE adder) is avoided. However, there is another advantage in a PIF-to-PIF PAU: the hardware for these conversions is naturally shared between different operations (such sharing is also possible in principle in a posit-to-posit PAU, but then it will restrict instruction-level parallelism).

Table 2.14 compares pipelined versions of the same operators, targeting a frequency of 333 MHz (3ns cycle time), and producing one output per clock cycle. As the latency estimated by the Vivado HLS tool is usually pessimistic, the reported latencies are obtained by an automated exploration that finds the smallest pipeline depth allowing the design to run with the target clock period. The script performing this exploration is open source, and is also accessible from the MArTo repository for reproducibility.

There is no PIF to PIF line in this table: for this setup, the PIF to PIF approach fails to provide any latency improvement (the arithmetic operators require the same number of cycles, and sometimes require one more cycle). We therefore choose not to report these results, which we consider synthesis artifacts as they are inconsistent with the expectations and with Table 2.13.

The cost of supporting all rounding modes in IEEE

Tables 2.13 and 2.14 report result for IEEE operators that only support round to nearest, ties to even. However, supporting the five IEEE-754 rounding modes increase only very slightly the hardware cost. For instance, adding this support to the 32-bit adder increases its area from 434 to 458 LUTs and actually decreases the delay from 11.9 to 11.7ns (another synthesis artifact). It remains well below the posit cost.

Quire versus standard operations

Synthesis results for the quire are given in Table 2. [START_REF] Caulfield | A cloud-scale acceleration architecture[END_REF], where we use MArTo to write a C++ loop that performs the sum of 1000 products and return the result as a posit. They are compared to a similar loop using floating-point Kulisch accumulator, and using regular floating-point hardware.

Quire is presented in unsegmented (U) version along with two segmented versions (S32 and S64 for segments of 32 or 64 bits). For 32 bits, the unsegmented version is not able to achieve 3ns cycle time, due to the long carry IEEE-754 vs posit hardware cost comparison propagation.

The Kulisch accumulator used here is also based on a 2's complement segmented accumulator [45, variant 3], with an IEEE-compliant final conversion to float (round to nearest, ties to even). The implementation has been validated against MFPR [START_REF] Fousse | MPFR: A multipleprecision binary floating-point library with correct rounding[END_REF] simulations.

Unsurprisingly, the cost and performance of a posit32 quire and a Kulisch accumulator for 32 bits floats are almost identical.

An exact accumulator consumes vastly more resources than standard operators: a factor 10 for 32-bit floats (a smaller factor for posits, but only due to the higher cost of the standard operators). Such factors should not come as a surprise: the 512 bits of the posit32 quire are indeed 18 times the 27 bits of the posit32 significand. This is the price of the accuracy of an exact accumulator.

Another advantage of exact accumulation is that it offers a latency reduction proportional to the latency of the floating-point or posit adder (here a factor 6-7). This is thanks to the fact that the accumulation loop is an exact fixed-point addition, which offers opportunities to exploit more parallelism in its computations [START_REF] Brunie | Modified Fused Multiply and Add for exact low precision product accumulation[END_REF], [START_REF] Uguen | Application-specific arithmetic in high-level synthesis tools[END_REF].

Detailed synthesis results of the quire sub-components are given in Table 2. [START_REF] Sarkar | Fpga implementation of dna sequence alignment with traceback[END_REF]. The quire addition line reports the cost for the architecture of Figure 2.9, including the large shifter and the fixed-point accumulation loop. This component accepts a new input every cycle. The two other lines describe the conversion of the quire result back to posit. The carry propagation is a loop component that adds zeroes, and is mostly merged with the quire addition component. However, there is an irreducible latency for the final carry propagation once the accumulation is over.

The latency overhead of the expensive conversion from quire to float or posits is easily amortized for large loops. However, it is also clear that a hardware quire will be very inefficient when used for small sequences of operations (e.g. fused multiply and add, complex arithmetic, small matrices or convolutions, etc).

Case study conclusion

This case study both allowed to demonstrate the relevance of having an HLS library for arithmetic components, and to compare the hardware cost of posit and IEEE-754 basic operation implementation.

Regarding the first point, it is simple with such a library to replace the arithmetic type in use for a given application by using a few typedefs. This allows easy experimentation to determine the format that gives the best accuracy/performance trade-off for a given application. On the second point, the take-away message of this study is that the indisputable complexity of the IEEE-754 standard, much attacked by posit proponents, does not necessarily translate into expensive hardware. Among the features that the posit system discards as useless, most (in particular overflow management, NaNs, and directed rounding mode) were designed to be implementable at very little cost. The only really expensive feature is subnormal support, due to rounding happening in a variable position of a bit vector. Posit arithmetic, despite the simplicity and elegance of the number system, involves such variable-position rounding, and therefore entail an overhead that is comparable in nature to subnormal support.

Limits and future work

A first issue with the current library design is the imposed default choice for IEEE-754 number rounding mode. This is however only a design issue and could be solved by adding a policy parameter to the IEEENumber template class determining the default rounding mode.

A more fundamental issue comes from the limits of C++ template metaprogramming. While the C++ template system has been shown to be Turingcomplete, many obstacles hinders the development of a full library based only on this system. Limits and future work 1. the C++ standard version accepted by the tools accepts only very simple constexpr constructs, limiting a lot what is practically writable, 2. there is no tool to move step-by-step into the template instantiation mechanism, making such code difficult to debug and limiting the manageable complexity of code using this system, 3. template instantiation mechanism has impact on compilation time (while this overhead is negligible compared to the time required to place and route an FPGA design, it is very significant when iterating on the software version of the design to debug it), 4. there is no way to reuse external runtime libraries to precompute useful values at compilation time.

This last point is quite annoying: a lot of libraries exists to determine interesting values for arithmetic operators (for instance "good" coefficients for polynomial approximation of arbitrary functions) which would require a huge amount of expertise and a lot of work to rewrite as full compile-time libraries (and suffer from the issues listed above). Next chapter describes two approaches to solve these issues by getting support from the compiler.

HLS library for arbitrary fixed-point function approximations

The previous chapter has demonstrated the possibility of creating fully parametrized HLS operators for basic arithmetic operation using only pure C++ meta-programming. This is an interesting first step toward the "ultimate goal" for arithmetic support in HLS, that would be to allow the construction of optimized hardware for arbitrary arithmetic operators (in the sense of section 1.3). A second step toward this goal consists in supporting operators approximating arbitrary unary function, with fixed-point input format and destination format. This is an interesting objective because the generation of such hardware is a well studied problem (a brief overview of the existing solutions is given in section 3.3), but no existing HLS tool support it. Besides, architectures for approximating arbitrary functions requires complex value pre-computation (for instance polynomial coefficients, or function values for tabulation). Computing these values fall behind the limit of what can be easily and efficiently done using C++ meta-programming. That is why alternative methods are required to reach this objective. This chapter introduces these methods.

On the user side, a library allows HLS developers to build operator specifications in their code. Two methods to implement the library backend that convert these specifications to hardware description are proposed. Both of them require a modification of the HLS compilation flow.

The first method, presented in section 3.4, allows doing so without modifying the HLS compiler, by introducing an intermediate compilation step [START_REF] Forget | A single-source C++20 HLS flow for function evaluation on FPGA and beyond[END_REF]. This approach is portable (as it is not bound to a specific HLS compiler) but requires a bit of extra work from the library user.

When modifying the compiler is an option, an alternative method, transparent to the user, is available. It is presented in section 3.5.

Custom C++ HLS compiler supporting C++20

The HLS compiler used to implement this method is a custom C++20 compiler targeting Xilinx devices. It is introduced in section 3.1.

Custom C++ HLS compiler supporting C++20

This section describes the custom C++ compiler that has been developed for enabling experimentation with C++20/C23 constructs on HLS targeting Xilinx FPGAs.

This compiler has been initially developed to bring experimental SYCL support to Xilinx devices [START_REF]SYCL for vitis: Experimental fusion of triSYCL with intel SYCL oneAPI DPC++ upstreaming effort into clang/llvm[END_REF]. SYCL is a programming standard from the Khronos Group [START_REF]SYCL 2020 specification[END_REF] to program applications using heterogeneous accelerators within a modern single-source C++ framework. Intel has developed a clang fork that support SYCL on various devices, rebranded the oneAPI DPC++1 compiler [START_REF]Intel[END_REF]. This fork has been itself forked to add support for Xilinx FPGAs. Both forks rebases on the latest Clang/LLVM developments on a regular basis, which allows using all C++ constructs supported by recent clang++.

Internally, this custom compiler uses Vitis HLS, as a backend. Vitis HLS consists in a derivative from Clang/LLVM [START_REF]The llvm compiler infrastructure[END_REF], with added support for custom pragmas and some specific optimization passes. The tool is frozen at LLVM version 7. Moving the Vitis frontend back from the past to align it with the current LLVM 15 seems to represent an enormous amount of work. Instead of doing so, the strategy is to feed the LLVM IR produced by a modern clang++ 15 compiler to Vitis HLS, bypassing the Vitis HLS old clang++ frontend. This architecture is visible on figure 3

A library to specify arithmetic operators

As stated earlier, the objective of this work consists in allowing the construction of hardware arithmetic operators having one fixed-point input and an arbitrary approximated function. This section introduces the library that is provided to the user for specifying such operators.

The library is actually made of two layers. A first layer describes elementary fixed-point operations. It can be seen as the counterpart of the MArTo library for fixed-point formats, and could actually have been part of it. However, as the compiler used for the experiments supports C++20 concepts and C23 BitInt, it was also the opportunity to test the design improvements enabled by these features. That is the reason why this fixed-point operations library is not a regular part of MArTo. In addition, C++20 support enables the use of functions of the standard library that became constexpr 2 with this standard. These functions would have otherwise needed to be reimplemented. The same work could however have been done using the ancient C++ stan-dard supported by main vendor HLS tools, as shown in the previous chapter. The fixed-point layer of the library is detailed in section 3.2.1.

The second component of the library allows the definition of mathematical expressions, that are used to specify the function that the operator approximates. This library is detailed in section 3.2.3. The same remark as for the first layer applies: C++20 concepts and constexpr standard functions have been used beneficially within the library, but it is perfectly possible to develop an equivalent library without them. This is only a matter of productivity of library development.

C++ types for fixed-point number

In the fixed-point library, a fixed-point format is represented by the template class FixedFormat. The two first template parameters are respectively the most significant and least significant bit position (similarly to what is described in section 1.2.1). A third template parameter specifies whether the format is signed. For instance, the type FixedFormat<31, 0, signed> represents a fixed-point format which is equivalent to a 32-bit signed integer.

A FixedFormat is used to parametrize the template class FixedNumber which represents a value of a given format.

A FixedNumber can be constructed from the representation of its value. For instance, in Listing 3.1, the variable X is initialized with the value 2 -2 because only its second bit (of weight -2) is set. A FixedNumber may also be constructed from the value of a standard arithmetic type, thanks to the static method get_from_value. This conversion returns the value of the target format closest to the source value. The tie breaking rule is toward +∞.

Listing 3.2 presents an example of computation performed with the fixedpoint library.

A user-defined literal (UDL) allows to create a FixedNumber from a constant value without having to specify its format. The narrower format that can represent exactly the expression will be automatically deduced from the operator. As decimal floating-point literals may have non-finite representation (for instance 0.1 10 = 0.000110011001100... 2 ), only hexadecimal floating-point literals are accepted by this UDL in the current library state. Finally, conversion from FixedNumber to standard arithmetic types are expressed with the get_as template method from FixedNumber. This method takes a target arithmetic type as template parameter and returns the value of this type which is the closest to the represented value (with the same tiebreaking rule as the other conversions).

Classical arithmetic computations

The proposed library defines standard arithmetic operations between FixedNumber of different FixedFormat, similarly to what the hint library does with arbitrary width integers. In the current experimental state, the modern C++ compiler presented in section 3.1 is the only target of this library. This compiler supports C23 _BitInt, so a backend using these types has been developed for hint. Fixed-point functions are built on this backend 3 to benefit from the strong semantic checking and useful API of hint 4 . Addition, subtraction and multiplication are computed exactly, the FixedFormat of the result being deduced by the library to be wide enough to hold the result. Rounding, overflow or saturation must be managed explicitly: FixedNumber provides additional methods to slice it, round to a specific precision, or extend the number to a wider format.

The proper way to handle division and remainder is still not clear. Some HLS tools provides Arbitrary fixed-point format libraries, as for instance ap_fixed and ac_fixed, the fixed-point counterparts of respectively ap_int and ac_int. These libraries offer quite similar functionalities in terms of elementary operation support. They also allow more flexibility in the format definition, by exposing policy parameters to set the overflow and rounding behaviors. They however suffer from the same issue as their integer counterparts concerning the lack of type checking and the transparent casting that can induce counter-intuitive behavior. The original feature from the proposed library is its ability to evaluate arbitrary mathematical expressions of fixed-point inputs. Next section describes how to represent such expressions.

Arbitrary mathematical function specification via its expression graph

Applications often require the evaluation of mathematical functions which are themselves defined as compositions of elementary operations and functions, for instance ( ) = log(1 + ). Such a function can be evaluated at run-time by a composition of library components (here an exponential, an addition and a logarithm). It is however often more efficient to get a single component to evaluate this expression, as illustrated by Figure 3.2. For this purpose, the proposed library also provides a mechanism to describe a function of one variable, then generate an operator to evaluate this function under an accuracy constraint.

In the proposed approach, a composite function may be defined as a C++ expression with a specific property: it has a free variable defined by the FreeVariable constructor. More precisely, this expression must be a tree of mathematical operations or elementary functions, and the leaves are either free variables or constants. A valid expression, for the purpose of building an evaluator, should have at most one free variable. The expression tree defines the function that the resulting operator should approximate. Its operand input format is the format of its FreeVariable leaf, and the output format is defined when requesting an evaluation. to evaluate it can be specified. At line 4, a FreeVariable node is created from a fixed-point variable. The FreeVariable format is deduced from the source FixedNumber, and corresponds to the operator input format in the formalism of section 1.3. At line 5, a Constant leaf is created using a user-defined literal 5 .

The library also has special constants to represent numbers like , which are kept symbolic (infinitely accurate) in the expression tree defining a function. The deduced type of variable f at Line 6 represents the expression tree corresponding to the function :

→ log(1 + ). The call to the evaluate template function specifies an arithmetic operator:

• the function it approximates is given by its argument (here f),

• the output format is given as template parameter,

• the input format is obtained from the function free variable leaf,

• the rounding relation is implicit, as only faithful rounding at the moment.

An arguably better design in terms of separation of concern would have consisted in having format-agnostic FreeVariable, and pass the operator input as a second argument to the evaluate function, that could then be called Fixed-point function approximation architectures apply. This might become the new API in future library versions.

Once a specification is obtained, Hardware matching it can be generated. Next section introduces briefly the architecture used to implement these hardware operators.

Fixed-point function approximation architectures

Multiple architecture exists for hardware implementation of arithmetic operators. Two main families exists for general function approximation: tablebased methods and polynomial methods.

Table-based hardware arithmetic operators

The simplest form of table-based function approximation operator is plain tabulation. A pre-computed table contains the function output value for each possible input value. The runtime computation consists only in reading the value associated to the input in the table. This architecture is interesting for small input width operators, as it has a low latency. However, as the storage requirement evolves exponentially with the input width, it becomes costly with wide inputs. In this case, an alternative consists in using a multipartite table decomposition.

The simplest form of multipartite method is the bipartite [START_REF] Sarma | Faithful bipartite rom reciprocal tables[END_REF], [START_REF] Sunderland | Cmos/sos frequency synthesizer lsi circuit for spread spectrum communications[END_REF] table method. As illustrated by Figure 3.3, this method is based on a piecewise linear approximation. The input domain is partitioned in 2 sub-intervals by splitting the input as = + 2 -, where (the interval index) consists of the leading bits of and (the index within one interval) consists of the least significant bits. The linear approximation on each subdomain indexed by is ( ) ≈ ( ) + 2 -× ( ) where ( ) is the slope of the approximation segment. The idea of the bipartite method is to tabulate the multiplications × ( ). To reduce the cost of this tabulation, the slopes are shared between 2 consecutive intervals, so instead of × ( ) it is possible to tabulate × ( ) where consists of theleading bits of .

Many evolutions to this method have been developed (see [START_REF] Hsiao | Table size reduction methods for faithfully rounded lookup-table-based multiplierless function evaluation[END_REF] and references therein, and [START_REF] Christ | Lossless differential table compression for hardware function evaluation[END_REF] for a method for compressing the table of initial values). Due to lack of time, the current flow only provides a basic bipartite approximation: its purpose is to demonstrate that the approach is not restricted to simple tabulation. The current implementation computes the parameters and , and fills the corresponding tables for a faithful operator. Plain tabulation is used when no bipartite approximation with a better table storage cost is found.

Polynomial approximation methods

Polynomial approximation methods consists in finding a polynomial that is close enough to the function to approximate on the operator input domain. The approximated function value is then obtained by evaluating the polynomial, using for instance the Horner evaluation scheme. Finding good polynomial for a given function is not trivial, as two optimization goal coexist: the approximation error should be small enough to respect the overall accuracy constraint, and at the same time the polynomial has to be efficiently evaluated. The later point implies in particular that the polynomial degree should not be too high, and that the coefficients should have a compact binary representation to avoid wide intermediate products.

Tools like sollya [START_REF] Chevillard | Sollya: An environment for the development of numerical codes[END_REF] allows finding good polynomial for a specified function and a specified error bound. These tools are already used by HDL generators such as FloPoCo [START_REF] De Dinechin | Reflections on 10 years of FloPoCo[END_REF]. The main contribution of this thesis is not to improve the underlying methods, but enable the usage of these tools in an HLS context. The following sections present two methods for achieving this purpose.

Compiler agnostic specialization generator method

The first method is useful when the HLS compiler is a black box, that cannot be modified. It consists in splitting the HLS compilation process in two steps. This process is detailed on figure 3.4. The first involves compiling the application for execution on traditional CPU. This compilation step can be done with any off-the-shelf C++ compiler. The executable resulting from this compilation is called the specialization generator.

Running the specialization generators produces C++ implementations for all the operators present in the source application. These specializations are dumped in a specialization header, called operators.hpp on figure 3.4. This generated header file can then be included in the application source file app.cpp, and be compiled with an HLS compiler.

Here again, there is no constraint on the HLS compiler, except that it should be able to accept both the application and specialization header code 6

C++ types for arithmetic operator evaluation

Listing 3.4 present a toy example that illustrates the main mechanisms that are used in the specialization generation method. One of these mechanism is template specialization. It allows specifying that a different code from the default template code should be generated when a template is instantiated with certain parameters. This is visible with the TemplatedClass specialization at line 9 of the listing. The class definition that follows will be selected when the template is instantiated with int as its parameter. For all other parameters, the default template class defined at line 2 is used. The second mechanism is the dispatching mechanism used to instantiate a template instantiation based on a function argument type. This is illustrated by the dispatcher templated function. When called with a float parameter (line 24), the template parameter is deduced to be float, so the argument is forwarded to TemplatedClass<float> which is instantiated using the default template. At line 27, the call with an int argument is dispatched to the int specialization of TemplatedClass.

The evaluate function of the library is used similarly to require a specific instantiation of a template class that depends on the expression type and output format. The outline of this templated class, ArithOp, is given in listing 3.5. Its first template parameter, ET, is the type that represents the expression tree to evaluate. The second template parameter, OutputType, is the destination type of the operator.

The evaluate method of ArithOp is responsible for computing the result of the operator evaluation. Instead of dispatching the expression to a static method of ArithOp, evaluate7 dispatches it to the evaluate member method of a templated ArithOp variable. This detail is important, as it has for consequence that the constructor of this variable will be called before entering the main function.

In the default ArithOp template, this constructor has two responsibilities:

• determine an architecture to evaluate the specific expression of its template parameters,

• generate code that implement this architecture and dump it in the specialization header.

The first point is achieved by using an architecture generation library, which query tools such as sollya and FloPoCo to determine an efficient hardware implementation for the given expression tree. In detail, a runtime representation the expression tree is built and converted to a representation understandable by the external tools.

In the bipartite generation case, sollya is used to build the highly accurate reference table of the function values. This table is then processed to find a good bipartite decomposition. If no good decomposition is found, then the tool falls back to a plain tabulation. The chosen algorithm and associated tables are then saved in the specialization header.

Listing 3.6 gives a simplified example of what the implementation of the specialization of ArithOp corresponding to the evaluate call at line 7 of listing 3.3. This specialization is generated when running the opgen.elf executable. It can be seen that the type specialization is fully specified by a type representing the expression tree log(1 + ) (with the free variable). ) with output format ufix(5, -17).

Such specializations are generated for each operator specification present in the program. As a consequence, when including the specialization header, there is no more use of the default template. As these specializations have an empty constructor, and their evaluate method are pure functions that only use fixed point constructs, they can be fed to the HLS compiler.

Application example

The next sections present two applications that have been built using the library in specialization generator mode [START_REF] Forget | A single-source C++20 HLS flow for function evaluation on FPGA and beyond[END_REF] and phases = 0. The amplitudes is an input of the kernel. In both versions, the oscillators are instantiated using template unrolling, to ensure that they are inlined.

In the FixedNumber implementation, the are fixed-point number of format FixedFormat<-1, -8, unsigned>. The sine is computed from a wave table of 2 12 entries generated by the library. The function thus evaluated is ( ) =

(2 ) for in [0; 1)8 . Two experiments, FixedNumber 8 and FixedNumber 16 have been performed with an evaluator accuracy constraint of 2 -8 and 2 -16 respectively.

In the binary32 implementation, the std::sin function from the mathematical library is used. All the computations are done in binary32.

Function values over one period for all the formats and difference with binary32 for FixedNumber 8 and FixedNumber 16 are plotted on Figure 3.5. The graphs for all the format coincide perfectly, and as expected, the absolute difference to float is smaller for FixedNumber 16 than for FixedNumber 8 .

The evaluation is performed using our compiler and targeting a Virtex ultrascale plus FPGA with part number xcvu13p-fhga2104-3-e. Both pipelined and unpipelined solutions have been generated for the three formats. Results are given after place and route. The design metrics are reported in table 3. 3.1: Area and timing metrics comparison between float and various FixedNumber for an additive synthesizer of 256 oscillators. II stands for Initiation Interval, the number of clock cycles that pass before the pipeline is ready to be fed a new input.

These results show that fixed-point is definitely a good option for this application. Three factors can explain this better quality of results. The first one is a question of bitwidth of the argument that have to be moved in the design. Indeed, the 256 8 bits amplitude coefficients input of the FixedNumber case can be implemented as a very wide input register while it is not possible for the 256 32bits coefficients of the input. Secondly, the intrinsic arithmetic of floating point is more complex than fixed-point on the operation involved in this application. Lastly, the fact that floating-point arithmetic is based on external IP integrated by the design can hinder some optimization opportunities.

Single-source end-to-end LNS with SYCL

As the compiler used in our developments also offers an experimental support for compiling SYCL programs targeting AMD FPGAs, we tried to use the library in a full end-to-end single-source application. In the SYCL compilation flow, the computation kernels are outlined by the compiler, and compiled separately for the hardware targets. Instead of producing a Vitis IP, it generates a x86 fat binary that embeds the hardware kernel binary, and contains all the glue to invoke the kernel execution on the FPGA. A small SYCL application was developed, consisting in one kernel that computes the sum of two LNS number having an exponent of format FixedFormat<7,-8>, using the adder shown on listing 3.7. This application has been built to run on Xilinx devices. This required a bit of manual adjustment of the SYCL compilation flow, but nothing that could not be automated in the future. It was then possible to run the program that moved the operands to the FPGA, launched the kernel, and get the result back. No resources are reported here, as the hardware synthesized to ensure communications between the FPGA and host CPU is way bigger than the small LNS adder.

Limits of the approach by specialization generation

Using the type system to represent the expression raise a fundamental issue. Indeed, it is not possible to distinguish two variables sharing the same type. This is in particular problematic with the FreeVariable class. Listing 3.8 illustrates this issue: the FreeVariable x and y share the same type due to the fact they are created with values of the same FixedNumber type. As such, there is no way to distinguish between a valid expression tree that uses multiple time the same FreeVariable (line 6) from an invalid expression that mix usages of x and y (line 8). The approach followed by the library is very conservative: only expression trees with at most one free variable leaf are considered valid. With this approach, code from listing 3.8 raise a compile-issue complaining Compiler agnostic specialization generator method of an invalid expression both for valid_expr and invalid_expr. This avoids error from the user at a price of reduced expressivity. An alternative would be to ask the user to manually specify an uniquing type as parameter to FreeVariable (e.g. FreeVariable<class MyUniqID>{}). However, this would clutter the library API, stays quite error-prone and would have the side effect that identical expression of two distinct variables will no more share the same architecture generation step.

Another limit of this approach is the relative difficulty to perform optimizations on expression tree. Supposing that the library accept the expression + (which is not the case as appears twice in the expression tree), it would be desirable to have a mechanism to optimize it as 2 . Similarly, some trigonometric identities could be applied to "canonicalize" the expression representation and limit the architectural exploration for mathematically identical expression of different expression graph. This would require a way to apply a collection of rewriting rules to optimize the mathematical expression. This is actually a good part of the job of optimizing compiler job. However, these compilers first build intermediate representation (IR) of the program that are easy to manipulate. Optimizations are then performed on this IR. Following section explores a method that follows this approach.

Compiler support for fixed-point functions in HLS

The second approach used to provide support for fixed point function in HLS using the same user interface relies on compiler support. The overall architecture is presented in section 3.5.1. In order to support the actual fixedfunction compilation process, intermediate representations have been developed to represent such expressions. These IR are described in section 3.5.2.

Fixed-point function compilation architecture

The architecture diagram of the solution is presented on figure 3.6. The user has only one compilation step to launch.

On this backend, operations on expression and FreeVariable constructor are implemented using custom compiler built-ins. When compiling application code, the compiler identify these built-ins, and is able to extract the operator specification from the rest of the code. The rest of the code that does not correspond to operator specification is directly compiled into the compiler internal representation (LLVM IR in this case). This corresponds to the IP.bc file on figure 3.6. In parallel, the operator specifications are compiled into a custom IR, which corresponds to ArithOps.ir on the figure. A specific compiler (the expression compiler on the figure) is used to convert this high level specification into low level LLVM IR (ArithOps.bc). This LLVM IR and IR.bc are fed to sycl_vxx.py, that will link them together and run optimization and downgrading passes on the resulting code before feeding it to Vitis HLS. The lowering done by the expression compiler is a multi-stage process, that involves multiple intermediate representations of the operator at different abstraction level. These intermediate representations and some optimizations that are (or could be) performed on them are detailed on figure 3.7. They consists in MLIR [START_REF] Lattner | MLIR: Scaling compiler infrastructure for domain specific computation[END_REF] dialects. MLIR, for multi-level intermediate representation, is a compilation framework that allows the definition of custom intermediate representations. These IR are called dialects, and consists in collections of custom types and operation on these types. The MLIR compiler framework provides all the infrastructure needed to define these dialects, apply rewriting rules for optimization and define conversion rules to lower a high-level dialect to a lower-level one. A standard MLIR dialect exists to represent LLVM code. The aim of the lowering block is then to translate the high-level operator specification given in our custom dialect (the approx dialect) to low-level implementation in LLVM IR dialect. Once this representation is obtained, it can be directly translated to LLVM IR, that can be linked with the rest of the application. Next section details these custom MLIR dialects.

Intermediate representations for fixed-point functions

Two custom MLIR dialects describing arithmetic operators have been developed. The first one, the approx dialect is an almost one to one mapping with the library constructs for representing mathematical expression. This is illustrated by listing 3.10 that gives the high-level representation produced by the compiler for the operator specified in listing 3.9. The main difference between the C++ code and the approx dialect representation of the arithmetic operator is the requirement of having a named node for each intermediate node in the expression graph.

A second custom dialect, the fixedpt dialect, is used to represent actual fixed-point operations. An MLIR conversion pass lowers the approx dialect

Current state of the expression compiler prototype

The described compiler supports plain tabulation and basic polynomial approximation using Horner scheme. FloPoCo offers other function approximation methods (the state of the art in multipartite approximation [START_REF] Hsiao | Table size reduction methods for faithfully rounded lookup-table-based multiplierless function evaluation[END_REF] and very good piecewise polynomial approximation [START_REF] De Dinechin | Automatic generation of polynomial-based hardware architectures for function evaluation[END_REF]). These methods rely on fixed-point operations that are already implemented in the fixedpt dialect, so adding their support to the compiler should not be complicated. However, this requires refactoring on the FloPoCo side, in order to better separate the approximation logic from the VHDL generation code (that the complete HLS compilation flow replaces). This refactoring has been initiated, but completing it is left for future work.

Correctness of execution has been tested exhaustively in relatively narrow formats, and QoR of synthesized solutions has been manually checked on some examples to ensure sound results. However, there is no automatic testing tools yet, and this compiler project is still at an early development stage.

As it can be seen on listing 3.10, the current approx dialect represents all computation nodes with the approx.math operation. The first argument of this operation determines the actual computation performed by the node. This design choice was made in order to allow quick addition of computation nodes, but this is not a good choice when it comes to applying high level expression rewriting rule.

Besides, a lot of work remains to be done to optimize the evaluation strategy. For instance, the lowering from the approx dialect to the fixedpt dialect could benefit from algebraic optimization, for instance detecting how the evaluation range can be reduced. The lowering from the fixedpt to the arith dialect could also benefit from low-level optimization, for instance by using bitheap compression [START_REF] Brunie | Arithmetic core generation using bit heaps[END_REF] for efficient sum computations.

In its current state the compiler can help an expert to build an optimized arithmetic operator, and still allow a non-expert to at least get a fixed-point operator for arbitrary function, but the ultimate goal of easily and automatically generate operator with excellent quality of result is not reached yet.

As a final remark, having an efficient expression compiler is not only beneficial to the integrated workflow. Indeed, there is no fundamental obstacle to developing a backend that generate HLS-supported C++ code from arith IR (for instance using the hint library). As such, it is perfectly possible to integrate the expression compiler in the specialization generator flow, by interfacing it with the "Architecture generation library" from figure 3.4. This will however not solve the issue of variable type identification. As such, the integrated flow is preferable when possible.

Conclusion and perspectives

The work presented on this thesis contributed to bring support for applicationspecific arithmetic operators to High Level Synthesis tools.

Part of this can be done using pure host language constructions (in this thesis, C++ was used as it is the language used by mainstream HLS tools), in a portable way. On this work, this has been illustrated in chapter 2 with the implementation of elementary operations for various floating-point formats. This allows application developers to write their code once and change the underlying numerical format with a simple typedef. This facilitates the choice of a format that provide a good accuracy vs performance trade-off on a perapplication basis. In the future, this source-only approach can be extended to other operators, as long as they do not require complex pre-computations. This includes for instance as mixed-precision or fused operators.

For operators requiring complex pre-computations, the HLS support is still possible, but requires adaptation of the HLS compilation flow. This is shown in chapter 3 with operators that evaluate arbitrary mathematical function with fixed-point input and output format. Implementations of arbitrary function evaluators use function-specific pre-computed values such as approximation polynomial coefficients. Computing these values is complex enough that special tools have been developed in this purpose, such as FloPoCo or Sollya. Two methods have been proposed to use these tools to compute function-specific values before injecting the function evaluation algorithm in the standard HLS compilation flow. The first is based on a sequence of compilation stages, using unmodified off-the-shelf compilers, and is therefore portable. It uses a pre-compilation stage that produces an executable that invoke the existing tools to produce the HLS-enabled operator implementation. That implementation can then be fed to the HLS compiler. This approach is however limited due to it being based on the C++ type system to extract the expression graph of the mathematical function to evaluate. The second method solves this issue by adding built-in support for these operators inside the HLS compiler.

While this breaks the portability of the solution, as it requires modifying the compiler, the greatest part of the work can be factored to be reused in different compilers by using a compiler-agnostic intermediate representation (IR) for the operators. The compiler should produce an IR describing the operator high-level specification. A compiler-independent tool, the expression compiler, can convert this specification to a low-level implementation, which can be imported back in the original compiler. This design has been tested using a custom HLS compiler based on clang++ and Vitis HLS. The complexity of developing the intermediate representation for operators and the mechanism to lower it to implementation details was greatly reduced by the MLIR framework.

While the work presented in this thesis focuses in providing high level functionalities that were previously unsupported by HLS tools, the same technics could be used to port other hardware-specific optimizations to HLS tools. Part of the reason that operators generated by specialized tools such as FloPoCo have good latency and area comes from the clever usage of primitives such as bit heaps [START_REF] Brunie | Arithmetic core generation using bit heaps[END_REF] or multiple constant multipliers [START_REF] Kumm | Hybrid Multiple Constant Multiplication for FPGAs[END_REF]. In number of cases, it is possible to automatically identify the places where such primitives could be used beneficially at the abstraction level of fixed point operations in the IR. Having an IR representation for such operations could then be of interest to improve area and latency of library-generated operators. While this is a medium term perspective, a shorter-term goal is the support of additional functions evaluation architectures.

In a nutshell, supporting a greater part of existing state of the art constitutes a first progression axis for the proposed system. A second progression axis consists in offering higher level constructs. An example of this would be to allow the user to create a filter by frequency response specification. This is in phase with the objective of allowing non-specialized programmers to write programs with good productivity, by specifying what should be computed and delegating to the tool the choice on how to compute it. Embedded domainspecific languages in the C++ code seems to be a good interface for specifying such computations. Optimizations can be performed at each abstraction level view of the operator: from the high level function description (with algebraic optimizations), to the low level implementation (by using efficient fixed-point primitives for instance). Having an IR and associated optimizations for each abstraction level enables efficient implementation of these optimizations. Be-sides, the optimizations on low-level implementation details are not specific to mathematical function evaluation, and can be reused in other contexts.

The current system has two main limitations:

1. Vendor HLS compilers backend are closed-source, and constitutes a blackbox part of the design. It often requires retro-engineering to understand which IR construct will give good implementation result, and it makes the system very fragile to backend version change.

2. There is no cost model to estimate reliably the QoR of an implementation without having to perform the full synthesis, placing and routing steps.

The first limitation could be lifted by integrating with open-source HDL backend. The CIRCT project [START_REF]CIRCT -circuit IR compilers and tools[END_REF] is a collection of MLIR dialects representing digital circuits, that can be lowered to Verilog. Targeting these IR does not pose fundamental difficulties. While the ideal would be to have an opensource toolchain for end-to-end source-to-bitstream conversion (as the HDL can be considered to be another IR in this process), the HDL mapping to FPGA primitive is a bit more predictable than the source-to-HDL conversion realized by HLS backends. In the current state of the compiler however, only the operator codes go through the MLIR pipeline. The rest of the code is delegated to the HLS backend that will, among others, perform all the transformation required to get hardware that respects the call graph semantics of the original program. With the CIRCT based solution, this should also be done inside the MLIR pipeline. This is however not a fundamental issue, as proof of concepts of HLS compilers using the CIRCT framework have been developed [START_REF] Ye | Scalehls: A new scalable high-level synthesis framework on multi-level intermediate representation[END_REF].

Lifting the second limitation is more complicated. Indeed, it requires not only good knowledge of the hardware internal, but also of the synthesis and placement algorithms used to produce the FPGA bitstream out of the HDL description. This information is often considered sensitive by hardware vendors, and kept internally. Besides, experiments have shown that current vendor cost model does not perform well on small designs such as arithmetic operators. For instance, some logical optimizations are not taken into account by HLS cost models, which can result in local overestimates of latency of component parts, leading the pipelining system to insert more pipeline stages than what would really be required.

Improving these models is probably seen by vendors as too costly compared to the benefit they get by providing HLS solutions. This is however a chicken and egg problem, as the interest of HLS is limited by the QoR it can produce. Getting a better QoR by providing easy to use and to test libraries with state-of-the-art implementations is a good first step. The development in recent year of high quality compilation framework and hardware construc-

A

Source code for floating-point adders

This annex reports the MArTo library code used to implement addition for IEEE-754 and posit. In this code, the Wrapper template template parameter specify the hint backend for which the operator is implemented. See section 2.1.3 page 48 for more details.

Listing 2.7 page 57 gives an example of user-written code that will trigger the instantiation of one of the adder presented here, depending on the actual type of my_fp_type.

A.1 IEEE-754 Adder

The IEEE adder uses the IEEENumber, which is a helper class built on hint wrappers to provides useful methods such as exponent or fraction field extraction. The code for this class is presented on listing A.1. Listing A.2 presents the complete adder algorithm. Les techniques de synthèse de haut niveau permettent aux programmeurs non spécialistes de générer des descriptions de circuits numériques en utilisant des langages de programmation généralistes. Cependant, les outils existants ne supportent qu'un petit nombre de formats numériques et un petit nombre d'opérateurs standards. Cette thèse présente plusieurs techniques pour rajouter le support de nouveaux formats et de nouveaux opérateurs. Dans un premier temps, l'étude se focalise sur ce qui est réalisable en se restreignant aux fonctionnalités de métaprogrammation du standard C++ supporté par les outils HLS. Une bibliothèque d'opérateurs élémentaires pour les formats IEEE-754 et posit de taille arbitraire est proposée. Elle sert de base à une étude de cas comparant le coût matériel de l'implémentation de ces deux formats. L'implémentation d'évaluateurs de fonctions mathématiques arbitraires se heurte aux limites de la première approche. Dans un second temps, l'étude se porte sur les possibilités offertes par la modification du flot de compilation HLS, avec comme objectif de supporter cette fonctionnalité. Une bibliothèque permettant au développeur de spécifier des opérateurs pour approximer des fonctions arbitraires en précision arbitraire est présentée. Deux approches pour l'interfaçage de cette bibliothèque avec les outils de HLS sont proposées, selon que l'on a ou pas accès aux sources des compilateurs HLS.

Mots-clés

Figure 1 .Figure 1 . 1 :

 111 Figure 1.1 shows the representable elements of the format defined by table 1.1, and illustrates the value returned by the two directed rounding function on this format.

Figure 1 . 2 :Figure 1 .

 121 Figure 1.2: Binary positional notation of 42.75

. 3 .Figure 1 . 3 :

 313 Figure 1.3: Superposition of two fixed-point representation windows with the infinite positional representation.

Figure 1 .Figure 1 . 4 :

 114 Figure 1.4: Visualization of representable values for two 5-bits unsigned fixed-point formats.

Figure 1 . 5 :

 15 Figure 1.5: Decomposition of a signed fixed-point format with = 5 and = -4

Figure 1 . 6 :

 16 Figure 1.6: Visualization of representable for the two signed fixed-point schemes with = 1 and = -3. Scheme-specific representable values are circled in red.

Figure 1 . 7 :

 17 Figure 1.7: Binary and binary-coded octal positional notation of 42.75

Figure 1 . 8 :

 18 Figure 1.8: Binary-coded decimal positional notation of 42.75

Figure 1 .Figure 1 . 9 :

 119 Figure 1.9 gives a possible representation for the value 42.75 in such a

Figure 1 . 10 :

 110 Figure 1.10: Unique IEEE-754 binary16 representation of 42.75.

Figure 1 . 11 :

 111 Figure 1.11: Subtraction of two floating-point values of minimal normal exponent min .

  Posit<8,2> representation of 2 -20 .

Figure 1 . 14 :

 114 Figure 1.14: Posit<8,2> representations decomposition.

Figure 1 . 16 :

 116 Figure 1.16: Representable positive values of an LNS system with representation of the exponent being a sfix(2, -3).

Figure 1 . 2 ,

 12 Figure 1.16 show the representable values for an LNS encoding scheme having the exponent defined as an sfix(2, -3). The values form a regular grid on a logarithmic scale. One example of representable irrational value for binary LNS is √ 2, represented as 2 12 . LNS has however the drawback of having almost no representable successive integers. For instance, binary LNS can represent 0, 1 and 2, but next representable integers are power of 2.The rationale behind LNS is that it simplifies some elementary operations computations. Indeed, the representation of the product of two values is just the sum of their representations, and the power function evaluation is reduced to the computation of a product. However, this comes to the price of a complex addition.Indeed, for 0 and 1 two values such that 0 > 1 , represented in LNS by their logarithm 0 and 1 , the computation of the representation of = 0 + 1 is evaluated as follow:

Figure 1 . 17 :

 117 Figure 1.17: Wide adder implementation by Full Adder chaining.

Figure 1 . 18 :

 118 Figure 1.18: Part of the schematic of the UltraScale fast carry chain Source: [12].

Figure 1 .FieldFigure 1 . 19 :

 1119 Figure 1.19 schematizes the different compilation stages for a simple com-

Listing 1 . 1 :

 11 Verilog description of a circuit computing ( + ) • for 12-bits inputs. preadd_res; assign io_result = preadd_res * C; always @(posedge clock) begin if (reset) begin preadd_res <= 13'h0; end else begin preadd_res <= A + B; end end endmodule

Listing 1 . 2 :

 12 Chisel description of a circuit computing ( + ) • for 12-bits inputs. class PreAddProd extends Module { val io = IO(new Bundle { val A = Input(UInt(12.W)) val B = Input(UInt(12.W)) val C = Input(UInt(12.W)) val result = Output(UInt(25.W)) }) val preadd_res = RegInit(0.U(13.W)) preadd_res := io.A + io.B io.result := preadd_res * io.C } println(getVerilog(new PreAddProd))

Listing 1 . 3 :

 13 HLS C++ description of a circuit computing ( + ) • . int preAddProd(int A, int B, int C) { return (A + B) * C; } Listing 1.3 shows an example of a pure C++ HLS code describing the same computation as on the HDL/HCL examples. This code does not contain any pipeline related information. It is the role of the compiler to find a good pipelining for the implementation in consideration with the compilation objective (clock frequency, area or latency minimization).

Field-Programmable Gate Arrays Listing 1 . 4 :

 14 Simultaneous definition and call of an FPGA kernel computing ( + ) • . int main() { // Declare buffer of one element sycl::buffer<int> abuff{{1}}, bbuff{{1}}, cbuff{{1}}, rbuff{{0}}; // Write values inside the buffer on host (CPU) side { abuff.get_access()[0] = 5; bbuff.get_access()[0] = 24; cbuff.get_access()[0] = 17; } // Create a task queue that will run on FPGAs sycl::queue Q{sycl::fpga_selector{}}; { Q.submit([&](sycl::handler cgh){ // Will be executed on FPGA auto a = abuff.get_access(cgh)[0]; auto b = bbuff.get_access(cgh)[0]; auto c = cbuff.get_access(cgh)[0]; rbuff.get_access(cgh)[0] = (a+b) * c; print the result of (24+5) * 17 std::cout << rbuff.get_access()[0] << "\n"; } return EXIT_SUCCESS;

11 }

 11 res = a*b; // print 4294967295 (= 2^(32) -1) std::cout << res <<'\n'; Listing 2.1: Example of a counter-intuitive C intermediate result type.

}Listing 2 . 2 :

 22 cout << b << "\n"; Undefined behavior on shifts.

3 Figure 2 . 1 :Figure 2 . 2 :

 32122 Figure 2.1: A generic shifter (left) instantiated (right) in a floating-point adder with fraction bits

Listing 2 . 3 :

 23 template<template<unsigned int, bool> class Wrapper> // Takes a 12-bit vector as input, // Consider it as the concatenation of two unsigned integer // Compute the LZC of their sum Wrapper<3, false> computation(Wrapper<12, false> input) { auto high = input.template slice<11, 6>(); auto low = input.template slice<5, 0>(); auto sum = high + low; // 7 bits, unsigned return lzc(sum); // 3 bits, unsigned } Example of hint code computing the leading-zero count of the sum of two 6-bits integers provided concatenated in a 12-bits vector.

}Listing 2 . 4 :

 24 template<unsigned int MSB, unsigned int LSB, unsigned int size, bool isSigned, template<unsigned int, bool> class Wrapper> auto slice(Wrapper<size, isSigned> input) { return input.template slice<MSB, LSB>(); Pure function wrapping call to the slice template method.

  template<unsigned int MSB, unsigned int LSB> auto slice(HintWrapper input) { return input.template slice<MSB, LSB>(); } Listing 2.5: Pure function wrapping call to the slice template method with cleaner interface using C++20 concepts.

using my_fp_type = IEEENumber<9, 14 ,}Listing 2 . 7 :

 1427 hint::VivadoWrapper>; auto myFunction(my_fp_type a, my_fp_type b, my_fp_type c) { return (a*b) + c; MArTo code example.

1 using 2 / 4 auto 6 } 2 . 8 :

 124628 my_ieee_type = IEEENumber<9, 14, hint::VivadoWrapper>; /computes ↓ N ( * ) + 3 auto myFunction(my_ieee_type a, my_ieee_type b, my_ieee_type c) { prod = a * b;// Default rounding to-nearest ties to even 5 return ieee_add(prod, c, IEEERounding::roundDown); Listing MArTo code example.

Figure 2 . 3 :

 23 Figure 2.3: Example of product accumulation on the carry-save banks of a Kulisch-like accumulator.

Figure 2 . 4 :

 24 Figure 2.4: Architecture of a posit operator in a Posit Arithmetic Unit that uses posit registers and posit-to-posit operators.

Figure 2 . 5 :

 25 Figure 2.5: Architecture of a posit to PIF decoder.

  7) and multiplier (Figure 2.8) first compute the exact result (top part of the figures) using the transposition to the PIF format of classical floating-point algorithms.

Figure 2 . 6 :

 26 Figure 2.6: Architecture of a UPIF to posit encoder. The PIF to posit encoder is similar, with the round and sticky logic (including the final adder) removed.

Figure 2 . 9 :

 29 Figure 2.9: Architecture of a posit quire addition/subtraction.

  detect. Underflow detect. '0' detect.

Figure 2 . 10 :

 210 Figure 2.10: The bits of a standard quire.

IEEEFigure 2 . 11 :

 211 Figure 2.11: Architecture of a PAU using posits as a memory-only encoding, with PIF registers and PIF-to-PIF operators.

Figure 3 . 1 : 1 ↩→

 311 Figure 3.1: Custom C++ HLS compiler stages

1 usingListing 3 . 1 :

 131 Construction of a FixedNumber from its representation.

}Listing 3 . 2 :

 32 using radius_fmt = FixedFormat<3, -4, unsigned>; using pi_fmt = FixedFormat<<1, -4, unsigned>; auto circumference(FixedNumber<radius_format> radius) { // Generate a constant of value two // The deduced format is FixedNumber<1, 1, unsigned>{1}; constexpr auto two = 2._fixed; // Generate an approximation of PI rounded to 2^(-4) // !! PI is already an approximation of !! auto pi = FixedNumber<pi_fmt>::get_from_value(PI); // Diameter is a FixedNumber<4, -3, unsigned>; auto diameter = two * radius; // circumference is a FixedNumber<6, -7, unsigned> auto circumference = diameter * pi; // Eventually we get rid of low bits // Return type is FixedNumber<6, -4, unsigned> return circumference.round_to<-4>(); Format inference in the fixed-point library.

FixedFormat<5Figure 3 . 2 :FixedNumber<FixedFormat<5, - 12 ,Listing 3 . 3 :

 321233 Figure 3.2: Example of function evaluator (here corresponding to listing 3.3).

Figure 3 . 4 :

 34 Figure 3.4: Architecture of the generator of template specializations

}Listing 3 . 4 :Listing 3 . 5 :

 3435 myAction(T) { std::cout << "Default template!\n"; our case, would be defined in a specialization header template<> class TemplatedClass<int> { static void myAction(int value) { std::cout << "Int specialization: " << value << "\n"; to TemplatedClass<float>: // Will print "Default specialization" dispatcher(3.14f); // Dispatch to TemplatedClass<int>: // Will print "Int specialization: 42" dispatcher(42); Simple example of template specialization and dispatching. template<ExpressionType ET, FixedFormatT OutputType> struct ArithOp { ArithOp(){ // Find function evaluation plan // and dump it to the specialization header } // empty default implementation FixedNumber<OutputType> evaluate(ET expr) { /*...*/ } }; Outline of the generic ArithOp template class. */FixedFormat<5, -17, unsigned>> { auto evaluate(LogExpr<...> expr) { // Detailed instantiation, for instance polynomial evaluation // or table read } }; Listing 3.6: Simplified example of generated ArithOp specialization for log(1 +

Figure 3 . 5 :

 35 Figure 3.5: Output of an additive synthesis with two frequencies of equal amplitudes and the computational error (as absolute difference on a logarithmic scale).

Listing 3 . 8 :

 38 using input_t = FixedNumber<FixedFormat<5, -12, unsigned>>; auto some_compute(input_t xVal, input_t yVal) { FreeVariable x{xVal} FreeVariable y{yVal}; // Valid because only one free variable auto valid_expr = sin(x) + cos(x); // Invalid because two distinct free variables auto invalid_expr = sin(x) + cos(y); using vexpr_t = decltype(valid_expr); using iexpr_t = decltype(invalid_expr); // The assert will pass static_assert(is_same_v<vexpr_t, iexpr_t>); } Example of expression sharing the same type but representing distinct expression tree due to the lack of variable instance type identifying mechanism.

Figure 3 . 6 :

 36 Figure 3.6: Architecture diagram of fixed-point function extracting compiler.

Figure 3 . 7 :

 37 Figure 3.7: IR and optimization used in the process of lowering high-level operator specifications to implementation details.

Listing A. 1 :

 1 MArTo code for the IEEENumber class template<unsigned int WE, unsigned int WF, template<unsigned int, bool> class Wrapper> class IEEENumber : public Wrapper<WE + WF + 1, false> { private: typedef Wrapper<WE+WF+1, false> basetype; // Underlying storage type template<unsigned int W> using us_wrapper = Wrapper<W, false>; public: static constexpr unsigned int _WE = WE; static constexpr unsigned int _WF = WF; using rounding_type_t = us_wrapper<3>; template<unsigned int, bool> class Wrapper> inline PositIntermediateFormat<N, WES, Wrapper, true> posit_decoder(PositEncoding<N, WES, Wrapper> positN) { //Sign bit auto s = positN.template get<N-1>(); //First regime bit auto count_type = positN.template get<N-2>(); //Remainder auto input_shift = positN.template slice<N-3, 0>(); auto zero_NAR = input_shift.or_reduction().bitwise_or(count_type).invert(); auto is_NAR = zero_NAR.bitwise_and(s); auto is_zero = zero_NAR.bitwise_and(s.invert()); auto implicit_bit = s.invert().bitwise_and(zero_NAR.invert()); constexpr int rangeCountSize = hint::Static_Val<N-2>::_storage; auto lzoc_shifted = hint::LZOC_shift<N-2, N-2>(input_shift, count_type); auto rangeCount = lzoc_shifted.lzoc.template leftpad<rangeCountSize+1>(); auto usefulBits = lzoc_shifted.shifted.template slice<N-4, 0>(); auto fraction = usefulBits.template slice<N-4-WES, 0>(); auto neg_count = s.bitwise_xor(count_type).invert(); auto extended_neg_count = Wrapper<rangeCountSize + 1, false>::generateSequence(neg_count); ↩→ auto comp2_range_count = rangeCount.bitwise_xor(extended_neg_count); auto exponent = getExponent<N, WES>( LYON, MEMBRE DE L'UNIVERSITÉ DE LYON NOM : FORGET Date de soutenance : 29/06/2023 Prénoms : Luc Titre : Description and compilation of ad-hoc arithmetic operators in the context of High-Level Synthesis Nature : Doctorat Numéro d'ordre : 2023ISAL0046 École doctorale : Informatique & Mathématique (InfoMath) Spécialité : Informatique Résumé :

  Wide adder implementation by Full Adder chaining. . . . . . . . 1.18 Part of the schematic of the UltraScale fast carry chain . . . . . . 1.19 Schematic representation of the compilation stages from computation graph to logical FPGA primitive netlist. . . . . . . . . . . 2.1 A generic shifter (left) instantiated (right) in a floating-point adder with fraction bits . . . . . . . . . . . . . . . . . . . . . . . 2.2 Generic leading-bit counters. . . . . . . . . . . . . . . . . . . . . . . 2.3 Example of product accumulation on the carry-save banks of a Kulisch-like accumulator. . . . . . . . . . . . . . . . . . . . . . . . . 2.4 Architecture of a posit operator in a Posit Arithmetic Unit that uses posit registers and posit-to-posit operators. . . . . . . . . . . 2.5 Architecture of a posit to PIF decoder. . . . . . . . . . . . . . . . . . 2.6 Architecture of a UPIF to posit encoder. The PIF to posit encoder is similar, with the round and sticky logic (including the final adder) removed. . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 2.7 Architecture of a PIF adder. Exponent comparison block denoted with ">" also takes the operand and bits to detect zero values, but wires have been omitted here for clarity. . . . . . . . 2.8 Architecture of a PIF multiplier. . . . . . . . . . . . . . . . . . . . . 2.9 Architecture of a posit quire addition/subtraction. . . . . . . . . . 2.10 The bits of a standard quire. . . . . . . . . . . . . . . . . . . . . . . . 2.11 Architecture of a PAU using posits as a memory-only encoding, Example bipartite approximation architecture, here replacing a table of 2 6 entries with two tables of 2 4 entries. . . . . . . . . . . . 3.4 Architecture of the generator of template specializations . . . . . 3.5 Output of an additive synthesis with two frequencies of equal amplitudes and the computational error (as absolute difference on a logarithmic scale). . . . . . . . . . . . . . . . . . . . . . . . . . . 3.6 Architecture diagram of fixed-point function extracting compiler.103 3.7 IR and optimization used in the process of lowering high-level operator specifications to implementation details. . . . . . . . . . 104 Segment sizes of 32 and 64 bits). . . . . . . . . . . . . 81 2.16 Detailed synthesis results of hardware posit quire . . . . . . . . . 83 3.1 Area and timing metrics comparison between float and various FixedNumber for an additive synthesizer of 256 oscillators. II stands for Initiation Interval, the number of clock cycles that pass before the pipeline is ready to be fed a new input. . . . . . . 100 Pure function wrapping call to the slice template method. . . . 2.5 Pure function wrapping call to the slice template method with cleaner interface using C++20 concepts. . . . . . . . . . . . . . . . 2.6 Rewriting of example from listing 2.3 with proposed interface improvement. . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 2.7 MArTo code example. . . . . . . . . . . . . . . . . . . . . . . . . . . . 2.8 MArTo code example. . . . . . . . . . . . . . . . . . . . . . . . . . . . Construction of the expression tree corresponding to ( ) = log(1 + ) (lines 4-6), and construction of an operator for this function (line 7). . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 91 3.4 Simple example of template specialization and dispatching. . . . 96 3.5 Outline of the generic ArithOp template class. . . . . . . . . . . . . 97 3.6 Simplified example of generated ArithOp specialization for log(1 + ) with output format ufix(5, -17). . . . . . . . . . . . . . . 97 3.7 A simplified LNS adder . . . . . . . . . . . . . . . . . . . . . . . . . . 98 3.8 Example of expression sharing the same type but representing distinct expression tree due to the lack of variable instance type identifying mechanism. . . . . . . . . . . . . . . . . . . . . . . . . . 3.9 C++ specification of an operator approximating ↦ → sin(0.5 * * ).105 3.10 High level representation of operator of listing 3.9 using the approx dialect. . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 3.11 Implementation of listing 3.9 operator in terms of fixed-point operations described in fixedpt dialect. . . . . . . . . . . . . . . . 3.12 LLVM IR corresponding to the implementation of listing 3.9. . .

	xii 1.17 xiv 2.15 Synthesis results for a sum of 1000 products (U: Unsegmented, LIST OF FIGURES LIST OF TABLES
	S32 and S64: List of Listings

with PIF registers and PIF-to-PIF operators. . . . . . . . . . . . . . 3.1 Custom C++ HLS compiler stages . . . . . . . . . . . . . . . . . . . . 3.2 Example of function evaluator (here corresponding to listing 3.3). 3.3 1.1 Verilog description of a circuit computing ( + ) • for 12-bits inputs. . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 1.2 Chisel description of a circuit computing ( + ) • for 12-bits inputs. 1.3 HLS C++ description of a circuit computing ( + ) • . . . . . . . . 1.4 Simultaneous definition and call of an FPGA kernel computing ( + ) • . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 1.5 HLS C++ description of a circuit computing ( + ) • with 12-bits inputs. . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 1.6 Flopoco command line to produce an IEEE-754 binary32 Fused Multiply Add operator for Kintex 7 architecture. . . . . . . . . . . 2.1 Example of a counter-intuitive C intermediate result type. . . . . 2.2 Undefined behavior on shifts. . . . . . . . . . . . . . . . . . . . . . . 2.3 Example of hint code computing the leading-zero count of the sum of two 6-bits integers provided concatenated in a 12-bits vector. . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 2.4 3.1 Construction of a FixedNumber from its representation. . . . . . . 3.2 Format inference in the fixed-point library. . . . . . . . . . . . . . xv Cette thèse est accessible à l'adresse : https://theses.insa-lyon.fr/publication/2023ISAL0046/these.pdf © [L. Forget], [2023], INSA Lyon, tous droits réservés xvi LIST OF LISTINGS 3.3

Table 1 .

 1 1: Example of an arbitrary numerical format on 2-bits vectors.

	Representation Associated value
	00	-2 3
	01	2
	10	-8
	11	12

Table 1

 1

	.2: Field widths for IEEE-754 standard formats.
	Format
	binary16	5	10
	binary32	8	23
	binary64	11	52
	binary128 15 112
	binary256 19 236

Table 1 .

 1

	3: Interpretation of IEEE-754 binary16 encoding.
	Representation	Value	Denomination
	11111 1111111111
	11111	…	Quiet NaN
	11111 1000000000
	11111 0111111111		Special values
	11111	…	Signaling NaN
	11111 0000000001
	11111 0000000000	-∞
	11110 1111111111 … … 00001 0000000000	-1. • 2 -	Negative normal range
	00000 1111111111
	00000	…	-0. • 2 min	Negative
	00000 0000000001		subnormal
	00000 0000000000	-0	range
	00000 0000000000	+0
	00000 0000000001 00000 … 00000 1111111111	0. • 2 min	Positive subnormal range
	00001 0000000000 … … 11110 1111111111	1. • 2 -	Positive normal range
	11111 0000000000	+∞
	11111 0000000001
	11111	…	Signaling NaN
	11111 0111111111		Special values
	11111 1000000000
	11111	…	Quiet NaN
	11111 1111111111

Table 1 .

 1

	4: Precision -dynamic -special-value representation trade-off for
	mainstream 16 -bit IEEE-like floating-point formats.
	Format	w E w F Precision # binades Special value repr.
	Binary16	5 10	2 -10		40	3.125 %
	DLFloat	6	9	2 -9		64	≈ 0.06 %
	BFloat16	8	7	2 -7		161	≈ 0.39 %
	portion of representation dedicated to special value encoding. Figure 1.13
	gives a visual representation of the precision trade-off offered by these for-
	mats.
			1	1.01	1.02	1.02	1.03
				Binary16 DLFloat BFloat
	Figure 1.13: Representable values for mainstream 16bit IEEE-like formats on
	[1, 1 + 2 -5 ]

Table 2 .

 2 1: Synthesis of LZC Arria 10 (achieved clock target of 240MHz)

		N	ALMs FFs MLABs cycles
		26	32.5	32	0	1
	ac_int LZC	55	86.5	91	1	5
		256 465.5 710	1	8
		26	32.5	32	0	1
	hint LZC	55	86.5	91	1	5
		256 465.5 710	1	8

Table 2 .

 2 2: Synthesis of LZC and shifters on Kintex 7 (achieved target delay of 3ns). The number are obtained with Vitis 2022.2.

	Implementation	N	LUTs FFs SRLs cycles
		26	26	33	1	2
	ap_int LZC	55	71 105	1	4
		256	278 546	3	9
		26	26	33	1	2
	hint LZC	55	71 105	1	4
		256	277 546	3	9

Table 2 .

 2 3: Synthesis of shifters+sticky on Arria 10 (achieved clock target of 240MHz)

	Implementation	N ALMs FFs MLABs
	ac_int	27	134	63	2
	(shifter then sticky) 56	277 212	3
	hint	27 56 179.5 128 82 40	0 0
	Table 2.4: Synthesis of shifters+sticky on Kintex 7 (achieved target of 3ns)
	Implementation	N LUTs FFs Cycles
	ap_int	27	113 110	3
	(shifter then sticky) 56	309 234	4
	hint	27 56	84 203 133 65	3 3

Table 2 .

 2

	5: Synthesis of normalizers Arria 10 (achieved clock target of 240MHz)
	Implementation	N	ALMs	FFs MLABs cycles
		28	93	106	0	2
		57	218.5	213	0	2
	hint LZC + native shift	64	296.5	340	0	3
		256	1487 1238	13	7
		279	1603 1322	14	7
		28	88.5	72	0	1
	hint fused Shift and LZC (normalizer)	57 64 256	212 279.5 1388 1960 209 308	0 0 0	2 3 6
		279 1455.5 1544	0	4
	Table 2.6: Synthesis of normalizers on Kintex 7 (achieved target delay of 3ns)
	Implementation	N	LUTs	FFs SRLs cycles
		28	96	81	0	8
		57	222	144	0	9
	hint LZC + native shift	64	264	142	0	8
		256 1532 1045	0	11
		279 1691 1131	0	12
		28	102	122	0	4
	hint fused shift + LZC (normalizer)	57 64 256 1164 1568 254 297 292 275	0 0 0	5 6 8
		279 1958 2265	0	10

Table 2 .

 2 7: Parameters of standard posit formats.

	Standard posit smallest FP superset internal formats
	N		max			pif	upif
	8	0	6	4	5	12	14
	16	1	28	6	12	21	23
	32	2	120	8	27	38	40
	64	3	496 10	58	71	73
	regime, therefore
					= -(3 +	)

Table 2 .

 2 8: Quire bit-width parameters for standard 3.2 posit formats.

	>
							+ 2	/	00
		-					+ 4	/
			/	+ 1			shifter+sticky
							/	+ 2	/	+ 5
						+
						/	+ 3	/	/ 2	+ 5	/ 1
	/
							LZOC + Shift
		-	log 2 (	/	+ 4)	/	+ 4
	isNaR	e					Significand / 3
	isNaR						round	sticky
	Figure 2.7: Architecture of a PIF adder. Exponent comparison block denoted
	with ">" also takes the operand and bits to detect zero values, but wires
	have been omitted here for clarity.

Table 2 .

 2 10: Comparison with[START_REF] Podobas | Hardware implementation of POSITs and their application in FPGAs[END_REF] on standard posit addition and product

	Op Format	ALM DSP Cycles FMax (MHz)
	[33]

Table 2 .

 2 MArTo synthesis have been performed using Vivado HLS/Vivado 2020.1 using part xc7vx330t-ffg1157-3. Experimental settings of[START_REF] Jaiswal | PACoGen: A hardware posit arithmetic core generator[END_REF] use the same part, but tool version is not reported.

	11: Comparison with [35] on posit<32,6> addition and product
		Op LUTs DSPs Cycles Delay (ns)
	[35]	+ ×	946 854	0 1	5 6	4.1 4.4
	MArTo	+ ×	792 435	0 2	5 6	3.9 4.1

Table 2 .

 2

	12: Comparison with [50] for standard posit addition and product
		Op Format LUT DSP Delay (ns)
	[50]	+ ×	<16, 1> 383 <32, 1> 939 <16,1> 201 <32, 1> 571	0 0 1 4	27.25 35.8 20.9 29.2
	MArTo	+ ×	<16, 1> 300 <32, 1> 672 <16,1> 205 <32, 1> 472	0 0 1 4	25.5 34.5 19.2 28.8
	MArTo synthesis have been performed using Vivado HLS/Vivado 2020.1 using part
	xc7z020clg484-1.

Table 2 .

 2 13: Synthesis results of combinatorial operators

		(a) Combinatorial adder
		N	LUT (ratio)	delay (ratio)
		16	312	1.33	11.1 ns	1.27
	posit→posit 32	647	1.49	15.8 ns	1.33
		64 1550	1.59	21.6 ns	1.35
		16	237	1.01	9.7 ns	1.10
	PIF→PIF	32	562	1.29	12.9 ns	1.08
		64 1244	1.27	14.7 ns	0.92
		16	234	1	8.8 ns	1
	IEEE→IEEE 32	434	1	11.9 ns	1
		64	976	1	16.0 ns	1
		(b) Combinatorial multiplier
		N	LUT (ratio) DSP	delay (ratio)
		16	182	1.03	1	11.3 ns	1.39
	posit→posit 32	466	1.37	4	15.8 ns	1.62
		64 1213	1.58	16	21.1 ns	1.48
		16	120	0.68	1	7.8 ns	0.96
	PIF→PIF	32	291	0.86	4	11.5 ns	1.17
		64	695	0.90	16	15.3 ns	1.08
		16	176	1		1	8.1 ns	1
	IEEE→IEEE 32	340	1		2	9.8 ns	1
		64	768	1		9	14.3 ns	1
		(c) Posit -PIF conversion operators
				N LUT delay
				16	61 2.59 ns
		posit→PIF	32 106 4.74 ns
				64 278 5.52 ns
				16	41 2.12 ns
		PIF→posit	32	98 2.50 ns
				64 301	2.83

Table 2 .

 2 [START_REF] Caulfield | A cloud-scale acceleration architecture[END_REF]: Synthesis results for a sum of 1000 products (U: Unsegmented, S32 and S64: Segment sizes of 32 and 64 bits).

			LUT	Reg. DSP cycles	delay
		U 1200 1026	1	1019	2.70 ns
	quire 16	S32	978 1062	1	1021	2.68 ns
		S64 1004	958	1	1019	2.36 ns
		U 5884 6235	4	1031	3.65 ns
	quire 32 (512 bits)	S32 3641 7237	4	1040	2.89 ns
		S64 3513 5189	4	1033	2.78 ns
	Kulisch 32	S32 3624 7632	2	1034	2.937
	(559 bits)	S64 3612 5165	2	1026	2.801
	IEEE Float 32		840	711	2	6012	2.92 ns
	IEEE Float 64		1798 1723	9	8015	3.33 ns
	Xilinx Float 32		445	544	3	6008	2.72 ns
	Xilinx Float 64		809 1386 11	8013	2.70 ns

Table 2 .

 2 16: Detailed synthesis results of hardware posit quire

				LUT	Reg. Cycles Delay (ns)
			U	618	885	4	2.576
		Quire addition	S32	403	585	3	1.886
	posit16	Carry propagation	S64 S32 S64	444 6 2*	606 390 261	3 3 2	1.984 1.539 1.651
		Quire to posit		480	166	3	2.735
			U 3609 4986	7	3.212
		Quire addition	S32 1305 2265	3	2.791
	posit32	Carry propagation	S64 1389 2276 S32 281 2874 S64 189 2391	3 8 7	2.791 2.851 2.183
		Quire to posit		1845 1457	17	2.878

  .1. The LLVM IR version produced by the recent Clang/LLVM is not directly compatible with Vitis HLS for two reasons:• Vitis HLS understands only LLVM 7 IR, and LLVM 15 IR has some constructs that do not exist or are invalid in LLVM 7 IR,• it expects some IR patterns that are no longer emitted by clang++. Without them, the quality of result can degrade abruptly. For instance, the canonical way to move objects in recent LLVM is via memcpy instructions, while Vitis works better with loads and stores.

	vitis_ip.cpp	clang++	IP.bc	sycl_vxx.py	IP.zip
				IR Downgrading
				Vitis HLS
	So a few LLVM passes that downgrade modern LLVM IR to Vitis understand-
	able IR have been developed. This compiler also offers a few custom C++ dec-
	orators to replace Vitis HLS pragmas.
	A new Vitis-IP target has been added to the Xilinx SYCL compiler in order
	to allow the usage of the tool outside SYCL environments. With this target,

  Example bipartite approximation architecture, here replacing a table of 26 entries with two tables of 2 4 entries. Source:[START_REF] De Dinechin | Application-Specific Arithmetic[END_REF]

		( )							=	C	A	B
	+	+	+	+	+	+	+	+	+ constant slope + + + + + + + on 2 intervals	initial value ≈ ( )	offset ≈ × ( ) +
	(a) The bipartite approximation (before discretiza-
	tion)								(b) Bipartite architecture
	Figure 3.3:

  .

		(any)			(any)
	app.cpp	CPU	opgen.elf	operators.hpp	HLS
		Compiler			Compiler
		Sollya.so	Architecture	Produced by	app.zip
			generation	execution of
		FloPoCo.so	library (.so)	opgen.elf

  . One of these examples is also shown to integrate well in a SYCL environment.

				2 8
	2 sin(2 t) + sin(4 t)	0.75 0.50 0.25 0.00 0.25 0.50 0.75	binary32 FixedNumber 16 FixedNumber 8	2 17 2 14 2 11 2 20 2 23	k = 16 k=8	10 5 10 4 k | 10 3 FixedNumber 10 7 10 6 |binary32
			0.00 0.25 0.50 0.75 1.00 t		0.00 0.25 0.50 0.75 1.00 t

Table operations

 operations Italic optimizationsare not yet implemented. Dashed IR frames are standard MLIR dialects.

			Bitheap compression
	LLVM	arith	Tiled multiplers
	Integer operations	Integer operations	Multiple Constant
	Memory references	Table operations	Multiplication
			…
			FloPoCo.so

  : Synthèse de haut niveau, formats numériques, IEEE-754, Approximation de fonctions, compilation, Posit Laboratoire(s) de recherche : Centre of Innovation in Telecommunications and Integration of Service (CITI)

	Directeur de thèse : Florent DE DINECHIN
	Président du jury : Fabrice RASTELLO
	Composition du jury :
	Roselyne CHOTIN (Examinatrice)
	Steven DERRIEN (Rapporteur)
	Fabrizio FERRANDI (Rapporteur)
	Anastasia VOLKOVA (Examinatrice)

The interested reader can discover the (terrible ((Ph.D. student) universe)) in the abundant related literature on the web.

Cette thèse est accessible à l'adresse : https://theses.insa-lyon.fr/publication/2023ISAL0046/these.pdf © [L. Forget], [2023], INSA Lyon, tous droits réservés

≤ ≤ -1 . Cette thèse est accessible à l'adresse : https://theses.insa-lyon.fr/publication/2023ISAL0046/these.pdf © [L. Forget], [2023], INSA Lyon, tous droits réservés

Cette thèse est accessible à l'adresse : https://theses.insa-lyon.fr/publication/2023ISAL0046/these.pdf © [L. Forget], [2023], INSA Lyon, tous droits réservés Cette thèse est accessible à l'adresse : https://theses.insa-lyon.fr/publication/2023ISAL0046/these.pdf © [L. Forget], [2023], INSA Lyon, tous droits réservés

It can make sense if the iteration space is known at compile time, and the slice indices depends only on it. In this case, the loop can be spatially unrolled with each iteration instance having bus of the required size. The loop unrolling can be required by a specific pragma in Vitis HLS. However, the intent for this kind of constructs is more clearly expressed with a proper template unrolling.Cette thèse est accessible à l'adresse : https://theses.insa-lyon.fr/publication/2023ISAL0046/these.pdf © [L. Forget], [2023], INSA Lyon, tous droits réservés

For data parallel C++. Cette thèse est accessible à l'adresse : https://theses.insa-lyon.fr/publication/2023ISAL0046/these.pdf © [L. Forget], [2023], INSA Lyon, tous droits réservés

Which means they can be evaluated at compile time and used inside template parameters Cette thèse est accessible à l'adresse : https://theses.insa-lyon.fr/publication/2023ISAL0046/these.pdf © [L. Forget], [2023], INSA Lyon, tous droits réservés

As the library uses languages features that are not supported by IntelHLS or Vitis HLS, it makes no sense to add genericity to function libraries to support these backends. However, this is a technical decision for the sake of efficiency of library development. It should be emphasized again that there is no conceptual obstacle to build a similar library using pre-C++14 constructs. This is however more tedious and error-prone, hence the technical decision of using C++20 in this experimental work.

[START_REF] Suzuki | Building modern javascript engine[END_REF] Another technical point is that in their current specification, 1-bit signed _BitInt are not supported natively. While this has been "fixed" by adding this support at the compiler level in the case of the custom compiler presented here, it will also be handled at the hint library Cette thèse est accessible à l'adresse : https://theses.insa-lyon.fr/publication/2023ISAL0046/these.pdf © [L. Forget], [2023], INSA Lyon, tous droits réservés

It is also possible to create it by specifying the exact type. Cette thèse est accessible à l'adresse : https://theses.insa-lyon.fr/publication/2023ISAL0046/these.pdf © [L. Forget], [2023], INSA Lyon, tous droits réservés

A last emphasis is made here to highlight that, while it means supporting C++20 with the proposed library, it is not a strong requirement. Indeed, the concept mechanisms used in the library can be implemented by exploiting the C++ SFINAE (Substitution Failure Is Not An Error) rule, and the constexpr functions used can be implemented manually with C++11 constructs.Cette thèse est accessible à l'adresse : https://theses.insa-lyon.fr/publication/2023ISAL0046/these.pdf © [L. Forget], [2023], INSA Lyon, tous droits réservés

The free library function.Cette thèse est accessible à l'adresse : https://theses.insa-lyon.fr/publication/2023ISAL0046/these.pdf © [L. Forget], [2023], INSA Lyon, tous droits réservés

Here trigonometric identities could be used to perform argument range reduction. See[START_REF] De Dinechin | Fixed-point trigonometric functions on FP-GAs[END_REF] for an illustration of this technique. The long-term objective of the library is that such optimizations are performed transparently to the user, so this optimization technique is not relevant in user code here.Cette thèse est accessible à l'adresse : https://theses.insa-lyon.fr/publication/2023ISAL0046/these.pdf © [L. Forget], [2023], INSA Lyon, tous droits réservés

With these two improvement, the toy operator of listing 2.3 could be expressed as in listing 2. Concepts have not been used due to the lack of support of C++20 standard by the targeted HLS toolchains. Using functional wrappers to improve the readability could have been done earlier, but the idea has only been identified recently, which explains why it is not implemented.

Evaluation

Usage of hint in real applications (floating-point operators) will be presented in section 2.2. Current section evaluates the individual building blocks provided by the library.

The evaluation is divided in two parts. The first consists in ensuring that the library does not add overhead compared to native library usage. The second part evaluates combined operators, ensuring that they indeed reduce resources consumption and/or latency compared to their sequential counterparts.

All the results presented section are given after place-and-route. At the time of evaluation, Vivado HLS (the ancestor of Vitis HLS) 2018.1, the latest available version, was quite unstable, with numerous crashes and silent production of incorrect hardware at rare occasions. That is why the more stable Vivado HLS 2016.3 version was used at this time. Most of the results presented in this section are obtained with this version. Experiments based on this library and using more recent tool version are presented in section 2.3.1. These experiment results seem to indicate that conclusions obtained with this old tool version still hold with recent tool. Nevertheless, in order to get recent numbers, the experiment on LZC comparison for Xilinx toolchain has been re-run using Vitis HLS 2022.2, the most recent version at time of writing. The FPGA architecture targeted for experiments on Xilinx HLS toolchain is Kintex 7. For the Intel toolchain, IntelHLS 19.1 is used to target Arria 10.

A narrow LNS adder

As seen in section 1.2.3, sum and subtraction in LNS format requires evaluating the function ⊕ : ↦ → log 2 (1 + 2 ). This function is a good example of composite functions that can be implemented using the techniques described in this chapter.

In order to demonstrate the capabilities of the library, a simple LNS adder has been developed. The simplified adder is presented in listing 3. The LNS Adder has been synthesized for an lns_t having an exponent of FixedFormat<5,-6, unsigned>. This requires to evaluate ⊕ with an accuracy constraint of 2 -8 . The report after synthesis and place-and-route shows that the design uses two BRAM, which matches the expectations.

Additive sound synthesis

Additive synthesis is a sound synthesis technique typically used in electronic music to create timbres by adding sine waves together [START_REF] Douglas | The electrical production of music[END_REF] as in the following computation :

It allows very rich sounds when a lot of sines are used, but then it is very compute-intensive. When used in an interactive configuration with hardware in the loop to simulate real instruments interacting in real-time with the real world physics, it requires very low latency. In order to check the impact of the library, an additive synthesis kernel was developed both with FixedNumber and binary32 computation. Both versions %2 = approx.variable %1 : <0, -4, s> %3 = approx.math "pi"() %cst = approx.constant <2, <2, 0, u>, "2.0"> %4 = approx.math "mul"(%2, %3) %5 = approx.math "div"(%4, %cst) %6 = approx.math "sin"(%5) %7 = approx.evaluate basic_poly of %6 as <0, -4, s> Listing 3.10: High level representation of operator of listing 3.9 using the approx dialect.

to the fixedpt dialect. This pass interacts with the external tools FloPoCo and Sollya to determine this plan, and generate the corresponding fixedpt IR. Listing 3.11 gives the low-level implementation description of the operator from listing 3.9 using this dialect. This dialect still exposes some level of abstraction. For instance, in this dialect addition or product are variadic operations. This leaves the implementation detail decision to be decided in lower-level dialects. Compiler support for fixed-point functions in HLS Dialects used for low-level implementation description are standard MLIR dialects. Fixed-point operations are lowered to integer operations represented using the arith dialect. Current state of conversion from fixedpt to arith dialect mainly consists in adding the good shifts and truncations to ensure good relative positioning of fixed point operands. It also handles the rounding logic of fixed point ops (with nearest tie breaking rule being towars +∞ in this case). This dialect is then converted to the llvm dialect and converted to LLVM IR that is linked with the rest of the application before starting the HLS process. As the llvm and arith dialect are very similar, and the former is a one to one mapping of LLVM IR, only the resulting LLVM IR for the example operator is reported on listing 3.12.

%4 = sext i5 %3 to i10 ; extend variable %3 which has type i5 (5-bit integer)

; to i10 (10-bit integer) %5 = mul i10 %4, -9

Listing 3.12: LLVM IR corresponding to the implementation of listing 3.9.

The tensor and memref standard dialects are used to represent the value tables at high and low abstraction level respectively, for table based approximation methods.

IEEE-754 Adder

A.2 Posit adder

In the posit case, two helper classes exists that play the same role as IEEENumber. The PositEncoding provide useful methods related to posit encodings. The PositIntermediateFormat is useful to manipulate posit intermediate format (PIF) values (see section 2.2.3 for more information). These class are not reported here as they are very similar to IEEENumber. Before performing the addition, posit encodings are converted to PIF using a posit decoder (presented on listing A.3). Then the two PIFs can be added, using the adder code from listing A.5. Finally, the result has to be rounded back using a posit encoder, presented on listing A.