# An Analytical Study of a Two-Sided Mobility Game

Ioannis Vasileios Chremos, Student Member, IEEE, and Andreas A. Malikopoulos,
Senior Member, IEEE This research was supported by the Sociotechnical Systems
Center (SSC) at the University of Delaware.The authors are with the Department
of Mechanical Engineering, University of Delaware, Newark, DE 19716 USA.
<EMAIL_ADDRESS>
###### Abstract

In this paper, we consider a mobility system of travelers and providers, and
propose a “mobility game” to study when a traveler is matched to a provider.
Each traveler seeks to travel using the services of only one provider, who
manages one specific mode of transportation (e.g., car, bus, train, bike). The
services of each provider are capacitated and can serve up to a fixed number
of travelers at any instant of time. Thus, our problem falls under the
category of many-to-one assignment problems, where the goal is to find the
conditions that guarantee the stability of assignments. We formulate a linear
program of maximizing the social welfare of travelers and providers and show
how it is equivalent to the original problem and relate its solutions to
stable assignments. We also investigate our results under informational
asymmetry and provide a “mechanism” that elicits the information of travelers
and providers. Finally, we investigate and validate the advantages of our
method by providing a numerical simulation example.

##  I INTRODUCTION

Commuters in big cities have continuously experienced the frustration of
congestion and traffic jams. Travel delays, accidents, and altercations have
consistently impacted the economy, society, and the natural environment by
playing a decisive role in the increase of idling vehicle engines on city
roads [1]. In addition, one of the pressing challenges of our time is the
increasing demand for energy, which requires us to make fundamental
transformations in how our societies use and access transportation. Thanks to
evolutionary developments that are currently afoot, it is highly expected that
we will be able to eliminate congestion entirely and significantly increase
mobility efficiency in terms of fuel consumption and travel time [2]. Self-
driving cars offer a most intriguing opportunity that will enable us to travel
safely and efficiently anywhere and anytime [3]. Several studies have shown
the benefits of _emerging mobility systems_ (e.g., ride-hailing, on-demand
mobility services, shared vehicles, autonomous vehicles) to reduce energy and
alleviate traffic congestion in a number of different transportation scenarios
[4, 5, 6, 7]. One question though that still remains unanswered is: _Can we
develop an efficient mobility system that can enhance accessibility while
controlling the ratio of travel demand over capacity and guarantee the welfare
of all travelers?_

Recently it was shown that when daily commuters were offered a convenient and
affordable taxi service for their travels, a change of behavior was noticed,
namely these commuters ended up using cars more often compared to when they
drove their own car [8]. Along with other studies [9, 10] this shows us that
in emerging mobility systems the travelers’ tendency to travel will drive
travelers to use cars more often and shift them away from public transit. So,
in this paper to address this problem, we study the game-theoretic
interactions of travelers seeking to travel and getting assigned to providers
(e.g., technology and transportation companies) where each operates one mode
of transportation. Inspired by the Mobility-as-a-Service concept, we consider
a system of multimodal mobility that handles user-centric information and
provides travel services (e.g., navigation, location, booking, payment) to a
number of travelers. The goal of such a mobility system is to guarantee
mobility as a seamless service across all modes of transportation accessible
to all. For our purposes, we consider a game to model the strategic and
economic interactions in a mobility system with two groups of agents, a
traveler and a provider group, where both have preferences and the objective
is to assign each traveler to only one provider.

One of the standard approaches to alleviate congestion in a transportation
system has been the management of demand size due to the shortage of space
availability and scarce economic resources in the form of congestion pricing
(alternatively called “tolling mechanisms” in [11]). Such an approach focuses
primarily on intelligent and scalable traffic routing, in which the objective
is to guide and coordinate decision-makers in path-choice decision-making.
Interestingly, by adopting a game-theoretic approach, advanced systems have
been proposed to assign users concrete routes or minimize travel time and
study the Nash equilibria under different tolling mechanisms [12, 13, 14, 15,
16, 17].

Partly related to our work are matching models which describe systems or
markets in which there are agents of disjoint groups and have preferences
regarding the “goods” of the opposite agent they associate with. Notable
examples are mechanisms for assigning students to schools [18], pickup and
delivery [19], electric vehicles [20]. It is easy to see that matching markets
are quite practical as they offer insights into the more general economic and
behavioral real-life situations. These examples are all centralized approaches
of determining who gets assigned to whom at what cost and benefit. One of the
very first studies was the marriage model which was analyzed by [21] and
existence of stable matchings between men and women was established. The
authors in [22] extended it by incorporating monetary transactions between the
agents to the marriage model and formulated it to the well-known “assignment
game.” They showed that there exists a set of stable assignments, called the
_core_ (no agent wants to deviate from their match) and it is identical to the
solutions of a dual linear program. However, no explicit mechanism was offered
on how to achieve a stable assignment in the core. Thanks to the natural
usefulness of matching markets, various extensions of the assignment game have
been developed focusing either on different behavioral settings or information
structures [23, 24].

The main contribution of this paper is the development of a game-theoretical
framework to study the economic interactions of travelers and providers in a
two-sided many-to-one assignment game. By two-sided we mean that we consider
the preferences of both the travelers and the providers. By many-to-one we
mean that we impose constraints of how many travelers can be assigned to one
provider and how many providers to one traveler. Our analysis can be divided
into two parts. First, we use linear programming arguments to showcase the
existence of an optimal assignment between travelers and providers that is
also stable, i.e., no one will seek to deviate from their match. Second, we
consider an asymmetric informational structure, where no traveler/provider is
expected to provide their private information willingly. We provide a pricing
mechanism (Algorithm 1) for this case and show how we can successfully elicit
the private information while also ensure efficiency (maximization of social
welfare).

The remainder of the paper is structured as follows. In Section II, we present
the mathematical formulation of the proposed mobility game which forms the
basis of our theoretical study for the rest of the paper. In Section III, we
derive the theoretical properties of the mobility game and, in Section IV, we
validate the game with a numerical simulation. Finally, in Section V, we draw
concluding remarks and offer a discussion of some future research directions.

##  II MODELING FRAMEWORK

###  II-A The “Mobility Game” Formulation

We consider a mobility system of two finite, disjoint, and non-empty groups of
agents of which one represents the travelers and the other the providers. We
denote the set of travelers by $\mathcal{I}$, $|\mathcal{I}|=I\in\mathbb{N}$
and the set of providers by $\mathcal{J}$, $|\mathcal{J}|=J\in\mathbb{N}$. In
a typical mobility system, we expect to have more travelers than available
providers, so $I\gg J$. Each provider represents a company (e.g., Uber, Lyft,
Amtrak, DART, Lime) that manages a fleet of vehicles (cars, trains, busses,
bicycles). We focus our study in static settings, in this paper, thus, each
provider can serve up to a fixed number of travelers within a fixed time
period. For example, in a generic city neighborhood, on any given weekday
morning, there are at most a certain number of ride-sharing vehicles available
(Uber/Lyft). Formally, for each provider $j\in\mathcal{J}$, we impose a
physical _traveler capacity_ , denoted by $\varepsilon_{j}\in\mathbb{N}$.
Naturally, each provider can serve a different number of travelers, so we
expect $\varepsilon_{j}$ to vary significantly. For example, a train company
can provide travel services per hour to hundreds of travelers compared to a
bikeshare company in a city. Next, travelers seek to travel using the services
of at most one provider. We do not focus our modeling in routing or path-
allocation (such problems have been studied extensively [25, 26]), rather we
are interested in an optimal collective assignment of travelers to providers.
Both travelers and providers have preferences and can be characterized by
their type; thus, this is a two-sided mobility game.

######  Remark 1.

Without loss of generality, we expect the aggregate travel demands of all the
travelers to be exactly met by the aggregate capacities of all the providers’
mobility services. Thus, we have
$\sum_{j\in\mathcal{J}}\varepsilon_{j}=|\mathcal{I}|=I$.

######  Remark 2.

Intuitively, via a smartphone app, travelers book in advance for their travel
needs and report their preferences and request a travel recommendation (which
provider to use). The app collects all requests from specific neighborhoods at
a fixed time, and then assigns each traveler to a provider by taking into
account both the traveler’s as well as the provider’s preferences.

######  Definition 1.

