#  AVEC: Accelerator Virtualization in Cloud-Edge Computing for Deep Learning
Libraries

Jason Kennedy, Blesson Varghese and Carlos Reaño  Queen’s University Belfast,
UK
{jkennedy49, b.varghese<EMAIL_ADDRESS>
###### Abstract

Edge computing offers the distinct advantage of harnessing compute
capabilities on resources located at the edge of the network to run workloads
of relatively weak user devices. This is achieved by offloading
computationally intensive workloads, such as deep learning from user devices
to the edge. Using the edge reduces the overall communication latency of
applications as workloads can be processed closer to where data is generated
on user devices rather than sending them to geographically distant clouds.
Specialised hardware accelerators, such as Graphics Processing Units (GPUs)
available in the cloud-edge network can enhance the performance of
computationally intensive workloads that are offloaded from devices on to the
edge. The underlying approach required to facilitate this is virtualization of
GPUs. This paper therefore sets out to investigate the potential of GPU
accelerator virtualization to improve the performance of deep learning
workloads in a cloud-edge environment. The AVEC accelerator virtualization
framework is proposed that incurs minimum overheads and requires no source-
code modification of the workload. AVEC intercepts local calls to a GPU on a
device and forwards them to an edge resource seamlessly. The feasibility of
AVEC is demonstrated on a real-world application, namely OpenPose using the
Caffe deep learning library. It is observed that on a lab-based experimental
test-bed AVEC delivers up to 7.48x speedup despite communication overheads
incurred due to data transfers.

###### Index Terms:

Edge Computing, Accelerators, Virtualization, Deep Learning

##  I Introduction

Edge computing offers multiple layers of resources at the edge of the network
between an end-user device and the cloud [1, 2]. The premise of edge computing
is to bring workloads closer to where data is generated for minimising
communication latencies between devices and geographically distant clouds and
for reducing the ingress bandwidth demand to the cloud [3, 4]. Edge resources
may be employed to offload workloads from the cloud to the edge or from
devices to the edge. Edge resources may vary in form factor depending on where
they are located. For example, a home router may be augmented with computing
resources to be an edge node or alternatively a dedicated micro-cloud with
more computing resources may be employed. Figure 1 provides an exemplar of an
edge computing architecture [5].

Figure 1: An example cloud-edge architecture

Workloads that execute on the edge, such as deep learning [6, 7] can benefit
from hardware accelerators that provide massive parallelism on relatively
small form factor processors. One example is the graphics processing unit
(GPU) that this paper focuses on. Alternate accelerators include FPGAs (field-
programmable gate arrays) or TPUs (tensor processing units). Co-locating
accelerators on the edge may reduce the execution time of workloads thereby
making it a more compelling proposition for the edge. Workloads will need to
typically rely on frameworks, such as CUDA [8] or OpenCL[9] to exploit the
parallelism offered by GPUs. The GPU executes functions that are referred to
as kernels, which are routines stored on the accelerator and are separate from
those executed on the CPU. The CPU will send an input to the accelerator for
the kernel and then receive the kernel output.

All end user devices may not have immediate access to an edge node with a GPU
since all edge resources will not be GPU powered. Therefore it is anticipated
that workloads will need access to GPUs that are located on other resources in
the cloud-edge network. For example, an accelerator that is hosted in a micro-
cloud rather than on an edge node. This is facilitated by the virtualization
technology; accessing remote resources and using them as if they were local
resources. The cloud-edge architecture mentioned above envisions that
resources in a cloud-edge network will be able to access virtualized
accelerators hosted on nodes in the network. Thus, in practice, resources with
no physical accelerators can potentially access accelerators on other nodes.

GPU virtualization in the cloud-edge network can be achieved by intercepting
local calls to a GPU and forwarding them to a resource hosting an accelerator,
executing the GPU kernel on the virtualized accelerator, and returning the
output of execution. This offers the possibility for a computationally weaker
device to send its data through the network so that a resource hosting a GPU
can execute the workload instead of the device (a high-level view is provided
in Figure 2).

Figure 2: Example of accelerator virtualization in cloud-edge

Although there is significant research on virtualizing GPUs for high-
performance computing systems, such as rCUDA [10] or gVirtuS [11], and on GPUs
for low-power embedded computers suited for the edge [12, 13], there is
currently no research that focuses on virtualization in a cloud-edge network.
Therefore, the research reported in this paper investigates the potential of
GPU accelerator virtualization to improve the performance of deep learning
workloads in a cloud-edge environment. More specifically, the following
research questions are addressed:

Q1: How can accelerator virtualization be offered at the edge given its
limited compute resources? To address this question, this paper will present a
novel prototype AVEC (accelerator virtualization in cloud-edge computing), a
framework that supports the use of virtual GPU accelerators in the cloud-edge
continuum. AVEC is portable and can be employed in heterogeneous cloud-edge
environments. The key advantage of the proposed framework is that it does not
require source code modification.

Q2: What overheads are incurred in virtualizing accelerators at the edge? The
key motivation in addressing this question is to determine whether the
computational benefits of having access to more powerful GPUs are offset by
additional overheads incurred due to communication. It is observed that the
proposed AVEC framework can deliver up to 7.48x speedup despite communications
overheads incurred due to data transfer to a remote GPU.

Q3: Can GPU accelerator virtualization at the edge improve the execution
performance of real-world applications? In this paper, we select a deep
learning library, namely Caffe [14] to demonstrate this. We show that it is
feasible to offload computationally intensive components of workloads that
require GPUs from weak devices and edge nodes to nodes hosting GPUs via
virtualization. The experimental evaluation is carried out in the context of a
deep learning application which uses the aforementioned Caffe library, namely
OpenPose [15].

The main research contributions of this paper are as follows:

  * •

The development of AVEC, a first prototype of a low overhead and edge
performance enhancing framework that supports the use of virtual GPU
accelerators in the cloud-edge continuum.

  * •

The demonstration of the benefits of accelerator virtualization at the edge by
transparent execution of the popular Caffe deep learning library kernels using
a real-world use case on remote GPUs.

This research confirms the feasibility of offloading workloads to virtualized
GPU accelerators in cloud-edge environments.

The remainder of this paper is organised as follows. Section II presents
existing techniques implementing virtualization and management of resources.
Section III highlights the motivation and the necessity for implementing
accelerator virtualization within a cloud-edge network. Section IV proposes
AVEC; the accelerator virtualization framework and provides its
implementation. Section V presents the results obtained from experimental
studies. Finally, Section VI concludes this paper by considering future work.

##  II Related Work

Existing literature presents accelerator virtualization solutions in the
context of large-scale systems, such as high-performance computing (HPC)
clusters and clouds, and low power embedded devices.

###  II-A Accelerator Virtualization in Large-scale Systems

There are multiple solutions available for general-purpose computing on
graphics processing units (GPGPU) virtualization in HPC clusters. These
solutions are based on middleware libraries or remote procedure calls (RPC).

Remote CUDA (rCUDA) [10] is one middleware-based solution that virtualizes
remote GPUs. The middleware intercepts CUDA calls on a client node and
forwards them to a server node hosting a physical GPU. This allows for the GPU
accelerator to be logically decoupled from the physical node, thereby allowing
for other clients to access and share the same physical GPU on the server. A
similar framework that is developed for OpenCL is VOCL [16].

vCUDA [17] is another CUDA based accelerator virtualization solution for HPC
clusters. This solution is RPC-based rather than middleware to achieve
accelerator virtualization. GViM [18] is an early implementation of GPU
virtualization in which multiple virtual machines are hosted on a single
physical node and access the same physical GPU. GViM achieves virtualization
by using API interception, so that the CUDA function calls from an application
running on a virtual machine, can be intercepted and sent to the host machine
in a privileged domain for execution.

Alternate solutions for GPU virtualization include DS-CUDA [19] and Grid-CUDA
[20]. DS-CUDA is a GPU virtualization solution designed for use in the cloud.
This solution incorporates a redundancy mechanism by mapping two physical
cloud accelerators on to a single virtual accelerator. The outputs of both
accelerators are compared; if the two results do not match, DS-CUDA
automatically re-runs the CUDA API calls until the same result is seen. This
mechanism improves accuracy, but the overheads remain unknown.

Grid-CUDA makes use of RPC to redirect workloads within a grid of nodes to
enable parallel execution. It is noted that the use of RPC incurs large
overhead costs, and the use of GRID-CUDA could lower the overall performance.

GVirtuS [11] is another accelerator virtualization solution that is based on
virtual machines operating over TCP/IP model, thereby offering remote
accelerator virtualization. GVirtuS is independent of the hypervisor. However,
the choice of hypervisor affects the performance. The hypervisor allows
multiple virtual machines or remote devices to interact with the same physical
hardware accelerator. The GVirtuS framework is designed so that the CUDA
library interacts with a front-end GPU virtualizer. The back end of GVirtuS
deals with the hardware, by unpacking the CUDA library call and then assigning
memory for it to be executed. The interception is made using a CUDA wrapper
library.

qCUDA [21] is proposed to improve the performance in areas such as bandwidth,
for local GPU virtualization. They achieve virtualization by incorporating API
interception and is designed for QEMU-KVM hypervisor. Their testing focused on
bandwidth performance and they report above 95% bandwidth efficiency when
compared to native execution. The main contribution qCUDA provides is
efficient memory transfers. This is made possible by eliminating the extra
memory copies between guest and host, by shifting a contiguous guest physical
address in memory as a host virtual address in the QEMU process via mapping.

Figure 3: Example output of the OpenPose application

###  II-B Accelerator Virtualization on Embedded Devices

Recently, GVirtuS incorporated accelerator virtualization for ARM based single
board computers (SBC) [22]. This is achieved by intercepting the front end API
stub that connects low powered SBCs to GPU accelerators located on x86
servers. The framework offloads workloads from the ARM based SBCs to remote
accelerators. Results highlight that the framework is best suited for longer
running workloads. This is because the latency factor due to communicating
between nodes decreases over time (performance increases as the communication
time compared to kernel execution time decreases).

qCUDA-ARM [23] is an extension of the qCUDA framework and incorporates SBCs.
This is a first attempt to virtualize GPUs to work on the ARM architecture.
While it uses the same conceptual architecture of qCUDA, the memory management
process is modified. A key CUDA function used by qCUDA is cudaHostRegister(),
which is not supported in ARM devices. As such, the team has implemented
memory copying in a different way, so that an additional copy of the pinned
memory section is created, mapping this new region to the memory in the guest
CUDA application. Although this utilizes twice the amount of memory, it is
noted that it accelerates the execution of other memory related CUDA calls
such as cudaMemcpy(). The authors report in their experimental findings that
they can achieve up to 90% of the native CUDA speed for pinned memory.
Furthermore, they achieve almost native performance on computation bound
applications such as matrix multiplication.

RAPID [24] project is a European Commission funded project that investigates
virtualization of lower power SBCs for offloading workloads in a heterogeneous
environment. A peer-to-peer sharing mechanism in which devices from
smartphones to cloud data centers are connected is envisioned to take
advantage of large accelerators in the network to augment the performance of
SBCs.

Alternate low-powered embedded FPGAs are virtualized in existing research [25,
26]. These FPGAs are designed for specific tasks and can be dynamically and
partially reconfigured for other tasks.

In summary, we can conclude that existing accelerator virtualization will not
work in a cloud-edge environment for a number of reasons. Firstly, each
solution is intended to operate within a specific environment comprising only
specific devices. Cloud-edge computing is a heterogeneous environment and this
needs to be accounted for. Furthermore, additional overheads can arise with
cloud-edge computing, such as latency from wireless communication, that these
solutions have not considered. Lastly, the research discussed is designed for
virtual machines unlike light-weight deployments, such as containers, employed
in a cloud-edge deployment.

##  III Motivation

Figure 4: Architecture of the AVEC framework

The following aspects in relation to design, environment, and usability have
been factored in for the accelerator virtualization framework AVEC that is
proposed in this paper for a cloud-edge computing environment:

(1) Computationally intensive workloads require hardware acceleration – Many
workloads that will be executed in the cloud-edge environment will be
computationally intensive. For example, deep learning is an important class of
workloads that has found multiple applications in modern mobile apps [3, 4]
and smart cities [27]. In the cloud-edge environment, these workloads can be
partitioned and executed across cloud and edge resources for performance gain
[7, 28]. However, the use of hardware accelerators can enhance the performance
of these workloads. In the design of AVEC, the execution of computationally
intensive workloads, such as deep learning, have been considered.

(2) Accelerators need to be brought to the edge – Although computationally
intensive workloads can rely on accelerators in the cloud, there is
performance gain when they are located at the edge of the network [5]. A
collection of end-user devices can access these accelerators on the edge and
enhance the performance of workloads. Or a cloud application can offload
selected services to an edge resource hosting accelerators for servicing user
requests closer to the source. In both cases, response time of an application
can be minimised and more computationally intensive workloads can get executed
closer to their source for privacy reasons. In AVEC, processing data nearer to
the source has been taken into account.

(3) Multiple workloads need to share edge accelerators – The edge environment
is anticipated to be busy with users since billions of end-user devices will
be connected. Cisco estimated that there would be around 50 billion connected
devices by 2020 [29]. Therefore, it is essential that multiple devices can
share an edge accelerator. While moving accelerators to the edge make them
more accessible to end devices, solutions designed to make them accessible
should be able to execute multiple workloads concurrently. These will need to
isolate the user space and provide safe memory access on the hardware
accelerator. In this paper, the virtualization solution has been designed to
offer isolation of user space in future versions, thus enabling the execution
of multiple workloads at the edge.

(4) Minimum source code modification should be required – An ideal
acceleration solution offered at the edge should be less intrusive (i.e.
should not require low-level source code modification). Giunta et al. [22]
note the importance of transparency in this manner in their work. Instead the
solution should be flexible so that existing distributed applications that
leverage accelerators can be scheduled on accelerators whether they are
located on an edge node, the micro-cloud or the cloud, based on availability
and proximity to the end-user for achieving an agreeable quality of service.
This design criteria has been taken into account while designing the AVEC
accelerator virtualization framework in which no source code modification is
required. The framework intercepts calls to the accelerator from a running
application and redirects them suitably to a remote accelerator.

(5) Overheads should not offset performance benefit – In providing an
accelerator virtualization solution overheads will be inevitably incurred [5].
These overheads are due to the additional operations, such as transferring
data from an end-user device to GPU kernels that execute remotely and
receiving the results of execution back on the device. However, these
overheads can be minimised by using efficient data transfer and execution
strategies and offsetting the overheads by performance gains. In this paper,
the virtualization solution attempts to minimise the overheads and offset them
by achieving a performance gain on deep learning workloads.

##  IV AVEC Framework

Figure 5: Sequence diagram of the AVEC framework

This section presents the first prototype of the AVEC framework. AVEC aims to
provide remote GPU acceleration for deep learning libraries, such as Caffe
[14]. The Caffe library is open source and offers deep learning algorithms and
models for C++ and Python. It uses the CUDA framework for acceleration on
Nvidia GPUs. AVEC executes Caffe kernels within cloud-edge computing using an
accelerator virtualization approach. AVEC is designed as a middleware that
allows calls to the Caffe library to be executed in a remote accelerator by
forwarding the input parameters of the GPU kernel to a remote destination in
the edge or the cloud.

The AVEC framework is demonstrated on the Caffe deep learning library using
the OpenPose application [15]. AVEC does not require that the source code of
OpenPose is modified, rather operates as an interception library that runs
outside the application level.

Figure 6: Organisation of experimental test-bed | Device | Edge Node | Cloud Node
---|---|---|---
GPU Resource | Jetson Nano | Jetson TX2 | RTX 2600
Flops (GPU) | 235 GFLOPS | 750 GFLOPS | 6.5 TFLOPS
GPU Memory | 4 GB LPDDR4 | 8 GB LPDDR4 | 6 GB
(Shared) | (Shared) | GDDR6
CPU Memory | 32 GB
DDR4
GPU Cores | 128 | 256 | 1920
CPU Cores | 4 | 4 | 8
TABLE I: Hardware used in the lab-based test environment

OpenPose takes media input, such as videos or images, and detects people
within the frames and then highlights the pose of the detected person. Figure
3 shows an example output. Person and pose detection requires the execution of
a GPU kernel using the CUDA language. AVEC aims to deploy the OpenPose
accelerator kernel on a remote destination node if a GPU is not available or
is resource constrained on a given device.

The architecture of the AVEC framework is shown in Figure 4. Its goal is to
execute a GPU kernel from a host node in a remote accelerator in a destination
node. For this AVEC extracts the data required to initiate the kernel in the
destination node, runs it there and then sends the output of the kernel back
to the host with the values needed to continue application execution. The
framework employs a communication module to send data between the host and
destination nodes. The original application on the host continues executing
the application using the output of the remote kernel as if it was locally
executed.

Figure 5 shows the sequence of activities between the host and destination
when using OpenPose on the AVEC framework. The communication between the host
and destination is via TCP/IP using Boost ASIO [30]. The host node in the
cloud-edge environment may be a device or an edge node that does not have an
accelerator or a sufficiently powerful accelerator. The host node executes the
original application that requires the deep learning library (Caffe) and
AVEC’s pre-loaded interception library to allow access to the virtualized GPU
accelerator. When the application calls functions within Caffe, those calls
are redirected to the interception library via API interception. The
interception library creates a connection between the host and destination
nodes (both in a cloud-edge network) using a TCP/IP connection. The
destination node hosts a physical GPU and the original Caffe library. Once the
destination node receives the forwarded requests, it executes the kernels
required by the host node on the physical GPU of the destination node. The
output of these functions is sent back to the user node in the same manner as
they were received. When the output is received, the application continues
executing on the host node.

Within the Caffe library that OpenPose uses, the call to the Caffe kernels has
been replaced with a function call to AVEC’s communications module. AVEC
extracts arguments from the host and sends them to the destination. To this
end, the Caffe prototxt file that contains the structure of the neural network
must be initially sent to the destination. A caching mechanism is implemented
so that these do not have to be sent for different executions of the kernel.
In addition to the Caffe files, the input video frames to be analysed must
also be sent with metadata regarding the frame, such as its size and aspect
ratio. The process of assembling the frame into an array is carried out on the
host to prevent unnecessary work being done in the destination node.

##  V Experimental Studies

Number of Images | Cloud Node | Edge Node | Device
---|---|---|---
64 | 8.13 | 69.47 | 130.77
128 | 13.82 | 134.02 | 256.64
256 | 25.98 | 258.19 | 497.06
TABLE II: Execution time (s) of the OpenPose application | Device | Edge Node | Cloud Node
---|---|---|---
Model transfer time (s) | 6.43 | 5.937 | 1.757
TABLE III: Time taken to transfer COCO model to the GPU

This section presents experimental studies on AVEC to implement GPU
virtualization in a cloud-edge environment. The capabilities of AVEC are
tested using the OpenPose application.

####  V-1 Experimental Setup

A lab-based test environment comprising resources that are comparable to a
user device, an edge node and a cloud resource are setup as shown in Table I.
Experiments are carried out with a one to one connection, such that the host
device (Jetson Nano) will only offload to one destination node during the
execution, as shown in Figure 6.

Tests are carried out on images and videos. The COCO dataset [31] is used to
test images (2017 version). For videos, a sample provided by OpenPose is used.
The images chosen from the COCO dataset are randomly selected and set into
batches of 64, 128, and 256 images. These images vary in number of objects
(i.e. persons) present in them (images that do not have persons are also
present). The same batches of images are used for all tests to ensure that the
results are comparable. OpenPose will superimpose the pose of persons present
on the image as shown in Figure 3.

(a) Image batches

(b) Video

Figure 7: Execution time for different image batches and video when using AVEC

####  V-2 AVEC for Images

The first tests were carried out using the randomly selected batches of 64,
128 and 256 images with different resolutions. Each time a selection of images
was taken from the COCO dataset, a different seeding was used so that when the
random selection of images to sort into each batch took place, the same images
would not be placed in all image batches. This was done to ensure that a
variety of images were used across the different tests. The original OpenPose
application was timed and tested with the image batches and the results are
shown in Table II. It is observed that the time taken to process each image
batch is approximately twice the time taken for a previous batch.

The same images were executed using AVEC on the same image batches as in the
above test. The results shown in Figure 7(a) indicate that an increase in
speed is achieved as AVEC offloads the computationally intensive components of
OpenPose to remote GPUs. The time is reduced by a significant amount when the
edge node is used, but when the device offloads to the cloud node a
substantial speedup is observed. This is as expected, since the resources
available on the GPU increases from the device through to the edge and the
cloud.

Table III shows the time taken to copy the deep learning model to the GPU. The
model needs to be transferred to the GPU. In the experiments, the COCO model
requires up about 5.5GB of memory on the GPU and is copied once during the
initialisation to the GPU and remains there throughout execution.

####  V-3 AVEC for Videos

The video used for this test is an 8 seconds clip consisting of 204 frames
with a resolution of 368 x 656 pixels. It has a large number of objects (i.e.
persons) present in each frame, therefore, rendering is required on the GPU.
The results obtained from offloading are shown in Figure 7(b). A similar trend
as seen for the images is noted, and AVEC provides larger amount of
improvement with the cloud node than with the edge node. The speedup obtained
on both images and video using AVEC is shown in Table IV. It is observed that
AVEC can deliver up to 7.48x speedup in the test with the video when
offloading the workload to the cloud.

Test type | Device-to-edge offload | Device-to-cloud offload
---|---|---
64 Images | 1.32x | 3.06x
128 Images | 1.32x | 2.83x
256 Images | 1.40x | 2.91x
Video | 1.45x | 7.48x
TABLE IV: Speedup for images and video using AVEC

Another dimension to quantify the performance benefit of AVEC is by measuring
the frames per second (FPS) processed. This shown in Table V. As expected, an
improvement in FPS is noted when offloading to more powerful GPUs. However,
the edge is still quite limited in its computational capabilities and
therefore there is limited acceleration.

Another observation is that the video performs better than the images. This
may be due to the video not being saved to memory after each frame is
processed, whereas each processed image is saved with the OpenPose prediction
transposed onto it.

Test | Device | Edge | Cloud | Device-to-edge | Device-to-cloud
---|---|---|---|---|---
type |  | using AVEC | using AVEC
Images | 0.5 | 1.1 | 10.5 | 0.65 | 2
Video | 0.4 | 0.7 | 9 | 0.6 | 3.1
TABLE V: Comparing frames per second processed on the device, edge, and cloud
resource natively and when offloading from the device to the edge and to the
cloud using AVEC

####  V-4 Profiling AVEC performance

The Nvidia profiler (nvprof) was used to measure the time spent for
computation of the GPU kernels, data transfers (i.e. communication), and other
overheads. It is noted that AVEC reduces the execution time on the GPU when
offloading, but incurs overheads due to transferring GPU kernel arguments and
results between nodes.

When running the OpenPose application natively without the use of AVEC, 27 GPU
kernels are used through the execution sequence. Using AVEC, 13 kernels are
executed on the host device and 17 kernel executions on the destination node
(edge or cloud). The three additional kernels are CUDA calls required on both
the host and destination. The host retains all OpenPose related kernels (e.g.
renderPoseCoco), but the destination node executes all Caffe related kernels.

Figure 8: Profiling single frame execution when using AVEC

(a) 64 images

(b) 128 images

(c) 256 images

(d) Video

Figure 9: Profiling of different image batches and video when using AVEC.
‘GPU’ refers to time taken for execution on the GPU, ‘Communication’ refers to
time taken to send data between nodes, and ‘Other’ refers to time spent on
OpenPose specific tasks carried by the CPU on the host device.

The time taken to complete a forward pass, and thus make a real time pose
estimation depends on the frame resolution (larger frames take a longer time
to process). The results discussed below are from using the sample video used
in previous experiments.

One execution cycle is defined as the time taken to (i) send the frame data as
float, (ii) send the resolution as an array, (iii) send the frame size as an
integer, (iv) execute the forward pass and return the results. The frame width
variable is two integers (8 bytes). The frame size is a single integer (4
bytes). The frame size will be the dimensions of the frame (1 x 3 x 368 x
656). Finally, the output result sent back to the host node will always be the
dimensions of the frame divided by a constant c. This constant depends on the
model used (in this case the constant for the model used is 3.368421). We can
formalise the amount of data transferred (DT) per frame as shown in Equation
1. In these experiments, the amount of data transferred per frame is
approximately 3.75 MB.

$DT=(2*4)+(1*4)+(Dims*4)+((Dims/c)*4)$ (1)

Figure 8 provides a breakdown of the profiled data when natively executing on
the device and when offloading using AVEC. When considering device-to-cloud
offloading, on average each frame requires 0.15 seconds of processing time for
a full execution cycle. Approximately 0.1 seconds are spent in the GPU
executing kernels to carry out the forward pass operation. An overhead of 0.05
seconds is incurred in transferring data between the device and the cloud.
When considering device-to-edge offloading approximately 1.48 seconds are
required per frame and the time taken in the GPU of the edge node is 1.24
seconds. An overhead of 0.24 seconds is incurred by transferring data. The
time taken for the host device to complete its own forward pass is around 2.5
seconds.

A breakdown of the time spent by AVEC on different image batches and video is
shown in Figure 9. Across all test types it is noted that the GPU execution
time is reduced with AVEC. When offloading to the cloud node the communication
overhead is larger than the execution time on the GPU. As expected, the
communication overhead incurred by the edge is greater than the cloud although
the same amount of data is transferred. This is due to the differences in
hardware, resulting in the Communication Module of AVEC being executed faster
on the CPU of the cloud node than on the edge.

##  VI Conclusions

In this paper is presented a first prototype of the AVEC framework
(accelerator virtualization in cloud-edge computing). The research
investigated the potential of GPU accelerator virtualization to improve the
performance of deep learning workloads in cloud-edge environments. More
specifically three research questions were addressed: (Q1) How can accelerator
virtualization be offered at the edge given its limited compute resources?
(Q2) What overheads are incurred in virtualizing accelerators at the edge?
(Q3) Can GPU accelerator virtualization at the edge improve the execution
performance of real-world applications?

AVEC is underpinned by a virtualization technique that employs API
interception and forwarding, that requires no source code modification of
applications. A thorough evaluation of the overheads that are incurred in
using AVEC was performed and its impact on the execution of workloads was
analysed. A real world use-case was employed, namely the OpenPose application
using the Caffe deep learning library. OpenPose is a computationally intensive
application that low powered embedded devices cannot accelerate. However,
performance acceleration was achieved when using AVEC. It was noted that
although there are communication overheads in transferring data from devices
to a remote GPU, AVEC still delivers a speedup reaching up to a maximum of
7.48x. The first implementation of AVEC has the following limitations
providing opportunities for further research:

(i) Only considers deep learning libraries: Currently, AVEC only incorporates
the Caffe deep learning library. In the future, it is envisioned that AVEC
will incorporate other deep learning libraries to deliver a more comprehensive
remote accelerator virtualization solution and thus enable AVEC to become more
mature.

(ii) No mechanism to enable migration of workloads between accelerators: Due
to the dynamic nature of cloud-edge environments and its benefit for mobile
users, AVEC must be capable of moving workloads to alternate locations based
on the requirements of users or their workloads. This is implemented by
packaging workloads into deployable containers and moving them across
accelerators in a cloud-edge network. Migration of workloads will also enhance
fault-tolerance in AVEC in the event of nodes that host accelerators failing.

(iii) Lack of device-aware scheduling: Given the heterogeneous nature of
cloud-edge computing, AVEC needs to be aware of the accelerators that are
available within the network that are capable of running a given workload.
Moreover, metrics such as accelerator usage, power usage or latency of these
accelerators should be taken into account for scheduling a given workload.

## Acknowledgements

We gratefully acknowledge the support of NVIDIA Corporation with the donation
of the Jetson TX2 Development Kit used for this research.

## References

  * [1] B. Varghese _et al._ , “Challenges and Opportunities in Edge Computing,” in _Proceedings of the IEEE International Conference on Smart Cloud_ , 2016, pp. 20–26.
  * [2] M. Satyanarayanan, “The Emergence of Edge Computing,” _Computer_ , vol. 50, no. 1, pp. 30–39, 2017.
  * [3] N. Wang _et al._ , “ENORM: A Framework for Edge NOde Resource Management,” _IEEE Transactions on Services Computing_ , vol. 13, no. 6, pp. 1086–1099, 2020.
  * [4] N. Wang _et al._ , “DYVERSE: DYnamic VERtical Scaling in Multi-tenant Edge Environments,” _Future Generation Computer Systems_ , vol. 108, pp. 598–612, 2020.
  * [5] B. Varghese _et al._ , “Accelerator Virtualization in Fog Computing: Moving From the Cloud to the Edge,” _IEEE Cloud Computing_ , vol. 5, no. 6, pp. 28–37, 2018.
  * [6] Y. Bengio, “Learning Deep Architectures for AI,” _Foundations and Trends in Machine Learning_ , vol. 2, no. 1, pp. 1–127, 2009.
  * [7] L. Lockhart _et al._ , “Scission: Performance-driven and Context-aware Cloud-Edge Distribution of Deep Neural Networks,” in _IEEE/ACM 13th International Conference on Utility and Cloud Computing_ , 2020, pp. 257–268.
  * [8] NVIDIA Corporation. CUDA Zone $|$ NVIDIA Developer. [Online]. https://developer.nvidia.com/CUDA-zone
  * [9] The Khronos Group Inc. OpenCL Overview. [Online]. https://www.khronos.org/opencl/
  * [10] F. Silla _et al._ , “On the Benefits of the Remote GPU Virtualization Mechanism: The rCUDA Case,” _Concurr. Comput. Pract. Exp._ , vol. 29, no. 13, p. e4072, 2017.
  * [11] G. Giunta _et al._ , “A GPGPU Transparent Virtualization Component for High Performance Computing Clouds,” in _Euro-Par 2010 - Parallel Processing_ , 2010, pp. 379–391.
  * [12] C. Reaño _et al._ , “Exploring the Use of Remote GPU Virtualization in Low-Power Systems for Bioinformatics Applications,” in _Proceedings of the 47th International Conference on Parallel Processing Companion_ , 2018, pp. 8:1–8:8.
  * [13] R. Montella _et al._ , “Accelerating Linux and Android Applications on Low-Power Devices through Remote GPGPU Offloading,” _Concurr. Comput. Pract. Exp._ , vol. 29, no. 24, p. e4286, 2017.
  * [14] Y. Jia _et al._ , “Caffe: Convolutional Architecture for Fast Feature Embedding,” in _Proceedings of the 22nd ACM International Conference on Multimedia_ , 2014, p. 675–678.
  * [15] Z. Cao _et al._ , “OpenPose: Realtime Multi-Person 2D Pose Estimation Using Part Affinity Fields,” _IEEE Transactions on Pattern Analysis and Machine Intelligence_ , vol. 43, no. 1, pp. 172–186, 2021.
  * [16] S. Xiao _et al._ , “VOCL: An Optimized Environment for Transparent Virtualization of Graphics Processing Units,” in _Innovative Parallel Computing (InPar)_ , 2012, pp. 1–12.
  * [17] L. Shi _et al._ , “vCUDA: GPU-Accelerated High-Performance Computing in Virtual Machines,” _IEEE Transactions on Computers_ , vol. 61, no. 6, pp. 804–816, 2012.
  * [18] V. Gupta _et al._ , “GViM: GPU-Accelerated Virtual Machines,” in _Proceedings of the 3rd ACM Workshop on System-Level Virtualization for High Performance Computing_ , 2009, p. 17–24.
  * [19] M. Oikawa _et al._ , “DS-CUDA: A Middleware to Use Many GPUs in the Cloud Environment,” in _SC Companion: High Performance Computing, Networking Storage and Analysis_ , 2012, pp. 1207–1214.
  * [20] T. Liang _et al._ , “GridCuda: A Grid-Enabled CUDA Programming Toolkit,” in _IEEE Workshops of International Conference on Advanced Information Networking and Applications_ , 2011, pp. 141–146.
  * [21] Y. Lin _et al._ , “qCUDA: GPGPU Virtualization for High Bandwidth Efficiency,” in _IEEE International Conference on Cloud Computing Technology and Science (CloudCom)_ , 2019, pp. 95–102.
  * [22] R. Montella _et al._ , “On the Virtualization of CUDA Based GPU Remoting on ARM and X86 Machines in the GVirtuS Framework,” _International Journal of Parallel Programming_ , vol. 45, no. 5, p. 1142–1163, 2017.
  * [23] B.-Y. Huang _et al._ , “qCUDA-ARM: Virtualization for Embedded GPU Architectures,” in _Internet of Vehicles. Technologies and Services Toward Smart Cities_ , 2020, pp. 284–302.
  * [24] RAPID project. Heterogeneous Secure Multi-level Remote Acceleration Service for Low Power Integrated Systems and Devices. [Online]. http://www.rapid-project.eu
  * [25] T. Xia _et al._ , “Hypervisor Mechanisms to Manage FPGA Reconfigurable Accelerators,” in _International Conference on Field-Programmable Technology_ , 2016, pp. 44–52.
  * [26] J. Mandebi Mbongue _et al._ , “FPGA Virtualization in Cloud-Based Infrastructures Over Virtio,” in _IEEE 36th International Conference on Computer Design_ , 2018, pp. 242–245.
  * [27] NVIDIA Corporation. NVIDIA Metropolis. [Online]. https://www.nvidia.com/en-gb/autonomous-machines/intelligent-video-analytics-platform/
  * [28] Y. Kang _et al._ , “Neurosurgeon: Collaborative Intelligence Between the Cloud and Mobile Edge,” in _Proceedings of the 22nd International Conference on Architectural Support for Programming Languages and Operating Systems_ , 2017, pp. 615–629.
  * [29] D. Evans, “How the Next Evolution of the Internet Is Changing Everything,” in _Cisco White Paper_ , 2011, pp. 1–11.
  * [30] C. Kohlhoff. Boost.Asio. [Online]. https://www.boost.org/doc/libs/1_69_0/doc/html/boost_asio.html
  * [31] T.-Y. Lin _et al._ , “Microsoft COCO: Common Objects in Context,” in _Computer Vision – ECCV_ , 2014, pp. 740–755.


Single-RF Multi-User Communication Through
Reconfigurable Intelligent Surfaces: An Information-Theoretic Analysis

Roy Karasik, Student Member, IEEE,
Osvaldo Simeone, Fellow, IEEE,
Marco Di Renzo, Fellow, IEEE,
and Shlomo Shamai (Shitz), Life Fellow, IEEE
This work has been supported by the European Research Council (ERC) and by the Information and Communication Technologies (ICT) under the European Union’s Horizon 2020 Research and Innovation Programme (Grant Agreement Nos. 694630, 725731, and 871464).
R. Karasik and S. Shamai are with the Department of Electrical Engineering, Technion — Israel Institute of Technology, Haifa 32000, Israel. (royk@campus.technion.ac.il).
O. Simeone is with the Centre for Telecommunications Research,
Department of Informatics, King’s College London, London WC2R 2LS, U.K.
M. Di Renzo is with Université Paris-Saclay, CNRS, CentraleSupélec, Laboratoire des Signaux et Systèmes, 3 Rue Joliot-Curie, 91192 Gif-sur-Yvette, France.

RISs are typically used in multi-user systems to mitigate interference among active transmitters. In contrast, this paper studies a setting with a conventional active encoder as well as a passive encoder that modulates the reflection pattern of the RIS. The RIS hence serves the dual purpose of improving the rate of the active encoder and of enabling communication from the second encoder. The capacity region is characterized, and information-theoretic insights regarding the trade-offs between the rates of the two encoders are derived by focusing on the high- and low-power regimes.

§ INTRODUCTION

A RIS is a nearly-passive device that can shape the wireless propagation channel by applying phase shifts to the incident signals <cit.>.
In MU systems, RISs can help mitigate inter-user interference and obtain beamforming gain for standard active transmitters <cit.>. To this end, the configuration of the RIS is kept fixed for the duration of a coherence interval and optimized to maximize some function of the achievable rates <cit.>. In this paper, we study a different use of RISs, whereby a single active transmitter coexists with a passive user, having no direct RF chain, that conveys its own message by modulating the reflection pattern of the RIS (see <ref>).

regular polygon,
regular polygon sides=3,
shape border rotate=180,

partial ellipse/.style args=#1:#2:#3
insert path=+ (#1:#3) arc (#1:#2:#3)

(EN) [thick,draw,minimum width=1cm,minimum height=1cm, font=] at (0,0) Encoder 1;
(corner1) [right=0.1cm of EN] ;
[draw] (EN.east) – (corner1.center);
(Ant1) [ant, above=0.2cm of corner1,scale=0.5] ;
[draw] (corner1.center) – (Ant1);

(user) [thick,draw,minimum width=1cm,minimum height=1cm,right = 4.3cm of EN.center, font=]  Decoder;
(RxP1) [left=0.15cm of user] ;
(RxP6) [left=0.40cm of user] ;
(RxP2) [above=0.05cm of RxP1.center] ;
(RxP3) [above=0.05cm of user.west] ;
(RxP4) [below=0.05cm of RxP6.center] ;
(RxP5) [below=0.05cm of user.west] ;
[draw] (RxP3.center) – (RxP2.center);
[draw] (RxP5.center) – (RxP4.center);
(Ant2) [ant, above =0.2cm of RxP2.center,scale=0.5] ;
[draw] (RxP2.center) – (Ant2);
(Ant3) [ant, above =0.2cm of RxP4.center,scale=0.5] ;
[draw] (RxP4.center) – (Ant3);
[font=, above=0.0mm of user.north,xshift=0mm] $N$ antennas;
[thick] ($(user.west)+(-1.2mm,0mm)$) [partial ellipse=50:310:0.7mm and 2.7mm];

(IRS) [thick,draw,minimum width=1.4cm,minimum height=1.4cm, font=] at ($(EN)!0.5!(user)+(0,1.8cm)$) ;
(R1) [thick,fill=blue!50,draw,minimum width=0.26cm,minimum height=0.26cm] at ([yshift=5.2mm,xshift=-5.2mm]IRS.center);
(R2) [thick,fill=blue!50,draw,minimum width=0.26cm,minimum height=0.26cm] at ([yshift=-5.2mm,xshift=-5.2mm]IRS.center);
(R3) [thick,fill=blue!50,draw,minimum width=0.26cm,minimum height=0.26cm] at ([yshift=5.2mm,xshift=+5.2mm]IRS.center);
(R4) [thick,fill=blue!50,draw,minimum width=0.26cm,minimum height=0.26cm] at ([yshift=-5.2mm,xshift=+5.2mm]IRS.center);
(R5) [thick,fill=blue!50,draw,minimum width=0.26cm,minimum height=0.26cm] at ([yshift=+5.2mm,xshift=+1.7mm]IRS.center);
(R6) [thick,fill=blue!50,draw,minimum width=0.26cm,minimum height=0.26cm] at ([yshift=+5.2mm,xshift=-1.7mm]IRS.center);
(R7) [thick,fill=blue!50,draw,minimum width=0.26cm,minimum height=0.26cm] at ([yshift=-5.2mm,xshift=+1.7mm]IRS.center);
(R8) [thick,fill=blue!50,draw,minimum width=0.26cm,minimum height=0.26cm] at ([yshift=-5.2mm,xshift=-1.7mm]IRS.center);
(R9) [thick,fill=blue!50,draw,minimum width=0.26cm,minimum height=0.26cm] at ([yshift=1.75mm,xshift=-5.2mm]IRS.center);
(R10) [thick,fill=blue!50,draw,minimum width=0.26cm,minimum height=0.26cm] at ([yshift=-1.75mm,xshift=-5.2mm]IRS.center);
(R11) [thick,fill=blue!50,draw,minimum width=0.26cm,minimum height=0.26cm] at ([yshift=1.75mm,xshift=+5.2mm]IRS.center);
(R12) [thick,fill=blue!50,draw,minimum width=0.26cm,minimum height=0.26cm] at ([yshift=-1.75mm,xshift=+5.2mm]IRS.center);
(R13) [thick,fill=blue!50,draw,minimum width=0.26cm,minimum height=0.26cm] at ([yshift=+1.75mm,xshift=+1.7mm]IRS.center);
(R14) [thick,fill=blue!50,draw,minimum width=0.26cm,minimum height=0.26cm] at ([yshift=+1.75mm,xshift=-1.7mm]IRS.center);
(R15) [thick,fill=blue!50,draw,minimum width=0.26cm,minimum height=0.26cm] at ([yshift=-1.75mm,xshift=+1.7mm]IRS.center);
(R16) [thick,fill=blue!50,draw,minimum width=0.26cm,minimum height=0.26cm] at ([yshift=-1.75mm,xshift=-1.7mm]IRS.center);
[font=, above=-0.5mm of IRS.north] RIS;
[font=, right=0.0mm of IRS.east,yshift=2mm,blue!50] $K$ elements;

(EN2P1) at ($(IRS.west)+(0,0mm)$) ;
(EN2) [thick,draw,minimum width=1cm,minimum height=1cm, font=, above = 0.76cm of EN] Encoder 2;

(A) at ($(Ant1)+(2mm,0)$);
[draw,->,line width=0.5mm,red] (A) – node[above,font=,pos=0.15,xshift=-1mm] $\mathbf{h}_i$ coordinate[pos=0.2] (P1) coordinate[pos=0.25] (P2) ($(IRS)+(0,0)$);

(B) at ($(Ant2)+(-2mm,-2mm)$);
[draw,->,line width=0.5mm,red] ($(IRS)+(0,0)$) – node[above,font=,pos=0.8,xshift=1mm] $\mathbf{H}_{r}$ coordinate[pos=0.75] (P3) coordinate[pos=0.8] (P4) (B);

[draw,->,line width=0.5mm,red] (A) – node[above,font=,pos=0.5,yshift=-1mm] $\mathbf{h}_d$ coordinate[pos=0.75] (P5) coordinate[pos=0.78] (P6) (B);

[font=, below=10mm of IRS,red] Wireless Link;

[draw,->,line width=0.2mm] ($(EN.west)-(15mm,0)$) – node[above,font=] $w_1$ node[below,font=] ($nR_1$ bits) ($(EN.west)+(0,0)$);
[draw,->,line width=0.2mm] ($(EN2.west)-(15mm,0)$) – node[above,font=] $w_2$ node[below,font=] ($nR_2$ bits) ($(EN2.west)+(0,0)$);
[draw,->,line width=0.2mm] ($(user.east)-(0mm,0)$) – node[above,font=] $(\hat{w}_1,\hat{w}_2)$ ($(user.east)+(13mm,0)$);

[line width=0.3mm,OliveGreen,->,dashed] (EN2.east) – (EN2P1.center);
[font=, above=0.0mm of EN2.north, OliveGreen, xshift=-3mm] Control Link ($\text{Rate}=1/m$);

In the system under study, Encoder 1 is active and it encodes its message $w_1$ into a codeword of $n$ symbols sent on the wireless link; whereas Encoder 2 is passive and it encodes the message $w_2$ into a control action, which is sent on the control link to the RIS at a rate of one action every $m$ channel symbols.

With reference to <ref>, the RIS is accordingly used for the dual purpose of enhancing the rate of the active encoder (Encoder 1) and of enabling communication for the passive encoder (Encoder 2). Unlike prior work <cit.> that focused on a specific transmission strategy, this paper concentrates on the information-theoretic analysis of the rate trade-offs between the two encoders, providing fundamental insights.

Related Work: A comprehensive survey of the state-of-the-art on RIS-aided MU systems is available in <cit.>. As notable representative examples of works involving active transmitters, the maximization of the weighted sum-rate in RIS-aided MU systems was studied in <cit.>, whereas references <cit.> focused on optimizing the energy efficiency, and papers <cit.> on physical-layer security and outage-probability enhancements. A MU system with an active transmitter and a passive encoder, akin to <ref>, was proposed in <cit.> by assuming binary modulation, a single receiver antenna, and a specific successive interference cancellation decoding strategy.

From an information-theoretic perspective, the single-RF MU communication system depicted in <ref> can be viewed as a MAC with both multiplicative and additive elements. The capacity of the Gaussian multiplicative MAC was derived in <cit.> for two active encoders. The capacity region of a backscatter multiplicative MAC, which can be viewed as a special case of the RIS-aided MU communication system in <ref> with one reflecting element, was studied in <cit.>. Under the assumptions of a single receiver antenna and Gaussian codebooks, this work shows that conventional time-sharing schemes are suboptimal in the high-power and weak-backscatter regimes. The capacity of an RIS-aided single-user channel was derived in <cit.>.

Main Contributions: In this paper, we study the RIS-aided MU system illustrated in <ref>, in which Encoder 1 is active, whereas Encoder 2 can only alter the reflection pattern of the RIS in a passive manner. We derive the capacity region under the practical assumptions of a multi-antenna decoder, a finite-input constellation, and a set of discrete phase shifts at the RIS. Then, we specialize the results for the high- and low-power regimes, showing that: (i) for sufficiently high transmission power, both encoders can simultaneously communicate at maximum rate; and (ii) in the low-power regime, Encoder 1 can achieve maximum rate if and only if Encoder 2 does not communicate, while Encoder 2 can achieve its maximum rate while still enabling non-zero-rate communication for Encoder 1. Finally, numerical examples demonstrate the dual role of the RIS as means to enhance the transmitted signal on the one hand and as the enabler of MU communication on the other hand.

Random variables, vectors, and matrices are denoted by lowercase, boldface lowercase, and boldface uppercase Roman-font letters, respectively. Realizations of random variables, vectors, and matrices are denoted by lowercase, boldface lowercase, and boldface uppercase italic-font letters, respectively. For example, $x$ is a realization of random variable $\mathrm{x}$, $\bm{x}$ is a realization of random vector $\mathbf{x}$, and $\bm{X}$ is a realization of random matrix $\mathbf{X}$.
For any positive integer $K$, we define the set $[K]\triangleq \{1,2,\ldots,K\}$.
The cardinality of a set $\mathcal A$ is denoted as $|\mathcal{A}|$.
The $\ell^2$-norm and the conjugate transpose of a vector $\bm{v}$ are denoted as $\lVert\bm{v}\rVert$ and $\bm{v}^\ast$, respectively.
$\diag(\bm{x})$ represents a diagonal matrix with diagonal given by the vector $\bm{x}$.
The vectorization of matrix $\bm{H}$, i.e., the operator that stacks the columns of $\bm{H}$ on top of one another, is denoted by $\stack(\bm{H})$.
The Kronecker product $\bm{I}_m\otimes\bm{B}$ of the identity matrix of size $m$ and matrix $\bm{B}$  is denoted as $\bm{B}^{m\otimes}$.

§ SYSTEM MODEL

We consider the system depicted in <ref>
in which two encoders communicate with a decoder equipped with $N$ antennas over a quasi-static fading channel in the presence of an RIS that comprises $K$ nearly-passive reconfigurable elements. Encoder 1 is equipped with a single-RF transmitter and encodes its message $w_1\in[2^{nR_1}]$ of rate $R_1$ [bits/symbol] into a codeword of $n$ symbols sent on the wireless link to the decoder. In contrast, Encoder 2 encodes its message $w_2\in[2^{nR_2}]$ of rate $R_2$ [bits/symbol] in a passive manner by modulating the reflection pattern of the RIS. The reflection pattern is controlled through a rate-limited control link, and is defined by the phase shifts that each of the $K$ RIS elements applies to the impinging wireless signal. Encoder 2 represents, for example, a sensor embedded in the RIS that applies metasurface-based modulation in order to convey its sensed data without emitting a new radio wave <cit.>.

A coding slot consists of $n$ symbols, which are divided into $n/m$ blocks of $m$ symbols each, with $n/m$ assumed to be integer. Specifically, the codeword transmitted by Encoder 1 as a function of message $w_1$ occupies the entire coding slot, and it includes $n$ symbols from a constellation $\mathcal{S}$ of $S=|\mathcal S|$ points. Furthermore, the RIS is controlled by Encoder 2 by selecting the phase shift applied by each of the $K$ elements of the RIS from a finite set $\mathcal{A}$ of $A=|\mathcal{A}|$ distinct hardware-determined values as a function of the message $w_2$.
Due to practical limitations on the RIS control rate, the phase shifts can only be modified once for each block of $m$ consecutive transmitted symbols.
During the $t$th block, the fraction of the codeword of Encoder 1 consisting of $m$ transmitted symbols is denoted by $\mathbf{s}(t)=(\mathrm{s}_{1}(t),\ldots,\mathrm{s}_{m}(t))^\intercal\in\mathcal S^{m\times 1}$, and is assumed to satisfy the power constraint

The phase shifts applied by the RIS in the $t$th block are denoted by the vector

with $\uptheta_{k}(t)\in\mathcal{A}$ being the phase shift applied by the $k$th RIS element, $k\in[K]$.

We assume quasi-static flat-fading wireless channels, which remain fixed throughout a coding slot. Specifically, the channel from Encoder 1 to the decoder is denoted by vector $\mathbf{h}_d\in\mathbb C^{M\times 1}$; the channel from Encoder 1 to the RIS is denoted by the vector $\mathbf{h}_i\in\mathbb C^{K\times 1}$; and the channel from the RIS to the $N$ receiving antennas is denoted by the matrix $\mathbf{H}_r\in\mathbb C^{N\times K}$. Furthermore, we assume that $\mathbf{h}_d$, $\mathbf{h}_i$, and $\mathbf{H}_r$ are drawn from a continuous distribution. Finally, we denote the signal received by the $N$ antennas for the $q$th transmitted symbol in block $t\in[n/m]$ by $\mathbf{y}_{q}(t)\in\mathbb C^{N\times 1}$, $q\in[m]$. The overall received signal matrix $\mathbf{Y}(t)=(\mathbf{y}_{1}(t),\ldots,\mathbf{y}_{m}(t))\in\mathbb C^{N\times m}$ in the $t$th sub-block can hence be written as
𝐘(t)   =   √(P)𝐇_r(e^j(t))𝐡_i𝐬^⊺(t)+𝐡_d𝐬^⊺(t)+𝐙(t)

=   √(P)𝐇_rie^j(t)+𝐡_d𝐬^⊺(t)+𝐙(t),

where $P>0$ denotes the transmission power of Encoder 1; the matrix $\mathbf{H}_{ri}\triangleq \mathbf{H}_r\diag(\mathbf{h}_i)\in\mathbb{C}^{N\times K}$, combines the channels $\mathbf{h}_i$ and $\mathbf{H}_r$;
and the matrix $\mathbf{Z}(t)\in\mathbb C^{N\times m}$, whose elements are iid as $\mathcal{CN}(0,1)$, denotes the additive white Gaussian noise at the receiving antennas.
In order to characterize the distribution of the output signal $\mathbf{Y}(t)$ in (<ref>), we vectorize it as

where we have defined the vector $\mathbf{z}(t)\triangleq\stack(\mathbf{Z}(t))\in\mathbb{C}^{Nm\times 1}$.

We assume that both the encoders and the decoder have perfect CSI, in the sense that the channel matrix $\mathbf{H}_{ri}$ and channel vector $\mathbf{h}_d$ are known.
Having received signal $\mathbf{y}(t)$ in (<ref>) for $t\in[n/m]$, the decoder produces the estimates $\hat{w}_\ell=\hat{w}_\ell(\mathbf{y}(1),\ldots,\mathbf{y}(n/m),\mathbf{H}_{ri},\mathbf{h}_d)$, for $\ell=1,2$, using knowledge of the CSI.
Given channel realizations $\bm{H}_{ri}$ and $\bm{h}_d$, a rate pair $\rb{R_1(\bm{H}_{ri},\bm{h}_d),R_2(\bm{H}_{ri},\bm{h}_d)}$ is said to be achievable if the probability of error satisfies the limit $\Pr(\hat{w}_1\neq w_1,\hat{w}_2\neq w_2)\rightarrow 0$ when the codeword length grows large, i.e., $n\rightarrow\infty$. The corresponding capacity region $\mathcal{C}(\bm{H}_{ri},\bm{h}_d)$ is defined as the closure of the set of achievable rate pairs.

§ CAPACITY REGION

In this section, we first derive a general characterization of the capacity region $\mathcal{C}(\bm{H}_{ri},\bm{h}_d)$ for the channel in (<ref>). Then, we leverage this result to provide theoretical insights into the trade-offs between the rate of the two encoders in <ref> by focusing on the low- and high-power regimes.

Most existing works on the multiplicative Gaussian MAC <cit.> and on RIS-aided systems (see, e.g., <cit.>) consider Gaussian codebooks for the transmitted signal $\mathbf{s}(t)$. This implies that the resulting achievable rates are formulated in the standard form “$\log_2(1+\text{SNR})$”. In contrast, as described in <ref>, we focus our attention on the more practical model in which the transmitted symbols and the RIS elements' phase response take values from finite sets <cit.>. Therefore, in a manner similar to <cit.>, the expressions for the achievable rates that we present in this section are more complex, and require the following definition.

The CGF of a random variable $\mathrm{u}$ conditioned on a random vector $\mathbf{x}$ is defined as
κ_r(u|𝐱)≜𝔼log_2𝔼e^ru|𝐱,  for r∈ℝ,

and the value of the conditional CGF for $r=1$ is denoted as $\kappa(\mathrm{u}|\mathbf{x})\triangleq\kappa_1(\mathrm{u}|\mathbf{x})$.

We now characterize the capacity region in the form of a union of rate regions, with each region corresponding to rates achievable for a specific choice of encoding distributions $p_{\mathbf{s}}(\bm{s})$ and $p_{\pmb{\uptheta}}(\pmb{\theta})$ of the transmitted symbols $\mathbf{s}(t)$ and RIS phase shifts $\pmb{\uptheta}(t)$ in (<ref>), respectively.

For input distributions $p_{\mathbf{s}}(\bm{s})$ and $p_{\pmb{\uptheta}}(\pmb{\theta})$, let
$\mathcal{R}(p_{\mathbf{s}},p_{\pmb{\uptheta}},\bm{H}_{ri},\bm{h}_d)$ be the set of rate pairs $\rb{R_1(\bm{H}_{ri},\bm{h}_d),R_2(\bm{H}_{ri},\bm{h}_d)}$ such that the inequalities
R_ℓ(H_ri,h_d)≤-Nlog_2(e)-1/mκ(u_ℓ|𝐬_1,_1,𝐳),  ℓ∈{1,2},
hold, where random variable $\mathrm{u}_1$, $\mathrm{u}_2$, and $\mathrm{u}_3$ are defined as
u_1   ≜    -‖𝐳+√(P)H_rie^j_1+h_d^m⊗𝐬_1-𝐬_2‖^2,

u_2   ≜    -‖𝐳+√(P)H_rie^j_1-e^j_2^m⊗𝐬_1‖^2,

u_3   ≜    -‖𝐳+√(P)H_rie^j_1+h_d^m⊗𝐬_1
respectively, with independent random vectors $\mathbf{s}_1,\mathbf{s}_2\sim p_{\mathbf{s}}(\bm{s})$, $\pmb{\uptheta}_1,\pmb{\uptheta}_2\sim p_{\pmb{\uptheta}}(\pmb{\theta})$, and $\mathbf{z}\sim\mathcal{CN}(\mathbf{0},\bm{I}_{Nm})$.
The capacity region $\mathcal{C}(\bm{H}_{ri},\bm{h}_d)$ is the convex hull of the union of the regions $\mathcal{R}(p_{\mathbf{s}},p_{\pmb{\uptheta}},\bm{H}_{ri},\bm{h}_d)$ over all input distributions $p_{\mathbf{s}}(\bm{s})$ and $p_{\pmb{\uptheta}}(\pmb{\theta})$ with $\bm{s}\in\mathcal{S}^{m\times 1}$, $\pmb{\theta}\in\mathcal{A}^{K\times 1}$, such that $\mathbb E[\mathbf{s}^\ast\mathbf{s}]\leq m$.

See Appendix <ref>.

Next, we specialize the results in <ref> to characterize the capacity region in the high- and low-power regimes.

§.§ High-Power Regime

The following corollary shows that the capacity region $\mathcal{C}(\bm{H}_{ri},\bm{h}_d)$ converges to a rectangle as the power of Encoder 1 increases.

For any finite constellation $\mathcal S$ of $S=|\mathcal S|$ points and any set $\mathcal{A}$ of $A=|\mathcal A|$ phases, let $\overline{\mathcal C}$ be the set of rate pairs $\rb{R_1,R_2}$ such that
𝒞≜R_1,R_2:R_1≤log_2(S),  R_2≤K/mlog_2(A).

The capacity region $\mathcal{C}(\bm{H}_{ri},\bm{h}_d)$ converges to $\overline{\mathcal C}$ as the power $P$ increases in the sense that $\mathcal{C}(\bm{H}_{ri},\bm{h}_d)\subseteq \overline{\mathcal C}$, and there exists a sequence of achievable rate pairs $\rb{R_1(\bm{H}_{ri},\bm{h}_d),R_2(\bm{H}_{ri},\bm{h}_d)}\in \mathcal{C}(\bm{H}_{ri},\bm{h}_d)$ such that, almost surely,
lim_P→∞R_1(H_ri,h_d)   =   log_2(S),

lim_P→∞R_2(H_ri,h_d)   =   K/mlog_2(A).

See Appendix <ref>.

<ref> implies that, for sufficiently high power $P$, both encoders can simultaneously achieve their maximum rates. As a result, while not useful in increasing the high-power rate of Encoder 1, the presence of the RIS enables communication at the maximum rate for Encoder 2 without creating deleterious interference on Encoder 1's transmission.

§.§ Low-Power Regime

In this section, we characterize the capacity region $\mathcal{C}(\bm{H}_{ri},\bm{h}_d)$ in the low-power regime. To simplify the analysis, we focus on a system with one receiver antenna, $N=1$, and an RIS control ratio of $m=1$. For this special case, the channel (<ref>) can be written as

where $\mathbf{h}_{ri}\in\mathbb{C}^{K\times 1}$ and $\mathrm{h}_d\in\mathbb{C}$ denote the reflected and direct channel paths, respectively, and $\mathrm{z}(t)\sim\mathcal{CN}(0,1)$ denotes the additive white Gaussian noise. Furthermore, we assume that the phase shift applied by each element of the RIS is chosen from a finite set of $A$ uniformly spaced phases, i.e., $\mathcal A=\{0,2\pi/A,\ldots,2\pi(A-1)/A\}$; and that $\mathcal{S}$ is a zero-mean input constellation, i.e.,

which is known to achieve the minimum energy per bit in many single-user channels <cit.>.

In order to formulate the capacity region in the low-power regime, we define the normalized rate $r_\ell(\bm{h}_{ri},h_d)$, $\ell\in\{1,2\}$, for unit of power as

The capacity region in the low-power regime $\underline{\mathcal{C}}(\bm{h}_{ri},h_d)$ is accordingly defined as the closure of the set of achievable normalized rate pairs (see, e.g., <cit.>).

For input distributions $p_{\mathrm{s}}(s)$ and $p_{\pmb{\uptheta}}(\pmb{\theta})$, let
$\underline{\mathcal{R}}(p_{\mathrm{s}},p_{\pmb{\uptheta}},\bm{h}_{ri},h_d)$ be the set of normalized rate pairs $\rb{r_1(\bm{h}_{ri},h_d),r_2(\bm{h}_{ri},h_d)}$ such that the inequalities
r_ℓ(h_ri,h_d)≤𝔼[u_ℓ]/ln(2),  ℓ∈{1,2},

and r_1(h_ri,h_d)+r_2(h_ri,h_d)≤𝔼[u_3]/ln(2)
hold, where random variable $\underline{\mathrm{u}}_1$, $\underline{\mathrm{u}}_2$, and $\underline{\mathrm{u}}_3$ are defined as
u_1   ≜    |h_ri^⊺e^j_1+h_ds_1-s_2|^2,

u_2   ≜    |h_ri^⊺e^j_1-e^j_2s_1|^2,

u_3   ≜    |h_ri^⊺e^j_1+h_ds_1-h_ri^⊺e^j_2+h_ds_2|^2,
respectively, with independent random variables $\mathrm{s}_1,\mathrm{s}_2\sim p_{\mathrm{s}}(s)$ and random vectors $\pmb{\uptheta}_1,\pmb{\uptheta}_2\sim p_{\pmb{\uptheta}}(\pmb{\theta})$.
The capacity region in the low-power regime $\underline{\mathcal{C}}(\bm{h}_{ri},h_d)$ is the convex hull of the union of the regions $\underline{\mathcal{R}}(p_{\mathrm{s}},p_{\pmb{\uptheta}},\bm{h}_{ri},h_d)$ over all input distributions $p_{\mathrm{s}}(s)$ and $p_{\pmb{\uptheta}}(\pmb{\theta})$ with $s\in\mathcal{S}$, $\pmb{\theta}\in\mathcal{A}^{K\times 1}$, such that $\mathbb E[|\mathrm{s}|^2]\leq 1$.

See Appendix <ref>.

Unlike the high-power regime, the low-power capacity region (<ref>) is not a rectangle, implying that it is not possible for both encoders to communicate at their respective maximum rates. The next corollary elaborates on this point.

Let $\tilde{\pmb{\theta}}$ be the beamforming phase-shift vector maximizing Encoder 1's rate, i.e.,

In the low-power regime, Encoder 1 can achieve its maximum normalized rate

if and only if Encoder 2 does not communicate, i.e., $r_2(\bm{h}_{ri},h_d)=0$. In contrast, if $\lvert\bm{h}_{ri}^\intercal e^{j\tilde{\pmb{\theta}}}+h_d\rvert^2>\lVert \bm{h}_{ri}\rVert^2$, Encoder 2 can achieve its maximum normalized rate

while Encoder 1 communicates at a normalized rate of

See Appendix <ref>.

The asymmetry between Encoder 1 and Encoder 2 revealed by <ref> stems from the fact the, in order for Encoder 1 to obtain its maximum rate in the low-power regime, Encoder 2 needs to steer its phases according to the beamforming solution (<ref>). This in turn makes it impossible to encode additional information for Encoder 2. In contrast, Encoder 2's maximum rate can be obtained as long as Encoder 1's signal is transmitted at the maximum power and can be decoded while treating the modulation of the RIS's phases by Encoder 2 as a nuisance.

We finally remark that <ref> and <ref> imply that time-sharing, which would yield a triangular rate region, is suboptimal in both high- and low-power regimes.
This is in contrast to the multiplicative MAC studied in <cit.> that assumes two standard active encoders with separate power constraints.

§ EXAMPLES

In this section, we provide numerical examples for the capacity region derived in <ref>. For the phase response set, we consider $A$ uniformly spaced phases in the set $\mathcal A=\{0,2\pi/A,\ldots,2\pi(A-1)/A\}$, whereas, for the input constellation, we consider ASK and PSK modulations. In addition, we assume a channel vector $\bm{h}_d$ with elements having amplitude $1$, and a channel matrix $\bm{H}_{ri}$ with elements having amplitude $\alpha>0$, where $\alpha$ denotes the path-loss ratio between the reflected and direct paths. The phases of $\bm{H}_{ri}$ and $\bm{h}_d$ used in this section are summarized in <ref>.

Phases of $\bm{H}_{ri}$ and $\bm{h}_d$ used for the numerical examples

Figure      $\angle \bm{H}_{ri}$ [rad]     $\angle\bm{h}_{d}$ [rad]

<ref>     $\begin{pmatrix}
1.11 & 0.71 & 2.92 & -2.29\\
2.52 & -0.72& 2.21 & 2.1
\end{pmatrix}$
3.11 \\
\end{pmatrix}$

<ref>     $\begin{pmatrix}
-2.63 & -1.22 & -2.92 & -1.52\\
1.85 & 0.36& -0.87 & -2.59
\end{pmatrix}$
2.82 \\
\end{pmatrix}$

Furthermore, the expectation over Gaussian random vectors, e.g., $\mathbf{z}$ in <ref>, is evaluated via a Monte Carlo empirical averages.

In <ref>, we plot the capacity region for an average power constraint of $P=-20$ dB, $N=2$ receiver antennas, $K=4$ RIS elements, $A=2$ available phase shifts, a symbol-to-RIS control rate $m=2$, input constellation given by BPSK, i.e., $\mathcal S=\{-1,1\}$, and a path-loss ratio of $\alpha=0.5$ or $\alpha=1$.

0.5 !

Capacity region for $P=-20$ dB, $N=2$, $K=4$, $A=2$, $m=2$, and BPSK input constellation. The dashed line illustrates the capacity of Encoder 1 for a channel with no RIS.

In addition, we plot for reference the maximum rate achievable by Encoder 1 for a channel with no RIS, i.e., for $\bm{H}_{ri}=\mathbf{0}$.
By comparing with the capacity of the channel with no RIS, <ref> illustrates the two roles of the RIS: The RIS can be used to increase the rate of Encoder 1 by beamforming the transmitted signal, and it can enable communication from a passive secondary user.
In this regard, <ref> demonstrates that the insights obtained in <ref> by studying the low-power regime carry over to more general conditions. In particular, the maximum rate for Encoder 1 is achieved if and only if Encoder 2 does not communicate, while Encoder 2's maximum rate can coexist with a non-zero rate for Encoder 1.

In contrast, by <ref>, for sufficiently high power $P$, both encoders can communicate with the decoder at their respective maximum rates. This is verified by <ref>, where we plot the capacity region for an average power constraint of $P=40$ dB, $N=2$ receiver antennas, $K=4$ RIS elements, $A=2$ available phase shifts, a symbol-to-RIS control rate $m=1$, input constellation given by 4-ASK, i.e., $\mathcal S=\{\sigma,3\sigma,5\sigma,7\sigma\}$ with $\sigma=1/\sqrt{21}$, and a path-loss ratio of $\alpha=1$.

0.5 !

Capacity region for $P=40$ dB, $N=2$, $K=4$, $A=2$, $m=1$, and 4-ASK input constellation. The dashed line illustrates the capacity of Encoder 1 for a channel with no RIS.

Although Encoder 1 does not gain from the existence of the RIS in the high-power regime, the RIS enables MU communication with a single transmitter in a manner that resembles the single-RF MIMO system <cit.>.

§ CONCLUSION

In this work, we have studied the finite-input capacity region of an RIS-aided MU communication system, in which the RIS is not used solely for increasing the rate of an active encoder, but also for enabling communication for a secondary passive encoder. The fundamental trade-offs between the rates of the two encoders were characterized. It was shown that, for sufficiently high power, both users can communicate at their respective maximum rates. Furthermore, in the low-power regime, the maximum rate for the active encoder is achieved if and only if the passive encoder does not communicate, while the passive encoder's maximum rate can coexist with a non-zero rate for the active encoder. Finally, time-sharing was demonstrated to be suboptimal.

§.§ Proof of Proposition <ref>

The model (<ref>) can be viewed as a MAC with inputs $(\mathbf{s},\pmb{\uptheta})$ and output $\mathbf{y}$. Therefore, it follows from the capacity region of the MAC <cit.> that $\mathcal{C}(\bm{H}_{ri},\bm{h}_d)$ is the convex hull of the union of regions $\tilde{\mathcal{R}}(p_{\mathbf{s}},p_{\pmb{\uptheta}},\bm{H}_{ri},\bm{h}_d)$ over all input distributions $p_{\mathbf{s}}(\bm{s})$ and $p_{\pmb{\uptheta}}(\pmb{\theta})$ such that $\mathbb{E}[\mathbf{s}^\ast\mathbf{s}]\leq m$, where $\tilde{\mathcal{R}}(p_{\mathbf{s}},p_{\pmb{\uptheta}},\bm{H}_{ri},\bm{h}_d)$ is the set of rate pairs $\rb{R_1(\bm{H}_{ri},\bm{h}_d),R_2(\bm{H}_{ri},\bm{h}_d)}$ such that inequalities
R_1(H_ri,h_d)   ≤    1/mI(𝐬;𝐲|),

R_2(H_ri,h_d)   ≤    1/mI(;𝐲|𝐬),

and R_1(H_ri,h_d)+R_2(H_ri,h_d)   ≤   1/mI(𝐬,;𝐲)
hold. Since inputs $\mathbf{s}$ and $\pmb{\uptheta}$ are selected from finite sets, the mutual information $I(\mathbf{s};\mathbf{y}|\pmb{\uptheta})$ in (<ref>) can be written as (see, e.g., <cit.>)
I(𝐬;𝐲|)   =   -NMlog_2(e)
with $\mathbf{z}\sim\mathcal{CN}(\bm{0},\bm{I}_{Nm})$ and where we have defined the scalar

By applying the conditional CGF definition in (<ref>) to (<ref>), we get

Similarly, we also have
I(;𝐲|𝐬)   =    -Nmlog_2(e)-κ(u_2|𝐬_1,_1,𝐳),

I(𝐬,;𝐲)   =   -Nmlog_2(e)-κ(u_3|𝐬_1,_1,𝐳).

Therefore, the region $\tilde{\mathcal{R}}(p_{\mathbf{s}},p_{\pmb{\uptheta}},\bm{H}_{ri},\bm{h}_d)$ in (<ref>) is identical to the region $\mathcal{R}(p_{\mathbf{s}},p_{\pmb{\uptheta}},\bm{H}_{ri},\bm{h}_d)$ in (<ref>).

§.§ Proof of Corollary <ref>

The inclusion $\mathcal{C}(\bm{H}_{ri},\bm{h}_d)\subseteq \overline{\mathcal C}$ is trivial since, for all input distributions $p_{\mathbf{s}}(\bm{s})$ and $p_{\pmb{\uptheta}}(\pmb{\theta})$ with $\bm{s}\in\mathcal S^{m\times 1}$ and $\pmb{\theta}\in\mathcal A^{K\times 1}$ we have $H(\mathbf{s})\leq m\log_2(S)$ and $H(\pmb{\uptheta})\leq K\log_2(A)$. In addition, in the high-power regime, we have the limits
I(𝐬;𝐲|)      H(𝐬)≤mlog_2(S),

I(;𝐲|𝐬)      H()≤Klog_2(A),
where equality is achieved for a uniform distributions $p_{\mathbf{s}}(\bm{s})$ and $p_{\pmb{\uptheta}}(\pmb{\theta})$.
Next, note that the noiseless received signal $\mathbf{y}(t)-\mathbf{z}(t)$ in (<ref>) takes values from a discrete set. Furthermore, since channel matrix $\mathbf{H}_{ri}$ and channel vector $\mathbf{h}_d$ are drawn from a continuous distribution, almost surely, for all $t\in[n/m]$, there exist unique inputs $\hat{\bm{s}}(t)\in\mathcal{S}^{m\times 1}$ and $\hat{\pmb{\theta}}(t)\in\mathcal{A}^{K\times 1}$ such that (see, e.g., <cit.>)

Therefore, for all input distributions $p_{\mathbf{s}}(\bm{s})$ and $p_{\pmb{\uptheta}}(\pmb{\theta})$, the transmitted signal $\mathbf{s}(t)$ and reflection pattern $\pmb{\uptheta}(t)$ can be correctly jointly decoded in the high-power regime, i.e., we have the limit

Let $\rb{R_1^u(\bm{H}_{ri},\bm{h}_d),R_2^u(\bm{H}_{ri},\bm{h}_d)}\in \mathcal{C}(\bm{H}_{ri},\bm{h}_d)$ be the rate pair achieved using uniform distributions $p_{\mathbf{s}}(\bm{s})$ and $p_{\pmb{\uptheta}}(\pmb{\theta})$.
It hence follows from the region in (<ref>) and limits (<ref>) and (<ref>) that, almost surely, we have the limits
lim_P→∞R_1^u(H_ri,h_d)   =   log_2(S),

lim_P→∞R_2^u(H_ri,h_d)   =   K/mlog_2(A).

§.§ Proof of Proposition <ref>

For input distributions $p_{\mathrm{s}}(s)$ and $p_{\pmb{\uptheta}}(\pmb{\theta})$, let functions $\tilde{R}_\ell(P,\bm{h}_{ri},h_d)$, $\ell\in\{1,2,3\}$, be defined as

where $\kappa(\mathrm{u}_\ell|\mathrm{s}_1,\pmb{\uptheta}_1,\mathrm{z})$ are the conditional CGFs in <ref> for the special case in which $N=m=1$. By calculating the derivative of $\tilde{R}_\ell(P,\bm{h}_{ri},h_d)$ with respect to the power $P$ and taking the limit $P\rightarrow0$, we get

where random variables $\underline{\mathrm{u}}_\ell$ are defined in (<ref>). Therefore, it follows from <ref> that the normalized rate pairs $\rb{r_1(\bm{h}_{ri},h_d),r_2(\bm{h}_{ri},h_d)}$ satisfy
r_ℓ(h_ri,h_d)   =   lim_P→0R_ℓ(h_ri,h_d)/P

≤   lim_P→0R̃_ℓ(P,h_ri,h_d)/P

=   lim_P→0∂R̃_ℓ(P,h_ri,h_d)/∂P

=   𝔼[u_ℓ]/ln(2),  ℓ∈{1,2},

and similarly we have

§.§ Proof of Corollary <ref>

Since $\mathrm{s}_1$, $\mathrm{s}_2$, and $\pmb{\uptheta}_1$ in <ref> are all independent, we have
𝔼[u_1]   =   𝔼|h_ri^⊺e^j_1+h_ds_1-s_2|^2

=   𝔼|h_ri^⊺e^j_1+h_d|^2𝔼|s_1-s_2|^2

≤    2|h_ri^⊺e^j+h_d|^2.

Similarly, we have the upper bounds
𝔼[u_2]   ≤   2‖h_ri‖^2,

𝔼[u_3]   ≤   2|h_ri^⊺e^j+h_d|^2.

Equality in (<ref>) and (<ref>) is achieved for fixed RIS reflection pattern $\pmb{\uptheta}=\tilde{\pmb{\theta}}$ with probability one and uniform input distribution $p_{\mathrm{s}}(s)=1/S$. Furthermore, since the upper bounds in (<ref>) and (<ref>) are equal, Encoder 1 can achieve the maximum normalized rate if and only if $\pmb{\uptheta}=\tilde{\pmb{\theta}}$ with probability one.
In contrast, equality in (<ref>) is achieved for uniform phase-shift distribution $p_{\pmb{\uptheta}}(\pmb{\theta})=1/A^K$ and any input distribution $p_{\mathrm{s}}(s)$ for which $\mathbb E[|\mathrm{s}|^2]=1$.
That is, Encoder 2 can achieve the maximum normalized rate, while Encoder 1 transmits at a positive normalized rate.