The _traveler-provider assignment_ is a vector
$\mathbf{X}=(x_{11},\dots,x_{ij},\dots,x_{IJ})=(x_{ij})_{{i\in\mathcal{I}},j\in\mathcal{J}}$,
where $x_{ij}$ is a binary variable of the form:

$x_{ij}=\begin{cases}1,\;&\text{if $i\in\mathcal{I}$ is assigned to
$j\in\mathcal{J}$},\\\ 0,\;&\text{otherwise}.\end{cases}$ (1)

We call $x_{ij}$ the _mobility outcome_ of each traveler $i$ and each provider
$j$ and denote by $\mathcal{X}$ the set of such outcomes.

######  Definition 2.

For any traveler $i\in\mathcal{I}$,
$\theta_{i}=\max_{j\in\mathcal{J}}\,\\{\theta_{ij}\\}\in\Theta_{i}$, where
$\theta_{ij}\in[0,1]$, is traveler $i$’s _personal predisposition_ of provider
$j\in\mathcal{J}$.

We denote by $\theta_{-i}$ for the personal predisposition profile of all
travelers except traveler $i$. Intuitively, a traveler might have a great
affinity towards a taxicab service and a lower affinity towards a bus service.
So, we expect different travelers to have different preferences on the mode of
transportation to use.

Next, we represent the preferences of each traveler with a utility function
consisted of two parts: the traveler’s valuation of the mobility outcome and
the associated payment required for the realization of the outcome. In other
words, any traveler is expected to pay a toll or ticket fee for the services
of a provider.

######  Definition 3.

Each traveler $i$’s preferences are summarized by a utility function
$u_{i}:\mathcal{X}\times\Theta_{i}\to\mathbb{R}$ that determines the monetary
value of the overall payoff realized by traveler $i$ from their assignment to
provider $j$. Let $t_{ij}\in[\underline{t},\bar{t}]\subset\mathbb{R}$ denote
traveler $i$’s mobility payment. Thus, traveler $i$ receives a total utility
in the form

$u_{i}(x_{ij},\theta_{i})=v_{i}(x_{ij},\theta_{i})-t_{ij},$ (2)

where $v_{i}:\mathcal{X}\times\Theta_{i}\to\mathbb{R}_{\geq 0}$ is a linear
valuation function that represents the maximum amount of money that traveler
$i$ is willing to pay for the mobility outcome $x_{ij}$.

######  Remark 3.

If for any traveler $i\in\mathcal{I}$, we have $x_{ij}=0$ for all
$j\in\mathcal{J}$, then $u_{i}=0$. Naturally, this means that for any traveler
$i$ with $x_{ij}=0$ for all $j\in\mathcal{J}$ we have $t_{ij}=0$.

On similar lines, we define the providers’ utility function.

######  Definition 4.

A provider $j$’s utility is given by

$u_{j}(x_{ij},\delta_{j})=t_{ij}-c_{j}(x_{ij},\delta_{j}),$ (3)

where $\delta_{j}\in(0,1]$ represents the type of provider $j$, and $c_{j}$ is
a linear cost function related to the operation of the mobility services
provided by $j\in\mathcal{J}$. We denote by $\delta_{-j}$ for the type profile
of all providers except provider $j$.

######  Remark 4.

Intuitively, $\delta_{j}$ can be interpreted as the “operational value” of
provider $j$ for the mobility services it provides and operates. In other
words, the monetary value of the entire process of its service to serve a
traveler on a given location and time.

In both (2) and (3), the “payment component” $t_{ij}$ is not expected to
dominate either the traveler’s or the provider’s utility function. This is
because $t_{ij}$ have an alternate sign in (2) and (3), so a high value (or
low) can lead to negative utility for the travelers (or the providers) leading
to a unfavorable match between traveler $i$ and provider $j$. We will see
later in the paper how we can ensure unfavorable matchings do not happen.

######  Definition 5.

Under the assignment $x_{ij}$ of traveler $i$ and provider $j$, their
_mobility $(i,j)$-matching payoff_ is given by

$a_{ij}(x_{ij})=u_{i}(x_{ij},\theta_{i})+u_{j}(x_{ij},\delta_{j}),$ (4)

where $a_{ij}$ measures the combined payoff or benefit measured in monetary
units of traveler $i$ being assigned to provider $j$.

######  Remark 5.

By Remark 3, if $x_{ij}=0$, then $a_{ij}=0$.

######  Definition 6.

The _utility assignment matrix_ $\mathbf{A}$ is constructed with
$|\mathcal{I}|$ rows and $|\mathcal{J}|$ columns and each entry represents the
$(i,j)$-matching utility $a_{ij}$ between traveler $i$ and provider $j$ for
all $i\in\mathcal{I}$ and all $j\in\mathcal{J}$.

Based on Definition 6, we can construct matrix $\mathbf{A}$ as follows:

$\mathbf{A}=\begin{bmatrix}a_{11}&a_{12}&a_{13}&\dots&a_{1J}\\\
a_{21}&a_{22}&a_{23}&\dots&a_{2J}\\\ \vdots&\vdots&\vdots&\ddots&\vdots\\\
a_{I1}&a_{I2}&a_{I3}&\dots&a_{IJ}\end{bmatrix}.$ (5)

The mobility game of travelers and providers is a collection of four objects,
namely set of agents, vector of assignments, matrix of utilities, and a vector
of capacities. Formally, we state the next definition.

######  Definition 7.

The mobility game can be fully characterized by the tuple
$\mathcal{M}=\langle\mathcal{I}\cup\mathcal{J},\mathbf{X}=(x_{ij})_{i\in\mathcal{I},j\in\mathcal{J}},\mathbf{A},(\varepsilon_{j})_{j\in\mathcal{J}}\rangle$.

######  Definition 8.

A _feasible_ assignment is a vector
$\mathbf{X}=(x_{ij})_{{i\in\mathcal{I}},j\in\mathcal{J}}$,
$x_{ij}\in\\{0,1\\}$ that satisfies constraints

$\displaystyle\sum_{j\in\mathcal{J}}x_{ij}$ $\displaystyle\leq 1,\quad\forall
i\in\mathcal{I},$ (6) $\displaystyle\sum_{i\in\mathcal{I}}x_{ij}$
$\displaystyle\leq\varepsilon_{j},\quad\forall j\in\mathcal{J},$ (7)

where (6) ensures that each traveler $i\in\mathcal{I}$ is assigned to only one
mobility service $j\in\mathcal{J}$, and (7) ensures that the traveler capacity
of each provider $j$ is not exceeded while its services are shared by multiple
travelers. An _optimal_ assignment is a feasible assignment
$(x_{ij})_{{i\in\mathcal{I}},j\in\mathcal{J}}$ such that

$\sum_{i\in\mathcal{I}}\sum_{j\in\mathcal{J}}a_{ij}(x_{ij})\geq\sum_{i\in\mathcal{I}}\sum_{j\in\mathcal{J}}a_{ij}(x_{ij}^{\prime}),$
(8)

for all feasible assignments $x_{ij}^{\prime}$.

######  Definition 9.

A feasible assignment
$\mathbf{X}=(x_{ij})_{{i\in\mathcal{I}},j\in\mathcal{J}}$,
$x_{ij}\in\\{0,1\\}$ is _stable_ if there exist non-negative vectors
$\phi=(\phi_{i})_{i\in\mathcal{I}}$ and $\psi=(\psi_{j})_{j\in\mathcal{J}}$
such that

$\sum_{i\in\mathcal{I}}\phi_{i}+\sum_{j\in\mathcal{J}}\varepsilon_{j}\cdot\psi_{j}=\sum_{i\in\mathcal{I}}\sum_{j\in\mathcal{J}}a_{ij}(x_{ij})$
(9)

with $\phi_{i}+\psi_{j}\geq a_{ij}$ for all $i\in\mathcal{I}$ and all
$j\in\mathcal{J}$.

We will see later in Section III the mathematical and physical interpretation
of $\phi$ and $\psi$.

######  Definition 10.

Let $(t_{ij}^{*})_{i\in\mathcal{I},j\in\mathcal{J}}$ denote the mobility
payments associated with the stable assignment, denoted by
$(x_{ij}^{*})_{{i\in\mathcal{I}},j\in\mathcal{J}}$. Then the equilibrium
$(x_{ij}^{*},t_{ij}^{*})_{i\in\mathcal{I},j\in\mathcal{J}}$ is called an
_ideal-mobility_ equilibrium.

From Definition 9, it is easy to see that Definition 10 implies that an ideal-
mobility equilibrium in mobility game $\mathcal{M}$ ensures that (i) providers
are assigned to travelers up to their capacity (thus, maximizing revenue), and
(ii) travelers receive the best-possible utility being assigned to a provider
(thus, maximizing welfare).

######  Assumption 1.

Every aspect of the mobility game $\mathcal{M}$ is considered known
information to every traveler and provider.

Assumption 1 seems strong but it will prove instrumental in Section III. Our
analysis will focus on how to show existence, optimality, and stability of
traveler-provider assignments and then in Subsection III-B, we will relax
Assumption 1 and show how we can elicit the private information of both
travelers and providers.

###  II-B The Optimization Problem

In the mobility game $\mathcal{M}$, we are interested to know what are its
stable assignments (alternatively called stable equilibria), whether they
exist and under what conditions.

######  Problem 1.

The maximization problem of $\mathcal{M}$ is

$\displaystyle\max_{x_{ij}}\sum_{i\in\mathcal{I}}\sum_{j\in\mathcal{J}}a_{ij}(x_{ij}),$
(10) $\displaystyle\text{subject
to:}\;\eqref{Constraint:Problem1-First},\eqref{Constraint:Problem1-Second},$

where $x_{ij}\in\\{0,1\\}$ for all $i\in\mathcal{I}$ and all
$j\in\mathcal{J}$.

We can relax the binary variable constraint to a non-negativity constraint
variable in Problem 1. We will show in the next section that this does not
affect the optimal solutions of Problem 1 as we can ensure all optimal
solutions of the equivalent linear program are binary valued.

######  Problem 2 (Linear Program).

The linear program formulation of mobility game $\mathcal{M}$ is

$\displaystyle\max_{x_{ij}}\sum_{i\in\mathcal{I}}\sum_{j\in\mathcal{J}}a_{ij}(x_{ij})$
(11) $\displaystyle\text{subject
to:}\;\eqref{Constraint:Problem1-First},\eqref{Constraint:Problem1-Second},\text{
and}$ $\displaystyle x_{ij}\geq 0,\quad\forall i\in\mathcal{I},\quad\forall
j\in\mathcal{J},$ (12)

where (12) transforms the (binary) assignment problem to a (continuous) linear
program of which $x_{ij}$ can be interpreted as the probability that traveler
$i$ is matched to provider $j$.

##  III ANALYSIS AND PROPERTIES OF THE MOBILITY GAME

###  III-A Existence, Optimality, and Stability of Assignments

In this subsection, we show that for Problems 1 and 2 at least one optimal
solution exists (thus, ensuring stability).

######  Theorem 1.

The stable assignments of mobility game $\mathcal{M}$ are the same with the
optimal solutions of Problem 1. Furthermore, the set of optimal solutions of
Problem 1 is non-empty.

###### Proof.

By relaxing the binary constraint of Problem 1, we get a linear program
(Problem 2). Its set of all real-valued solutions is a polytope whose vertices
have all integer-valued coordinates. Since the solutions are also guaranteed
to be non-negative, the set of solutions is non-empty [27]. Thus, Problem 2
has at least one solution with integer components (in our case 0-1
components). Hence, the set of all optimal solution of Problem 1 is non-empty
[27]. By Definition 8, assignments are stable as long as no agent in
$\mathcal{I}\cup\mathcal{J}$ has an incentive (e.g., higher utility) to break
their matching pair. So, finding a stable assignment is equivalent to finding
the best in terms of aggregate utility among all possible feasible
assignments. Mathematically, (8) naturally leads to a maximization problem.
Therefore, the existence of a stable assignment of mobility game $\mathcal{M}$
is guaranteed. ∎

Next, we derive the dual of Problem 2.

######  Problem 3.

The dual of Problem 2 is given below:

$\displaystyle\min_{\phi,\psi}\sum_{i\in\mathcal{I}}\phi_{i}+\sum_{j\in\mathcal{J}}\varepsilon_{j}\cdot\psi_{j},$
(13) subject to: $\displaystyle\phi_{i}+\psi_{j}\geq a_{ij},\quad\forall
i\in\mathcal{I},\quad\forall j\in\mathcal{J},$ (14) $\displaystyle\phi_{i}\geq
0,\quad\forall i\in\mathcal{I},$ (15) $\displaystyle\psi_{i}\geq
0,\quad\forall j\in\mathcal{J},$ (16)

where $\phi$ is a $|\mathcal{I}|$-dimensional vector and $\psi$ is a
$|\mathcal{J}|$-dimensional vector.

Our objective is to establish a method for the mobility game $\mathcal{M}$’s
stable assignments by solving Problem 2. In turn, we want to solve Problem 3
to find the stable assignments. This is possible only if we can guarantee
strong duality (satisfying the conditions of complementary slackness).
Formally, a feasible assignment $x_{ij}$ and a feasible solution $(\phi,\psi)$
are optimal if and only if

$\sum_{i\in\mathcal{I}}\sum_{j\in\mathcal{J}}a_{ij}(x_{ij})=\sum_{i\in\mathcal{I}}\phi_{i}+\sum_{j\in\mathcal{J}}\varepsilon_{j}\cdot\psi_{j}.$
(17)

The conditions that guarantee optimality are given by the theorem of
complementary slackness, i.e.,

$\displaystyle\phi_{i}+\psi_{j}-a_{ij}$ $\displaystyle=0,\quad\forall
i\in\mathcal{I},\quad\forall j\in\mathcal{J},$ (18)
$\displaystyle\sum_{j\in\mathcal{J}}(x_{ij}-1)\cdot\phi_{i}$
$\displaystyle=0,\quad\forall i\in\mathcal{I},$ (19)
$\displaystyle\sum_{i\in\mathcal{I}}(x_{ij}-\varepsilon_{j})\cdot\psi_{j}$
$\displaystyle=0,\quad\forall j\in\mathcal{J}.$ (20)

######  Lemma 1.

The set of solutions of Problem 3 is non-empty and convex.

###### Proof.

We have already established that Problem 2 has at least one solution. Thus, it
follows easily that Problem 3 has at least one solution too. Any solution of
Problem 3 has a specific structure due to the geometry of the constraint set
(14) - (16). Since at least one solution will be in 0-1 components, the
constraints will force this solution to be in at a corner of a polyhedra.
Thus, the set of solutions of Problem 3 is non-empty and has to be convex. ∎

######  Remark 6.

Intuitively, a dual solution $(\phi,\psi)$ can be seen as a method to share
the “gains of mobility” among travelers and providers at an ideal-mobility
equilibrium (see Definition 10). For example, component of vector $\phi$
describes the realized gain of traveler $i$ when assigned to provider $j$
(thus enjoying the mobility services of provider $j$). A component of vector
$\psi$ describes the per unit gain of provider $j$.

######  Corollary 1.

The set of solutions of Problem 3 is a compact subset of
$\mathbb{R}^{|\mathcal{I}|}\times\mathbb{R}^{|\mathcal{J}|}$.

###### Proof.

By Lemma 1 and Remark 6, it is straightforward to show that the set of
solutions of Problem 3 is compact. ∎

######  Corollary 2.

There always exists at least one profile of mobility payments
$(t_{ij})_{i\in\mathcal{I},j\in\mathcal{J}}$ under assignment
$(x_{ij})_{i\in\mathcal{I},j\in\mathcal{J}}$.

###### Proof.

By definition of the mobility game $\mathcal{M}$ for any (feasible) assignment
$(x_{ij})_{i\in\mathcal{I},j\in\mathcal{J}}$, there must be an associated
profile of mobility payments $(t_{ij})_{i\in\mathcal{I},j\in\mathcal{J}}$. ∎

Next, we show that the existence of an optimal profile of mobility payments
$(t_{ij})_{i\in\mathcal{I},j\in\mathcal{J}}$ can be guaranteed by the
formulation of the dual program of Problem 2 and the computation of its
solutions.

######  Theorem 2.

There exists an optimal profile of mobility payments
$(t_{ij}^{*})_{i\in\mathcal{I},j\in\mathcal{J}}$ under stable assignment
$(x_{ij}^{*})_{i\in\mathcal{I},j\in\mathcal{J}}$. Furthermore, we must have
$\phi_{i}=u_{i}$ and $\psi_{j}=u_{j}$ for all $i\in\mathcal{I}$ and all
$j\in\mathcal{J}$.

###### Proof.

Suppose $\mathbf{x}^{*}=(x_{ij}^{*})_{i\in\mathcal{I},j\in\mathcal{J}}$ is a
stable assignment for mobility game $\mathcal{M}$. Under $\mathbf{x}^{*}$, we
can calculate $\sum_{i\in\mathcal{I}}\sum_{j\in\mathcal{J}}a_{ij}(x_{ij})$,
which by the theorem of strong duality and the definition of stability, (17)
holds true. This is because (18) - (20) are equivalent to the conditions that
ensure stability. Thus, there exist vectors $(\phi^{*},\psi^{*})$ from Problem
3 that feasible and optimal. By Definition 10, it follows that the mobility
payments associated with the optimal assignments of travelers and providers
are essentially the same with the optimal solutions of Problems 2 and 3.
Therefore, by the established existence of solutions to Problems 2 and 3 as
long as there exists a stable assignment
$(x_{ij}^{*})_{i\in\mathcal{I},j\in\mathcal{J}}$ an optimal profile of
mobility payment $(t_{ij}^{*})_{i\in\mathcal{I},j\in\mathcal{J}}$ must exist.
By Remark 6 and Definition 9 at an optimal assignment we have $\phi_{i}=u_{i}$
and $\psi_{j}=u_{j}$ for all $i\in\mathcal{I}$ and all $j\in\mathcal{J}$. ∎

###  III-B Asymmetric Information in the Mobility Game

So far, we have implicitly assumed that both the travelers and providers have
complete information of the entire information structure of the mobility game
$\mathcal{M}$. In other words, each traveler knows every other travelers’ and
providers’ information, i.e., travelers know each others’ utilities and
valuations, providers know each other providers’ types and cost functions. In
a realistic setting, this implicit assumption is unreasonably restrictive.
Thus, for the rest of the paper, we focus on a “mechanism” that _induces the
mobility game_ $\mathcal{M}$ by eliciting the private information of all the
travelers and providers. First, we relax Assumption 1 and consider that the
types of travelers, i.e., $\theta=(\theta_{i})$ and of providers, i.e.,
$\delta=(\delta_{j})$ are private information, i.e., known only to themselves.
Next, we denote by $\mathbf{X}_{-i}$ the assignment of travelers in
$\mathcal{I}\setminus\\{i\\}$ to providers in $\mathcal{J}$. Similarly, we
denote by $\mathbf{X}_{-j}$ the assignment of travelers in $\mathcal{I}$ to
providers in $\mathcal{J}\setminus\\{j\\}$. Furthermore, we assume that
travelers are charged by the mechanism, say $t_{i}\in\mathbb{R}$, and
providers are compensated by the mechanism, say $t_{j}\in\mathbb{R}$. The
proposed mechanism ensures to collect all funds from the travelers and
compensate accordingly the providers.

Data:
$\mathcal{I},\mathcal{J},(\theta_{i})_{i\in\mathcal{I}},(\delta_{j})_{j\in\mathcal{J}}$

Result: $\mathbf{x}^{*}$ and $(t_{ij})_{i\in\mathcal{I},j\in\mathcal{J}}$

Define the valuation functions of every traveler and provider and use them to
construct matrix $\mathbf{A}$. Solve for the optimal solution $\mathbf{x}^{*}$
of Problem 1;

for  _$i\in\mathcal{I}$_ do

Solve for the optimal solution $\mathbf{X}_{-i}^{*}$ of Problem 1;

Set the mobility payment for each traveler $i$:
$t_{i}=\sum_{\ell\in\mathcal{I}\setminus\\{i\\}}\sum_{j\in\mathcal{J}}u_{\ell}(x_{ij},\theta_{-i})\\\
-\sum_{\ell\in\mathcal{I}\setminus\\{i\\}}\sum_{j\in\mathcal{J}}u_{\ell}(x_{ij},\theta_{\ell})$

end for

for  _$j\in\mathcal{J}$_ do

Solve for the optimal solution $\mathbf{x}_{-j}^{*}$ of Problem 1;

Set the mobility payment for each provider $j$:
$t_{j}=\sum_{i\in\mathcal{I}}\sum_{\kappa\in\mathcal{J}\setminus\\{j\\}}u_{\kappa}(x_{ij},\delta_{-j})\\\
-\sum_{i\in\mathcal{I}}\sum_{\kappa\in\mathcal{J}\setminus\\{j\\}}u_{\kappa}(x_{ij},\delta_{\kappa})$

end for

Algorithm 1 Pricing Mechanism

######  Theorem 3 (Voluntary Participation).

No traveler $i\in\mathcal{I}$ and no provider $j\in\mathcal{J}$ can gain for
better individual utility by matching externally compared to the utility
gained by participating in the induced mobility game $\mathcal{M}$.

###### Proof.

It is sufficient to show that no agent in $\mathcal{I}\cup\mathcal{J}$ can
gain negative utility by participating in the induced game $\mathcal{M}$,
i.e., we must have $u_{i},u_{j}\geq 0$ for all
$i,j\in\mathcal{I}\cup\mathcal{J}$. First, note that the maximization of
$\sum_{i\in\mathcal{I}}\sum_{j\in\mathcal{J}}a_{ij}(x_{ij})$ is the highest
possible value we can achieve. Removing even one agent, does not increase this
value under any scenario. Thus, by definition, both payments $t_{i}$ and
$t_{j}$ are non-negative. At equilibrium, the utilities of any traveler and
provider are equivalent to the solutions $(\phi,\psi)$ of Problem 3 (as we
showed in Theorem 2). Since $(\phi,\psi)$ ensures non-negativity it follows
that $u_{i},u_{j}\geq 0$ for all $i,j\in\mathcal{I}\cup\mathcal{J}$. ∎

######  Theorem 4 (Truthfulness).

Misreporting does not benefit any traveler or provider.

###### Proof.

Let us assume that all agents in $\mathcal{I}\cup\mathcal{J}$ voluntary
participate in the mechanism. We consider two cases when traveler $i$
misreports their true type, i.e., $\hat{\theta}_{i}\geq\theta_{i}$ and
$\hat{\theta}_{i}\leq\theta_{i}$, where $\hat{\theta}_{i}$ is traveler $i$’s
report. If traveler $i$ reports $\hat{\theta}_{i}$ that is lower than their
true type, then traveler $i$ cannot improve their utility as $t_{i}$ is
necessarily non-negative and a lower misreporting
$\hat{\theta}_{i}\leq\theta_{i}$ can only lead to lower utilities. Suppose now
that traveler $i$ reports $\hat{\theta}_{i}$ that is higher than their true
type. The exact value of $\hat{\theta}_{i}$ can be chosen by the following
maximization problem

$\max_{\hat{\theta}_{i}}u_{i}(x_{ij},\hat{\theta}_{i})=\max_{\hat{\theta}_{i}}v_{i}(x_{ij},\hat{\theta}_{i})-t_{i},$
(21)

where $t_{i}$ is given by Algorithm 1. Note though that the maximization of
$v_{i}(x_{ij},\hat{\theta}_{i})-\sum_{\ell\in\mathcal{I}\setminus\\{i\\}}\sum_{j\in\mathcal{J}}u_{\ell}(x_{ij},\hat{\theta}_{-i})-\sum_{\ell\in\mathcal{I}\setminus\\{i\\}}\sum_{j\in\mathcal{J}}u_{\ell}(x_{ij},\hat{\theta}_{\ell})$
is equivalent to the maximization of
$\sum_{i\in\mathcal{I}}\sum_{j\in\mathcal{J}}v_{i}(x_{ij},\hat{\theta}_{i})$
with respect to the assignment $x_{ij}$. After all that is the goal of
traveler $i$, namely by misreporting their type to lead the mechanism to a
better assignment. Let $\hat{x}_{ij}^{*}$ and $x_{ij}^{*}$ denote the optimal
assignment with $\hat{\theta}_{i}$ and $\theta_{i}$, respectively. Hence, we
have

$\arg\max_{\hat{x}_{ij}^{*}}\sum_{i\in\mathcal{I}}\sum_{j\in\mathcal{J}}v_{i}(\hat{x}_{ij}^{*},\hat{\theta}_{i})\leq\arg\max_{x_{ij}^{*}}\sum_{i\in\mathcal{I}}\sum_{j\in\mathcal{J}}v_{i}(x_{ij}^{*},\theta_{i})$
(22)

for each traveler $i$. It follows immediately from (22) that traveler $i$ can
maximize their utility by minimizing the mobility payments as defined in
Algorithm 1. This is only possible when $\hat{\theta}_{i}=\theta_{i}$. Thus,
traveler $i$ cannot improve their utility by misreporting. Therefore, we
conclude that, with the proposed pricing mechanism (Algorithm 1), under no
circumstance can traveler $i$ improve their utility by misreporting about
their type $\theta_{i}$. In other words, any traveler $i$ has a strategy to
always truthfully report their type to the mechanism. We can follow the same
arguments to show this for the providers. Therefore, the proof is completed. ∎

######  Proposition 1.

If traveler $i$ is matched to provider $j$ while having misreported their type
to the mechanism, then traveler $i$ does not gain a better utility compared to
the utility gained under the true type.

###### Proof.

We show this only for the travelers as the arguments are similar for the
providers. By construction of the mobility payments in Algorithm 1 non-
negativity of the payments for each traveler is guaranteed, i.e., $t_{i}\geq
0$. By definition, the valuation of each traveler is non-negative under any
assignment. Thus, the utility defined in (2) is also non-negative. At
equilibrium, we have

$\max_{x_{ij}}\sum_{i\in\mathcal{I}}v_{i}(x_{ij},\theta_{i})\leq\sum_{i\in\mathcal{I}}v_{i}(x_{ij}^{*}).$
(23)

Thus, it follows that

$\displaystyle u_{i}(x_{ij},\theta_{i})$
$\displaystyle=v_{i}(x_{ij},\theta_{i})-t_{i}$
$\displaystyle=\sum_{i\in\mathcal{I}}v_{i}(x_{ij}^{*})-\sum_{\ell\in\mathcal{I}\setminus\\{i\\}}v_{\ell}(x_{ij},\theta_{\ell})\geq
0.$ (24)

Therefore, no traveler can hope for better utility by misreporting. ∎

######  Proposition 2 (Social Efficiency).

The proposed mechanism satisfies social efficiency as it ensures the
maximization of the aggregate social welfare of both travelers and providers.

###### Proof.

By construction of the mobility payments in Algorithm 1, it follows
immediately that the optimal solution maximizes the social welfare. ∎

##  IV SIMULATION RESULTS

In this section, we present a numerical example, its solution and discuss its
physical interpretation. Consider a mobility system of four providers
$\mathcal{J}=\\{\text{bike},\text{car},\text{bus},\text{train}\\}$ and twenty
travelers $\mathcal{I}=\\{1,2,\dots,20\\}$. Each provider has traveler
capacities, namely we have $\varepsilon_{bike}=1$, $\varepsilon_{car}=4$,
$\varepsilon_{bus}=5$, and $\varepsilon_{train}=10$. Moreover, we partition
the set of travelers $\mathcal{I}$ into four types, i.e., students, commuters,
tourists, consumers with sizes $|\mathcal{I}_{\text{students}}|=3$,
$|\mathcal{I}_{\text{commuters}}|=5$, $|\mathcal{I}_{\text{tourists}}|=4$,
$|\mathcal{I}_{\text{consumers}}|=8$. Each type of travelers can represent the
personal predisposition $\theta=(\theta_{i})_{i\in\mathcal{I}}$. Next, with a
slight abuse of notation, we generate a random utility assignment matrix

$\mathbf{A}=\\\ \begin{bmatrix}1&2&0.5\\\ 2.5&2&1.5\\\ 2.5&4&1.5\\\
2.5&5&6.5\end{bmatrix},$ (25)

where each row represents a type of travelers and each column represents a
provider. The entry $a_{ij}$ of $\mathbf{A}$ represents the overall utility of
assignment $x_{ij}$.

We solve Problem 2 and compute with an optimal solution that maximizes the
aggregate utilities of each traveler and each provider according to the
travelers’ preferences and maximizing the capacities of each provider. The
computational complexity of the proposed method is relatively low as long as
the number of travelers and providers remain small. This is reasonable to
expect as at any given moment there can only less than five different travel
options (so, the number of providers is always small). By ensuring we
partition the travelers’ requests according to origin, destination, and type,
we can make certain that the number of travelers does not make the
optimization problems untrackable.

Figure 1: Optimal assignment of travelers to providers with 20 travelers and 4
providers.

We can see from Fig. 1 that an efficient allocation of the providers’
resources and services to different types of travelers can be attained using a
game-theoretic framework. The assignment shown in Fig. 1 is stable, maximizes
the social welfare, and maximizes the capacity of the providers (thus,
maximizing their utilities too).

##  V CONCLUDING REMARKS

In this paper, we provided a theoretical study of a two-sided game for a
mobility system of travelers and providers focusing on how to discretely
assign travelers to providers when both have preferences. We formulated a
binary program and its equivalent linear program and showed that at least one
optimal solution exists and derived the conditions for such solution to be
stable. We then allowed informational asymmetry in the proposed mobility game
and provided a pricing mechanism to ensure we can elicit the private
information of all travelers and providers. We showed that our mechanism
guarantees economic efficiency in terms of maximizing the social welfare, and
ensures voluntary participation, thus making sure that all agents have a
unique dominant strategy.

Ongoing work includes relaxing our assumption of linearity in the utility
functions and also investigating our model under the behavioral model of
_prospect theory_ [28]. An interesting research direction should involve an
extension of our model with a sophisticated construction of the travelers’
utilities and preferences using data gathered from a behavioral survey. This
would helps us observe any correlations between the travelers’ tendencies or
attitudes and mobility preferences (which mode of transportation to use).

## References

  * [1] R. Colini-Baldeschi, R. Cominetti, P. Mertikopoulos, and M. Scarsini, “The asymptotic behavior of the price of anarchy,” _In International Conference on Web and Internet Economics_ , pp. 133–145, 2017.
  * [2] L. Zhao and A. A. Malikopoulos, “Enhanced mobility with connectivity and automation: A review of shared autonomous vehicle systems,” _IEEE Intelligent Transportation Systems Magazine_ , vol. 14, no. 1, pp. 87–102, 2022\.
  * [3] P. Barnes and E. Turkel. Autonomous vehicles in delaware: Analyzing the impact and readiness for the first state. 2021, October 1. [Online]. Available: http://udspace.udel.edu/handle/19716/21596
  * [4] R. Sarkar and J. Ward, “DOE SMART mobility: Systems and modeling for accelerated research in transportation,” in _Road Vehicle Automation_ , G. Meyer and S. Beiker, Eds. Springer, 2016, pp. 39–52.
  * [5] A. I. Mahbub, A. A. Malikopoulos, and L. Zhao, “Decentralized optimal coordination of connected and automated vehicles for multiple traffic scenarios,” _Automatica_ , vol. 117, no. 108958, 2020.
  * [6] G. Zardini, N. Lanzetti, M. Salazar, A. Censi, E. Frazzoli, and M. Pavone, “Towards a co-design framework for future mobility systems,” _Annual Meeting of the Transportation Research Board_ , 2020.
  * [7] A. A. Malikopoulos, L. E. Beaver, and I. V. Chremos, “Optimal time trajectory and coordination for connected and automated vehicles,” _Automatica_ , vol. 125, no. 109469, 2021.
  * [8] M. Harb, Y. Xiao, G. Circella, P. L. Mokhtarian, and J. L. Walker, “Projecting travelers into a world of self-driving vehicles: estimating travel behavior implications via a naturalistic experiment,” _Transportation_ , vol. 45(6), pp. 1671–1685, 2018.
  * [9] D. Bissell, T. Birtchnell, A. Elliott, and E. L. Hsu, “Autonomous automobilities: The social impacts of driverless vehicles,” _Current Sociology_ , 2018.
  * [10] P. A. Singleton, “Discussing the “positive utilities” of autonomous vehicles: will travellers really use their time productively?” _Transport Reviews_ , vol. 39.1, pp. 1–16, 2018.
  * [11] A. C. Pigou, _The Economics of Welfare_. Macmillan London, 2013.
  * [12] K. Wada and T. Akamatsu, “Auction mechanisms for implementing tradable network permit markets,” _Journal of Japan Society of Civil Engineers_ , vol. 67.3, pp. 376–389, 2011.
  * [13] M. Salazar, N. Lanzetti, F. Rossi, M. Schiffer, and M. Pavone, “Intermodal autonomous mobility-on-demand,” _IEEE Transactions on Intelligent Transportation Systems_ , vol. 21(9), pp. 3946–3960, 2019.
  * [14] I. V. Chremos and A. A. Malikopoulos, “Social resource allocation in a mobility system with connected and automated vehicles: A mechanism design problem,” in _Proceedings of the 59th IEEE Conference on Decision and Control, 2020_. IEEE, 2020, pp. 2642–2647.
  * [15] ——, “Design and stability analysis of a shared mobility market,” in _2021 European Control Conference (ECC)_ , 2021, pp. 374–379.
  * [16] I. V. Chremos, L. E. Beaver, and A. A. Malikopoulos, “A game-theoretic analysis of the social impact of connected and automated vehicles,” in _2020 23rd International Conference on Intelligent Transportation Systems (ITSC)_. IEEE, 2020, pp. 2214–2219.
  * [17] I. V. Chremos and A. A. Malikopoulos, “Socioeconomic impact of emerging mobility markets and implementation strategies,” in _AI-enabled Technologies for Autonomous and Connected Vehicles_ , I. Kolmanovsky, Y. Murphey, and P. Watta, Eds. Springer, 2022 (in press).
  * [18] A. Abdulkadiroǧlu, P. A. Pathak, A. E. Roth, and T. Sönmez, “The boston public school match,” _American Economic Review_ , vol. 95(2), pp. 368–371, 2005.
  * [19] K. Treleaven, M. Pavone, and E. Frazzoli, “An asymptotically optimal algorithm for pickup and delivery problems,” in _2011 50th IEEE Conference on Decision and Control and European Control Conference_ , 2011, pp. 584–590.
  * [20] P. You, J. Z. Pang, M. Chen, S. H. Low, and Y. Sun, “Battery swapping assignment for electric vehicles: A bipartite matching approach,” in _2017 IEEE 56th Annual Conference on Decision and Control (CDC)_ , 2017, pp. 1421–1426.
  * [21] D. Gale and L. S. Shapley, “College admissions and the stability of marriage,” _The American Mathematical Monthly_ , vol. 69(1), pp. 9–15, 1962\.
  * [22] L. S. Shapley and M. Shubik, “The assignment game I: The core,” _International Journal of Game Theory_ , vol. 1(1), pp. 111–130, 1971.
  * [23] G. Demange, D. Gale, and M. Sotomayor, “Multi-item auctions,” _Journal of Political Economy_ , vol. 94(4), pp. 863–872, 1986.
  * [24] E. Anshelevich, S. Das, and Y. Naamad, “Anarchy, stability, and utopia: Creating better matchings,” in _Autonomous Agents and Multi-Agent Systems_ , vol. 26(1), 2013, pp. 120–140.
  * [25] W. Krichene, B. Drighès, and A. M. Bayen, “Online learning of nash equilibria in congestion games,” _SIAM Journal on Control and Optimization_ , vol. 53(2), pp. 1056–1081, 2015.
  * [26] P. N. Brown and J. R. Marden, “Optimal mechanisms for robust coordination in congestion games,” _IEEE Transactions on Automatic Control_ , vol. 63(8), pp. 2437–2448, 2017.
  * [27] A. Schrijver, _Theory of Linear and Integer Programming_. John Wiley & Sons, 1996.
  * [28] A. Tversky and D. Kahneman, “Advances in prospect theory: Cumulative representation of uncertainty,” _Journal of Risk and Uncertainty_ , vol. 5(4), pp. 297–323, 1992.


11institutetext:  Department of Computer Science, Stony Brook University, USA
11email: {shilei.tian<EMAIL_ADDRESS>22institutetext:
Advanced Micro Devices, UK
22email<EMAIL_ADDRESS>33institutetext: Mathematics and Computer Science,
Argonne National Laboratory, USA
33email<EMAIL_ADDRESS>
# Experience Report: Writing A Portable GPU Runtime with OpenMP 5.1

Shilei Tian 11 0000-0001-6468-6839 Jon Chesterfield 22 0000-0002-8546-2014
Johannes Doerfert 33 0000-0001-7870-8963 Barbara Chapman 11
0000-0001-8449-8579

###### Abstract

GPU runtimes are historically implemented in CUDA or other vendor specific
languages dedicated to GPU programming. In this work we show that OpenMP 5.1,
with minor compiler extensions, is capable of replacing existing solutions
without a performance penalty. The result is a performant and portable GPU
runtime that can be compiled with LLVM/Clang to Nvidia and AMD GPUs without
the need for CUDA or HIP during its development and compilation.

While we tried to be OpenMP compliant, we identified the need for compiler
extensions to achieve the CUDA performance with our OpenMP runtime. We hope
that future versions of OpenMP adopt our extensions to make device programming
in OpenMP also portable across compilers, not only across execution platforms.

The library we ported to OpenMP is the OpenMP device runtime that provides
OpenMP functionality on the GPU. This work opens the door for shipping OpenMP
offloading with a Linux distribution’s LLVM package as the package manager
would not need a vendor SDK to build the compiler and runtimes. Furthermore,
our OpenMP device runtime can support a new GPU target through the use of a
few compiler intrinsics rather than requiring a reimplementation of the entire
runtime.

###### Keywords:

OpenMP LLVM Portability Target offloading Runtimes Accelerator

##  1 Introduction

In this paper, we describe how we ported the LLVM OpenMP device runtime
library to OpenMP 5.1 using only minor extensions not available in the
standard. The OpenMP device runtime provides the OpenMP functionalities to the
user and implementation code on the device, which in this context means on the
GPU. As an example, it provides the OpenMP API routines as well as routines
utilized by the compiler e.g., for worksharing loops.

Our work replaced the original LLVM OpenMP device runtime implemented in CUDA
to allow for code reusibility between different targets, e.g. AMD and Nvidia.
It further lowers the bar to entry for future targets that only need to
provide a few target specific intrinsics and minimal glue code.

The OpenMP device runtime library can now be shipped with pre-build LLVM
packages as they only need LLVM/Clang to build it; neither a target device nor
vendor SDKs are required, which lowers the barrier to entry for OpenMP
offloading. This work is a proof of concept for writing device runtime
libraries in OpenMP, with identical functionality and performance to that
available from CUDA or HIP compiled with the same LLVM version.

The remainder of the paper is organized as follows. We discuss background and
motivation in Section 2. Section 3 presents our approach, which is followed by
an evaluation in Section 4. Finally, we conclude the paper in Section 5.

##  2 Background

When compiling from one language to another, there are usually constructs that
are straightforward in the former and complicated or verbose in the latter.
For example, a single OpenMP construct #pragma omp parallel for is lowered
into a non-trivial amount of newly introduced code in the application,
including calls into a runtime that provides certain functionality, like
dividing loop iterations. In this work, the input is OpenMP target offloading
code, that is the OpenMP target directive and the associated code, and the
output is ultimately Nvidia’s PTX or AMD’s GCN assembler.

###  2.1 Device Runtime Library

The LLVM OpenMP device runtime library contains the various functions the
compiler needs to implement OpenMP semantics when the target is an Nvidia or
AMD GPU. It is basically libomp for the GPU. The original implementation in
LLVM was in CUDA [8], compiled with Nvidia’s NVCC to PTX assembler which was
linked with the application code to yield a complete program. We later adapted
that source to compile alternatively as HIP, which is close enough to CUDA
syntax for the differences to be worked around with macros. Prior to this work
the device runtime was hence comprised of sources in a common and target
dependent part. In order to let the target dependent compiler recognize the
code, target dependent keywords (such as __device__ and __shared__ in CUDA)
are replaced with macros (DEVICE and SHARED), and the header where these
macros are defined will be included accordingly depending on the target. The
basic idea is visualized in in 1.

⬇

// Common part

DEVICE void *__kmpc_alloc_shared(uint64_t bytes);

SHARED int shared_var;

// CUDA header

#define DEVICE __device__

#define SHARED __shared__

// AMDGCN header

#define DEVICE __attribute__((device))

#define SHARED __attribute__((shared))

Listing 1: Macros in current device runtime.

This strategy works. For Nvidia offloading the source is compiled as CUDA, for
AMDGPU offloading it is compiled as HIP. Both produce LLVM bitcode but with
different final targets, Nvidia’s PTX and AMD’s GCN respectively. However, if
a programming model does not adequately resemble CUDA, such as OpenCL or
Intel’s DPC++ [3], the approach will become less straight forward.

What’s more, this setup requires vendor SDKs (such as CUDA Toolkit or ROCm
Developer Tools) to compile the device runtime, which creates a barrier for
the package managers of Linux distributions. In practice that means the LLVM
OpenMP installed from Linux distributions does not support offloading out of
the box because the package would require a dependence on the CUDA or ROCm
package, among other things.

###  2.2 Compilation Flow of OpenMP Target Offloading in LLVM/Clang

The compilation of an OpenMP program with target offloading directives
contains the following two passes (as shown in Fig. 1):

Host Code Compilation.

This pass includes the regular compilation of code for the host and OpenMP
offloading code recognition as preparation for the second pass. Offloading
regions are replaced by calls to the corresponding host runtime library
functions (e.g. __tgt_target for the directive target in LLVM OpenMP) with
suitable arguments, such as the kernel function identifier, base pointers of
each captured variables and the number of kernel function arguments. In
addition, a fallback host version of the kernel function will be emitted in
case target offloading fails at runtime.

Device Code Compilation.

This pass utilizes the recognized OpenMP target offload regions, as well as
related functions and captured variables, and then emits target dependent
device code. This includes one entry kernel function per target region, global
variables (potentially in different address spaces), and device functions, as
well as some target dependent metadata. As part of this compilation the OpenMP
device runtime library is linked into the user code as an LLVM bitcode library
(dev.rtl.bc in the Fig. 1).

foo.cHostCompilerDeviceCompilerhost.odevice.imgdev.rtl.bcOffloadBundlerfoo.o
Figure 1: Compilation flow of an OpenMP program with target offloading.

In addition to the target construct (as well as its combined variants), OpenMP
provides the declare target directive which specifies that variables and
functions are mapped onto a target device, and should hence be usable in
device code. The declare variant directive can be used to specify a context,
e.g., the compilation for a specific target, in which a specialized function
variant should replace the base version.

###  2.3 Motivation

While the OpenMP device runtime library can be implemented in any language it
should be linked into the application in LLVM bitcode format for performance
reasons. This setup, shown in Fig. 1, allows to optimize the runtime together
with the application, effectively specializing a generic runtime as needed.

Given that the base language is irrelevant as long as we can compile to LLVM
bitcode, OpenMP comes to mind as a portable and performant way to write code
for different accelerators. As almost the entire device library can be
interpreted as C++ code, rather than a CUDA or HIP code base, the compilation
as OpenMP is feasible, in particular because LLVM/Clang is a working C++ and
OpenMP compiler already.

Since OpenMP 5.1 all conceptually necessary building blocks are present in the
language specification:

  * •

The declare target directive can be used to compile for a device, hence to
generate LLVM bitcode that is targeting Nvidia’s PTX or AMD’s GCN. As we do
not need a host version at all, we can even use the LLVM/Clang flag -fopenmp-
is-device to invoke only the device compilation pass described in Section 2.2.

  * •

The declare variant directive can be used if a target requires a function
implementation or global variable definition different from the default.

  * •

The allocate directive provides access to the different kinds of memory on the
GPU.

For an additional target architecture, the work done in the compiler backend
to emit code for that architecture will allow to retarget an OpenMP
implemented device runtime almost for free. The incremental development cost
is reduced from (re)implementing the device runtime in a language that can be
compiled to the new architecture to providing a few declare variant
specialisations.

Finally, if the port uses compiler intrinsics instead of CUDA or HIP functions
for the small target dependent part, it can be compiled without a vendor
specific SDK present. This unblocks shipping offloading as part of Linux
distributions.

##  3 Implementation

In this section, we describe the new LLVM OpenMP device runtime implemented
with OpenMP 5.1. First, we talk about the common part, and then discuss how
target dependent parts are implemented and why extensions were necessary. Only
AMD and Nvidia platforms are discussed as other GPU architectures cannot be
targeted by the community LLVM version at this time.

###  3.1 Common Part

##### Device Code

Using the declare target directive around all source files causes all
functions and data to be emitted for the target device. Macros to indicate
that functions or globals are for the device, as shown in 1, are not needed.

##### Global Shared Variables

The implementation of the device runtime maps an OpenMP team to a thread
block111We are using CUDA terminology here. For AMD platforms it is wavefront.
on the target device. Therefore, a shared variable visible to all threads in
the same thread block is equivalent to a variable that can be accessed within
the same OpenMP team. The allocate directive specifies how to allocate
variables in different memory spaces. Uses with an
allocator(omp_cgroup_mem_alloc)222The implementation currently uses
allocator(omp_pteam_mem_alloc) which is equivalent given the current mapping
of parallelism. we can place global variables in local shared memory, the
equivalent of the CUDA __shared__ shown in 1.

In contrast to shared CUDA or HIP variables, C++ specifies that global
variables are default initialized. While we can technically do this for global
shared variables defined with OpenMP, it is not supported by LLVM/Clang at
this time. Furthermore, the performance is likely to suffer as the device
runtime is designed to initialize these variables explicitly on demand. To
this end, we extended LLVM/Clang with a variable attribute for this work:
loader_uninitialized [1]. The effect is that annotated variables will not have
a default initialized value but instead be uninitialized like the CUDA or HIP
shared variables are as well.

2 shows device code and global shared variable declaration as it is used in
our OpenMP device runtime.

⬇

#pragma omp begin declare target

// Function declaration

extern __kmpc_impl_threadfence();

// Function definition

void __kmpc_flush(kmp_Ident *loc) {

__kmpc_impl_threadfence();

}

// Global variable

int global_var;

// Shared variable

int shared_var;

#pragma omp allocate(shared_var). \

allocator(omp_pteam_mem_alloc)

// Shared variable declaration

extern int other_shared_var;

#pragma omp allocate(other_shared_var) \

allocator(omp_pteam_mem_alloc)

#pragma omp end declare target

Listing 2: An example of new device runtime code.

##### Atomic Operations

The device runtime uses five atomic operations, add, inc, max, exchange, and
cas, implemented in target dependent parts with LLVM/Clang builtin functions.

OpenMP 5.1 [4] introduces the compare clause, which supports conditional
update statements. When combined with the capture clause, all of these atomic
operations except inc can be implemented via OpenMP, as shown in 3. We
implemented the support of the compare clause and its combination with the
capture clause for LLVM/Clang but the it has not been merged into the
community version yet. With the updated requirements for flush333OpenMP 5.1
removes the requirement for a flush operation at the entry and exit of an
atomic operation if write, update, or capture is specified and the memory
ordering is seq_cst, which we also implemented for this work, our OpenMP
versions of atomic operations can generate LLVM-IR that is identical to the
original target dependent implementation via compiler intrinsics.

⬇

uint32_t atomic_add(uint32_t *X, uint32_t E) {

uint32_t V;

#pragma omp atomic capture seq_cst

{ V = *X; *X += E; }

return V;

}

uint32_t atomic_max(uint32_t *X, uint32_t E) {

uint32_t V;

#pragma omp atomic compare capture seq_cst

{ V = *X; if (*X < E) { *X = E; } }

return V;

}

uint32_t atomic_exchange(uint32_t *X, uint32_t E) {

uint32_t V;

#pragma omp atomic capture seq_cst

{ V = *X; *X = E; }

return V;

}

uint32_t atomic_cas(uint32_t *X, uint32_t E, uint32_t D) {

uint32_t V;

#pragma omp atomic compare capture seq_cst

{ V = *X; if (*X == E) { *X = D; } }

return V;

}

Listing 3: Atomic operations implemented in OpenMP 5.1.

The missing atomic operation is inc. According to the CUDA specification [2],
inc implements:

⬇

{ v = x; x = x >= e ? 0 : x + 1; }

and returns v. This atomic operation can not be represented in a form that
OpenMP 5.1 supports because OpenMP 5.1 requires that the order operation be
either < or >, and the alternative statement of the conditional expression
statement must be x itself. Therefore, we still keep it in the target
dependent part implemented with LLVM intrinsics as shown in 4.

###  3.2 Target Specific Part

Target dependent global functions and variables are currently declared in a
header and implemented in target dependent source files which are only
compiled for the specific target, either as CUDA or HIP. A drawback of this
method is that the creation of a device runtime for a new target might require
us to remove a function from the common part and insert it into the target
specific part if the existing (common) implementation is not suited for the
new device.

Since OpenMP 5.0, the declare variant directive declares a specialized variant
of a base function and specifies the context in which that specialized variant
is used. It supports various context selector with the match clause, one of
which is device selector. For example, with match(device={arch(arch_name)}),
the code wrapped in a begin/end declare variant region will be only generated
if the target architecture matches the arch_name.

4 shows how the atomic inc function is implemented with target dependent
compiler intrinsics selected via the begin/end declare variant directive for
both Nvidia and AMD GPU targets.

Note that we use the match_any extension for Nvidia platforms as we support
two distinct architectures, nvptx and nvptx64, but we do not want to
distinguish between them in the device runtime. While this can be handled by
duplicating the code, our new context selector changes the semantic of the
matching to produce a match if any architecture in arch(nvptx, nvptx64) is
targeted. By default a match would require all architectures to be targeted.
In addition to match_any we extended LLVM/Clang with other useful context
selectors, e.g, match_none and allow_templates444See:
https://clang.llvm.org/docs/AttributeReference.html#pragma-omp-declare-
variant.

⬇

#pragma omp declare target

// Fallback version, which raises a compilation error

uint32_t atomic_inc(uint32_t *X, uint32_t E) {

error("target␣dependent␣implementation␣missing");

}

// AMDGCN implementation

#pragma omp begin declare variant \

match(device={arch(amdgcn)})

uint32_t atomic_inc(uint32_t *X, uint32_t E) {

return __builtin_amdgcn_atomic_inc32(X, E,

__ATOMIC_SEQ_CST, "");

}

#pragma omp end declare variant

// NVPTX implementation

#pragma omp begin declare variant \

match(device={arch(nvptx,nvptx64)}, \

implementation={extension(match_any)})

uint32_t atomic_inc(uint32_t *X, uint32_t E) {

return __nvvm_atom_inc_gen_ui(X, E);

}

#pragma omp end declare variant

#pragma omp end declare target

Listing 4: Atomic inc implementation with the match_any clause.

Other target dependent functions are required to handle synchronization,
thread hierarchy, etc. These are implemented via compiler intrinsics, function
calls to the corresponding native runtime library, or inline assembly.

##  4 Evaluation

In this section, we evaluated our proposed method in three ways: code
comparison, functional testing, and performance evaluation.

###  4.1 Code Comparison

The previous implementation compiled CUDA to LLVM-IR, and HIP to LLVM-IR,
while our proposed method compiles OpenMP to LLVM-IR for both platforms. The
accuracy of the port to OpenMP was assessed by comparing the text form of the
library before and after changing over to OpenMP. If the text forms were
identical, we would be certain the language change made no difference. This
was not quite the case. The differences were in semantically unimportant
metadata, symbol name mangling for variant functions, and the order of
inlining as preferred by the language front end which had minor reordering
effects on PTX and GCN generation.

###  4.2 Functional Testing

There are a number of OpenMP test suites and applications in use for checking
the behaviour of the compiler, including SOLLVE V&V [7], and Ovo [5]. All ran
identically with the new OpenMP runtime as they had using the previous device
runtime.

###  4.3 Performance Evaluation

####  4.3.1 Systems Configuration

We evaluate the performance of our method experimentally on the Summit
supercomputer. Each Summit node contains two IBM POWER9 processors and six
Nvidia Volta V100 GPUs. CUDA 10.1.243 was used, which is the version loaded by
default.

####  4.3.2 Benchmarks

The SPEC ACCEL benchmark suite V1.3 was used to evaluate the new device
runtime. Because support for Fortran is still in progress, we chose those
benchmarks written in C. There are 15 OpenMP enabled benchmarks in SPEC ACCEL.
Seven of them are in C, namely 503.postencil, 504.polbm, 514.pomriq, 552.pep,
554.pcg, 557.pcsp, and 570.pbt. 557.pcsp can not be compiled, therefore we
only ran the other six benchmarks. We also chose a C++ proxy application,
miniQMC [6].

O2 was used when compiling the benchmarks and application. Each test case was
executed five times, and the execution time was averaged. miniQMC was measured
through the miniqmc_sync_move benchmark executed as follows: miniqmc_sync_move
-g "2 2 1".

####  4.3.3 Results

Fig. 2 compares the execution time when the original device runtime is used
with the execution time obtained using our proposed new device runtime. We can
see that the execution times are almost identical, and for those cases where
they are not same, the variance is less than $1\%$ and assumed to be noise.

503.postencil504.polbm514.pomriq552.pep554.pcg570.pbt$10$$20$$30$$40$$50$$60$$70$$13.8$$32.2$$28.2$$68.6$$98.1$$513$$13.6$$32.2$$28.2$$68.6$$98.1$$514$Execution
Time (s) Figure 2: Comparison between execution time of original device
runtime ( ) and that of our proposed new device runtime ( ) on Nvidia
platform.

The proxy application benchmark miniqmc_sync_move contains two target regions,
evaluate_vgh and evaluateDetRatios. They are executed multiple times. Table 1
shows the profiling results (execution time) of each target region from
Nvidia’s profiler nvprof. There is no performance difference between the two
versions.

Target Region | Version | Time (ms) | # Calls | Avg ($\mu$s) | Min ($\mu$s) | Max ($\mu$s)
---|---|---|---|---|---|---
evaluate_vgh | Original | 1374.72 | 64512 | 21.309 | 19.744 | 32.384
New | 1376.59 | 21.338 | 19.776 | 33.760
evaluateDetRatios | Orignal | 573.46 | 18202 | 31.505 | 25.247 | 44.480
New | 573.93 | 31.531 | 24.544 | 47.103
Table 1: Comparison of execution time of the two target regions in
miniqmc_sync_move on Nvidia platform.

All the results above demonstrate that our proposed portable OpenMP device
runtime can provide the same performance as the current CUDA-like version on
the Nvidia platform. Based on the code comparison, functional testing and some
AMD internal performance testing results, the portable runtime is believed to
show no performance change from its HIP predecessor either.

##  5 Conclusions and Future Work

OpenMP works well as a language to implementing GPU-only code libraries. The
direct support for memory allocators and the precise dispatch through declare
variant are clear advantages over C++. While minimal compiler modifications
were required to match the CUDA and HIP semantics to the fullest, we expect
those to be incorporated into the OpenMP standard over time.

Using OpenMP is especially suitable as the vehicle for implementing an OpenMP
runtime library since the main prerequisite is an OpenMP compiler which needs
to be implemented all targets in any case. Since the library ships with the
LLVM repository, it can be built by any distribution which has built Clang.
Vendor SDKs or compilers are no longer required.

Since the host and device runtime libraries can build as part of LLVM, we will
coordinate with Linux distribution developers to ensure that people who
install the distribution LLVM package onto a system that has a target device
and driver available will be able to get this working “out of the box”.

## Acknowledgement

This research was supported by the Exascale Computing Project (17-SC-20-SC), a
collaborative effort of two U.S. Department of Energy organizations (Office of
Science and the National Nuclear Security Administration) responsible for the
planning and preparation of a capable exascale ecosystem, including software,
applications, hardware, advanced system engineering, and early testbed
platforms, in support of the nation’s exascale computing imperative.

## References

  * [1] Attributes in clang, https://clang.llvm.org/docs/AttributeReference.html#loader-uninitialized
  * [2] Cuda toolkit documentation v11.3.0, https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#atomicinc
  * [3] Intel® oneapi dpc++/c++ compiler, https://software.intel.com/content/www/us/en/develop/tools/oneapi/components/dpc-compiler.html
  * [4] Openmp application programming interface version 5.1, https://www.openmp.org/spec-html/5.1/openmp.html
  * [5] Ovo: Openmp vs offload, https://github.com/TApplencourt/OvO
  * [6] Qmcpack/miniqmc, https://github.com/QMCPACK/miniqmc
  * [7] Sollve/sollve_vv, https://github.com/SOLLVE/sollve_vv
  * [8] Jacob, A.C., Eichenberger, A.E., Sung, H., Antao, S.F., Bercea, G.T., Bertolli, C., Bataev, A., Jin, T., Chen, T., Sura, Z., Rokos, G., O’Brien, K.: Efficient fork-join on gpus through warp specialization. In: 2017 IEEE 24th International Conference on High Performance Computing (HiPC). pp. 358–367 (2017). https://doi.org/10.1109/HiPC.2017.00048