# LeftoverLocals: Listening to LLM Responses Through Leaked GPU Local Memory

Tyler Sorensen  Trail of Bits
University of California, Santa CruzSanta CruzCaliforniaUSA and  Heidy Khlaaf
Trail of BitsNew York CityNew YorkUSA

###### Abstract.

This paper describes LeftoverLocals: a vulnerability that allows data recovery
from GPU memory created by another process on Apple, Qualcomm, and AMD GPUs.
LeftoverLocals impacts the security posture of GPU applications, with
particular significance to LLMs and ML models that run on impacted GPUs. By
recovering local memory – an optimized GPU memory region – we built a PoC
where an attacker can listen into another user’s interactive LLM session
(e.g., llama.cpp) across process or container boundaries.

††copyright: none

##  1\. Introduction

This paper is essentially a direct port (by the authors) from the official
Trail of bits blog post, which can be found here:
https://blog.trailofbits.com/2024/01/16/leftoverlocals-listening-to-llm-
responses-through-leaked-gpu-local-memory/. This document was created because
an archived paper may be better suited for distribution and citation in some
cases.

Given the rise of ML applications, especially for privacy sensitive
application domains, it is imperative to rigorously examine security
properties throughout the ML stack. This includes the computational work horse
of ML applications, the GPU. While NVIDIA GPUs currently appear to dominate
the ML market, most other hardware vendors create their own GPUs, which are
becoming more widely used for ML applications, especially for local
computation. This work discusses a data leak vulnerability found on many of
these GPUs, specifically, from AMD, Apple, Qualcomm and Imagination. This
leak, dubbed LeftoverLocals, was disclosed earlier this year
(CVE-2023-4969111https://kb.cert.org/vuls/id/446598) and allows a co-resident
process to read leftover data in an optimized GPU memory region called local
memory.

Figure 1. An example of an LLM response that an attacker was able to
reconstruct utilizing LeftoverLocals. The victim terminal is on the left
(black background), and the attacker terminal is on the right (white
backgroun). We can see that the attacker is able to reconstruct the response
with relatively high fidelity. We believe that the listener could be more
finally tuned to be even more accurate. Details proof-of-concept attack
application, along with the system, can be found in Sec. 5.

LeftoverLocals can leak approximately 5.5 MB per GPU invocation on an AMD
Radeon RX 7900 XT which, when running a 7B model on llama.cpp, adds up to 181
MB for each LLM query. This is enough information to reconstruct the LLM
response with high precision, an example is shown in Fig. 1. The vulnerability
highlights that many parts of the ML development stack have unknown security
risks and have not been rigorously reviewed by security experts.

Trail of Bits worked with CERT Coordination Center on a large coordinated
disclosure effort involving all major GPU vendors, including: NVIDIA, Apple,
AMD, Arm, Intel, Qualcomm, and Imagination. As of writing (Feb. 1, 2024), the
status of the impacted vendors, Apple, AMD, and Qualcomm are as follows:

  * •

Apple: Despite multiple efforts to establish contact through CERT/CC, we only
received a response from Apple on January 13, 2024. We re-tested the
vulnerability on January 10 where it appears that some devices have been
patched, i.e., Apple iPad Air 3rd G (A12). However, the issue still appears to
be present on the Apple MacBook Air (M2). Furthermore, the recently released
Apple iPhone 15 does not appear to be impacted as previous versions have been.
Apple has confirmed that the A17 and M3 series processors contain fixes, but
we have not been notified of the specific patches deployed across their
devices.

  * •

AMD: We have confirmed with AMD that their devices remain impacted, although
they continue to investigate potential mitigation plans. Their statement on
the issue can be read here222https://www.amd.com/en/resources/product-
security/bulletin/amd-sb-6010.html.

  * •

Qualcomm: We received notice that there is a patch to Qualcomm firmware
v2.07333https://lore.kernel.org/linux-
<EMAIL_ADDRESS>that addresses
LeftoverLocals for some devices. However, there may still be other devices
impacted at this time. A Qualcomm representative has provided the following
comment: “Developing technologies that endeavor to support robust security and
privacy is a priority for Qualcomm Technologies. We commend Dr. Tyler Sorensen
and Dr. Heidy Khlaaf from the AI/ML Assurance group at Trail of Bits for using
coordinated disclosure practices and are in the process of providing security
updates to our customers. We encourage end users to apply security updates as
they become available from their device makers.”

  * •

Imagination: Despite not observing LeftoverLocals ourselves across the
Imagination GPUs that we tested (see Sec. 6), Google has confirmed that some
Imagination GPUs are indeed impacted. Imagination released a fix in their
latest DDK release, 23.3, made available to customers in December
2023444https://www.imaginationtech.com/gpu-driver-vulnerabilities/.

A list of tested and impacted devices can be found in Sec. 6 Other vendors
have provided us the following details:

  * •

NVIDIA: confirmed that their devices are not currently impacted. One reason
for this could be that researchers have explored various memory leaks on
NVIDIA GPUs previously, e.g., in (Pietro et al., 2016), and thus, they are
aware of these types of issues.

  * •

ARM: also confirmed that their devices are not currently impacted.

We did not hear a response from the remaining vendor, Intel, we however, we
tested two GPUs from them and did not observe that they were impacted (see
Sec. 6).

##  2\. Exploit Brief

GPUs were initially developed to accelerate graphics computations. In this
domain, performance is critical, and previously uncovered security issues have
generally not had any significant consequences on applications. Historically,
this entailed that GPU hardware and software stacks iterated rapidly, with
frequent major architecture and programming model changes. This has led to
complex system stacks and vague specifications. For example, while CPU ISAs
have volumes of documentation, NVIDIA simply provides a few short
tables555https://docs.nvidia.com/cuda/cuda-binary-
utilities/index.html#instruction-set-reference. This type of vague
specification has led to alarming issues, both previously (Alglave et al.,
2015) and currently, as LeftoverLocals exemplifies.

###  2.1. Exploitation requirements

This is a co-resident exploit, meaning that a threat actor’s avenue of attack
could be implemented as another application, app, or user on a shared machine.
The attacker only requires the ability to run GPU compute applications, e.g.,
through OpenCL, Vulkan, or Metal. These frameworks are well-supported and
typically do not require escalated privileges. Using these, the attacker can
read data that the victim has left in the GPU local memory simply by writing a
GPU kernel that dumps uninitialized local memory. These attack programs, as
our code demonstrates, can be less than 10 lines of code. Implementing these
attacks is thus not difficult and is accessible to amateur programmers (at
least in obtaining stolen data). We note that it appears that browser GPU
frameworks (e.g., WebGPU) are not currently impacted, as they insert dynamic
memory checks into GPU kernels.

Unless the user inspects the application’s low-level GPU source-code, it is
not possible for them to uncover if their application is utilizing GPU local
memory; this matter is further complicated as the GPU code is often hidden
deep in library calls, at low levels of deep software stacks (e.g., for ML).
Overall, there are very limited ways to observe that an attacker is currently
stealing data, or has stolen data. This attack hinges on the attacker reading
uninitialized memory on the GPU, and while this is technically undefined
behavior, it is not currently checked dynamically, or logged. Any additional
defenses would be quite invasive, e.g., performing code analysis on GPU
kernels to check for undefined behavior. We have released a PoC that
vulnerability shows how this vulnerability can be
exploited666https://github.com/trailofbits/LeftoverLocalsRelease, and Sec. 5
describes how it works.

###  2.2. Mitigations

Given the lack of comprehensive patches across impacted GPU vendors,
LeftoverLocals can be defended by modifying the source code of all GPU kernels
that use local memory. Before the kernel ends, the GPU threads should clear
memory (e.g., store 0s) to any local memory memory locations that were used in
the kernel. Additionally, the users should ensure the compiler doesn’t remove
these memory-clearing instructions away (e.g., by annotating their local
memory as volatile), as the compiler may detect that the cleared memory is not
used later in the kernel. This is difficult to verify because GPU binaries are
typically not stored explicitly, and there are very few GPU binary analysis
tools. Because of reasons like this, we note that this mitigation may be
difficult for many users, and we discuss this further in Sec. 6.

##  3\. Background: How GPUs work

Figure 2. A simplified view of the GPU architecture: processing elements are
partitioned into compute unites. All processing elements have access to global
memory (often located in VRAM for discrete GPUs), while only processing
elements in the same compute unit share the same local memory.

GPUs are massively parallel, throughput-oriented co-processors. While
originally designed to accelerate graphics workloads, their design, which
balances flexible programming and high computational throughput, has been
highly effective in a variety of applications. Perhaps the most impactful
current application domain is machine learning, where GPUs are the
computational workhorse and are used to achieve nearly all major results in
this area.

GPUs are not only in large servers; they are in our phones, our tablets, and
our laptops. These GPUs come from a variety of vendors, with almost all major
hardware vendors (Apple, AMD, Arm, Qualcomm, Intel, and Imagination) producing
their own GPU architecture. These GPUs are increasingly used for ML tasks,
especially because doing ML locally can preserve users’ privacy, achieve lower
latency, and reduce computational burdens on service providers.

#### GPU architecture

GPU architecture has a parallel, hierarchical structure, shown in Fig. 2. At
the top level, a GPU is made up of Compute Units (sometimes called Streaming
Multiprocessors in NVIDIA literature). Large, discrete GPUs contain many
compute units, and smaller, mobile GPUs have fewer. For example, the large AMD
Radeon RX 7900 XT discrete GPU has 84 compute units, while the mobile Qualcomm
Adreno 740 GPU has 8. All compute units have access to global memory. On
discrete GPUs, global memory is implemented using VRAM; on integrated GPUs,
global memory simply uses the CPU’s main memory.

Compute units encapsulate both compute and memory components. Compute units
contain an array of processing elements; these simple cores are the
fundamental units of computation and execute a stream of GPU instructions. In
terms of memory, compute units often contain a cache for global memory, but
they also contain a special region of memory called local memory. This is an
optimized memory region that is shared only across processing elements in the
same compute unit. This memory can be accessed with significantly less latency
than global memory, but also has much smaller capacity. Different GPUs have
varying amounts of local memory per compute unit, typically ranging from 16KB
to 64KB. For example, the AMD Radeon RX 7900 XT GPU has 84 compute units and a
local memory size of 64KB; thus, the total amount of local memory on the GPU
is 5MB. Local memory is a software-managed cache: the program executing on the
processing elements is responsible for loading values into local memory (e.g.,
values that will be repeatedly used from global memory).

#### GPU execution model

A GPU program, called a (GPU) kernel, is written in a shader language. Common
examples are SPIR-V (Vulkan), OpenCL C, (OpenCL), and Metal Shading Language
(Metal). These kernels specify a single entry point function, called the
kernel function, which is executed by many invocations (i.e., GPU threads).
Invocations have unique built-in identifiers (such as a global ID), which can
be used to index a unique data element in a data-parallel program. Invocations
are further partitioned into workgroups. Each workgroup is mapped to a compute
unit (although many workgroups may execute on the same compute unit, depending
on resource requirements). All invocations have access to the same global
memory, but only invocations in the same workgroup will share the same local
memory.

Applications that use the GPU often launch many short-running kernels. These
kernels often correspond to basic operations, such as matrix multiplication or
convolution. Kernels can then be executed in sequence; for example, each layer
in a deep neural network will be a kernel execution. Local memory is
statically allocated at each kernel launch and is not specified to persist
across kernel calls.

Platforms generally do not time-multiplex different GPU kernels. That is, if
multiple kernels are launched simultaneously (e.g., by different users), the
GPU will execute one kernel to competition before the next kernel starts.
Because GPU kernels are typically short running, sharing GPU resources at
kernel boundaries saves expensive preemption overhead while also maintaining
acceptable latency in practice.

#### Terminology

Because this blog post focuses on portable GPU computing, it uses OpenCL GPU
terminology. For readers more familiar with GPU terminology from a different
framework (e.g., CUDA or Metal), we provide a translation table in Tab. 1.

Table 1. A translation table between terminology used in different GPU programming frameworks. This work uses OpenCL terminology, but readers may by more familiar with CUDA or Metal terminology. CUDA |  OpenCL |  Metal
---|---|---
thread |  work-item |  thread
thread block |  workgroup |  threadgroup
shared memory |  local memory |  threadgroup memory

##  4\. The Vulnerability: LeftoverLocals

In this section we describe the vulnerability, named LeftoverLocals, in more
detail. Section 6 details our testing campaign across a wide variety of GPU
devices, which found that GPUs from AMD, Apple, and Qualcomm are vulnerable to
LeftoverLocals. We also note that while GPU memory leaks are not new, e.g.,
see (Pietro et al., 2016), LeftoverLocals has demonstrated both deeper impact
and wider breadth than previously discovered vulnerabilities.

At a high level, we found that several GPU frameworks do not sufficiently
isolate memory in the same way that it is traditionally expected in CPU-based
frameworks. We have observed that on impacted GPUs, it is possible for one
kernel—potentially from another user that is co-resident on the same
machine—to observe values in local memory that were written by another kernel.
Thus, an attacker who has access to a shared GPU through its programmable
interface (e.g., OpenCL) can steal memory from other users and processes,
violating traditional process isolation properties. This data leaking can have
severe security consequences, especially given the rise of ML systems, where
local memory is used to store model inputs, outputs, and weights.

Previous academic work showed that NVIDIA GPUs leaked memory across processes
through a variety of memory regions, including local memory and across
virtualized environments (Pietro et al., 2016; Maurice et al., 2014). However,
they examined only GPUs from NVIDIA (and we speculate that the results from
this paper may be part of the reason why we didn’t observe LocalLeftovers on
NVIDIA GPUs). They also did not discuss the impact on widely deployed use-
cases, such as ML. Other works have shown how GPUs leak graphics data, and
that a co-resident attacker can reconstruct partial visual information from
another process (Lee et al., 2014; Zhou et al., 2016; Wang et al., 2024). A
good survey of existing GPU vulnerabilities can be found in (Mittal et al.,
2018). Despite these prior works, LeftoverLocals shows that many GPUs remain
vulnerable to local memory leaks and that this vulnerability can be exploited
in co-resident attacks on important ML applications.

Overall, this vulnerability can be illustrated using two simple programs: a
listener and a writer, where the writer stores canary values in local memory,
while a listener reads uninitialized local memory to check for the canary
values. The listener repeatedly launches a GPU kernel that reads from
uninitialized local memory. The writer repeatedly launches a GPU kernel that
writes canary values to local memory. Below, we demonstrate how each of these
operations is carried out.

#### The listener

The listener launches a GPU kernel that reads from uninitialized local memory
and stores the result in a persistent main memory region (i.e., global
memory). This can be accomplished with the OpenCL kernel shown in Fig. 3.

⬇

1__kernel void listener(__global volatile

2 int *dump)

3{

4

5 local volatile int lm[L_SIZE];

6

7 for (int i = get_local_id(0);

8 i < L_SIZE;

9 i += get_local_size(0))

10 {

11

12 dump[L_SIZE*get_group_id(0)+i]=lm[i];

13 }

14}

Figure 3. An OpenCL kernel showing how to implement a LeftoverLocals listener.
Essentially the kernel dumps uninitialized local memory (from the lm array)
into a persistent (global) memory region (in dump) so that it can be examined
later by the host. The OpenCL builtin ids (e.g., get_local_id(0)), allow the
entire local memory dump to be done efficiently in parallel.

The keyword __kernel denotes that this is the GPU kernel function. We pass a
global memory array dump to the function. Whatever the kernel writes to this
array can be read later by the CPU. We statically declare a local memory array
lm with a predefined size L_SIZE (which we set to be the max size of local
memory for each GPU we test). This program technically contains undefined
behavior, as it reads from uninitialized local memory. Because of this, we use
the volatile qualifier to suppress aggressive compiler optimizations that
might optimize away the memory accesses. In fact, the code in the github we
provide contains a few more code patterns included to further stop the
compiler from optimizing away our memory dump. This process is guided by
trial-and-error, rather than a systematic approach.

For each loop iteration, the work-item (thread) reads from a location in local
memory, and then stores that value to a unique location in the global dump
array. The only tricky part of this code is the indexing, because local memory
is disjointed across workgroups, so workgroup local IDs need to be mapped to a
unique global ID in dump. The process utilizes OpenCL built-in identifiers to
achieve this, which are documented in the OpenCL C
specification777https://registry.khronos.org/OpenCL/specs/3.0-unified/html/OpenCL_C.html#work-
item-functions. At the end of the kernel, dump contains every value that was
stored in local memory when the listener kernel started executing. Because
dump is in the global memory region, it can be copied back to CPU memory and
examined to check for canary values.

#### The writer

On the other hand, the writer launches a kernel that writes a canary value to
local memory (for example, this work uses the value 123). We show an example
of the OpenCL kernel code in Fig. 4.

⬇

1__kernel void writer(__global volatile

2 int *canary)

3{

4 local volatile int lm[L_SIZE];

5 for (uint i = get_local_id(0);

6 i < L_SIZE;

7 i += get_local_size(0))

8 {

9 lm[i] = canary[i];

10 }

11}

Figure 4. An OpenCL kernel showing how to implement a LeftoverLocals writer.
Essentially the kernel writes a canary value to all of local memory (in the lm
array) so that a listener can later check to see if it observes canary values.
Similar to the listener, the OpenCL builtin ids (e.g., get_local_id(0)), allow
the writer to write to the entirety of local memory efficiently in parallel.

This code is very similar to the listener, except that rather than dumping
local memory, it writes to local memory. In this case, the code is writing a
value from a memory location called canary. We use an extra location so that
the compiler does not optimize away the memory write (as it is prone to do
with constant values). At the end of the kernel, the writer has filled all
available local memory with the canary value.

#### The Host

The CPU host programs for both the listener and the writer launch their
respective kernels repeatedly. In the case of the listener, at each iteration,
the CPU analyzes the values observed in the local memory and checks for the
canary value. On a server, these two programs can be run by different users or
in different Docker containers. On a mobile device, these routines can be run
in different apps. The apps can be swapped in and out of focus to alternate
reading and writing. If the listener can reliably read the canary values, then
we say that the platform is vulnerable to LeftoverLocals.

Figure 5 shows a series of images on how the listener and the writer interact,
and how the listener may observe values from the writer if local memory is not
cleared.

(a) Initial state of the GPU, where the writer process and the listener
process are co-resident, sharing one GPU device

(b) The writer process launches it’s writer kernel, which becomes the sole
kernel that is executing on the GPU

(c) The writer kernel writes a canary value, in this case 123 to all the local
memory locations across all the compute units

(d) The writer kernel finishes execution. In a system that is impacted by
LeftoverLocals, the canary values stay in local memory

(e) The listener now gets a turn to execute its kernel, sololy occupying the
GPU

(f) The listener kernel copies all local memory on the GPU into global memory

(g) The listener kernel finishes. At this point, global memory (owned by the
listener process) contains a dump of local memory

(h) The listener can now copy global memory back to the CPU and check for
leaked data, e.g., canary values

Figure 5. A series of images illustrating how the listener and the writer
interact, and how they can test for the LocalLeftover vulnerability.

##  5\. Listening to LLM Responses

Figure 6. A flow chart outlining the attack where a malicious listener (the
attacker) steals the output of another process (the victim) executing an
interactive LLM application. These steps are exclusively followed by the
attacker.

In this section, we provide an overview of how LeftoverLocals can be exploited
by a malicious actor (an attacker) to listen to another user’s (the victim)
LLM responses on a multi-tenant GPU machine, followed by a detailed
description of our proof-of-concept (PoC) attack.

At a high level, both actors are executed as co-resident processes. The attack
process implements the listener described above, with the additional steps of
comparing the stolen values to various fingerprints. The victim process is
unknowingly the writer, where instead of canary values, the values being
written are sensitive components of an interactive LLM chat session. The
attack ultimately follows two steps, outlined in Fig. 6:

  1. (1)

The attack process fingerprints the model that the victim process is using by
repeatedly dumping (i.e., stealing) the leftover local memory, which, in this
scenario, consists of sensitive components of linear algebra operations used
by the victim in the LLM model architecture. This requires that the victim is
using an open-source model that the attacker can obtain.

  2. (2)

The attacker then repeatedly dumps (i.e., steals) the local memory again,
specifically searching for the last, i.e., output, layer in the LLM’s DNN,
which can be identified using weights or memory layout patterns from the
earlier fingerprinting.

Note that the output layer is a matrix-vector multiplication with two inputs:
the model weights, and the layer input—in other words, the values derived from
the user input that propagated through the earlier levels of the deep neural
network (DNN). Given that the model weights of the output layer are too large
to comprehensively steal, an attacker can inspect available open-source models
to fully obtain the weights through the exposed model fingerprint. We found
that the second input to the last layer (i.e., the layer input) is
subsequently small enough to fit into local memory. Thus, the entire layer
input can be stolen, and the attacker can reproduce the final layer
computation to uncover the final result of the DNN.

We note that this is a fairly straightforward attack, and with further
creativity and ingenuity, a threat actor may be able to construct further
complex and sophisticated malicious scenarios that may compromise ML
applications in more severe ways. Below we provide a detailed description of
the PoC, and the configuration and testing carried out on various GPU
platforms to uncover their susceptibility to LeftoverLocals.

#### Our configuration

We outline our configuration in Tab. 2. Our attack builds on the llama.cpp LLM
due to its simplicity and cross-vendor GPU acceleration support. In our
example we use a large discrete GPU that we found to be susceptible to
LeftoverLocals: the AMD Radeon RX 7900 XT. We configure llama.cpp to use
OpenCL for GPU acceleration, which uses the CLBLAST linear algebra
library888https://github.com/CNugteren/CLBlast. We use the
wizardLM-7B.ggmlv3.q5_0.bin model, which is open source and can be obtained
from Hugging
Face999https://huggingface.co/TheBloke/wizardLM-7B-GGML/tree/main. This model
was selected due to its reasonable size, which enabled rapid prototyping and
analysis; however, this attack is transferable to many different models. In
our threat model, we assume that the victim is using the LLM in an interactive
chat session.

Table 2. The details of our PoC: code base and system LLM code base | llama.cpp
---|---
Model | wizardLM-7B.ggmlv3.q5_0.bin
Layers | 33
GPU backend | OpenCL
GPU BLAS Library | CLBLAST
GPU | AMD Radeon RX 7900 XT
Total local memory | 5MB

#### Modification

The attack requires an optimized GPU implementation of matrix-vector
multiplication. We found that the current matrix-vector multiplication in
llama.cpp (which does not call into CLBLAST) is not implemented in an
optimized idiomatic way. It stores partial dot product results in local memory
and then combines them at the end. While there is a more complex approach
using linear algebra to achieve our same results, for the simplicity of our
PoC and demonstration, we replace the llama.cpp matrix-vector multiplication
with our own that is more idiomatic (following best GPU programming
programming practices).

#### Step 1—Fingerprinting the model

An attacker can fingerprint a model if it can listen to several inference
queries from the victim. In our configuration, the GPU contains roughly 5MB of
local memory. The model has roughly 33 layers, each of them consisting of a
matrix multiplication operation. Matrix multiplication is often optimized on
GPUs by using tiling: an approach that subdivides the matrices into small
matrices, performs the multiplication, and then combines the
results101010https://cnugteren.github.io/tutorial/pages/page4.html. In many
optimized libraries, including CLBLAST, local memory is used to cache the
smaller matrices. Thus, for every layer, the attacker can steal 2.5MB of
weights, and 2.5MB of the inputs. While this is a significant amount of data,
we note that it is not enough to reconstruct the entire computation. Many of
these layers have weights and inputs that are 100s of MB large.

However, for a whole inference computation (33 layers), the attacker can steal
around 80MB of the weights, which is sufficient to fingerprint the model
(assuming the user is using an open-source model, such as one that can be
found on Hugging Face). Given this, we assume that it is a straightforward
task to fingerprint the model, and thus for the attacker to obtain the full
model being used by the victim.

#### Step 2—Listening to the LLM output

The attacker can then turn their attention to the output layer of the DNN. In
our configuration, we found that the output layer is a matrix-vector
multiplication, rather than a matrix-matrix multiplication. The weights matrix
is large ( 128MB), but the input vector is quite small ( 4KB). However, given
that the attacker has fingerprinted the model in step 1, the attacker does not
need to comprehensively steal the weights as they are available from the
fingerprinted model.

Matrix-vector multiplication has a different GPU implementation than matrix-
matrix multiplication. In the case where the input vector fits in local
memory, the most performant implementation is often to cache the input vector
in local memory, as it is used repeatedly (i.e., for repeated dot products).
Because the input vector is stored entirely in local memory, the attacker can
steal this entire vector. In determining whether the attacker has found local
memory from the output layer, we discovered that the attacker could simply
look for 4KB of floating point values with zeros on either side. In our
testing, this unique fingerprint was associated with the output layer nearly
every single time. For different models and different GPUs, this fingerprint
will likely have to be recalibrated.

#### Putting it together

With an attacker in possession of both the weights and the input vector to the
final layer of the DNN, they can perform the matrix-vector multiplication and
obtain the resulting token returned by the inference. This allows the attacker
to reproduce the output of the victim’s LLM chat session with high fidelity,
as demonstrated in the introduction. In practice, we tuned the attacker to
dump the local memory very efficiently (that is, by using only a small number
of work-items and allocating only a small amount of memory). This allows the
attacker to listen to long chat queries with only a small number of noticeable
output artifacts. Some of the artifacts observed include:

  * •

Duplicate tokens: This occurs when the attacker steals the same output layer
twice due to circumstances such as the attacker process being scheduled twice
in a row, thus the LLM was not scheduled to compute its next token.

  * •

Missing tokens: This occurs when the attacker kernel isn’t scheduled at the
right time, i.e., immediately after the output layer computation kernel.

  * •

Incorrect tokens This occurs due to the attacker mis-identifying a stolen set
of data to be the last layer. In this case, it will print a junk token.

  * •

Similar tokens These tokens are “close” to the original output, even if they
are not exact. That is, the attacker may be unable to steal the exact token
embedding at the target layer. This results in a corrupted token embedding
which, when decoded, is semantically similar (in the word2vec sense) to the
original token. As an example, in Fig. 1, the attacker extracts the incorrect
word “Facebook”, which is semantically similar to other Named Entities tokens
(like “Google”, and “Amazon”) in the generated text.

Despite these discrepant artifacts, the stolen text is more than sufficient to
uncover the LLM response. Additionally, the attacker can be further tuned by,
for example, having multiple threads launch the listener kernel or by having a
more precise fingerprint of the last layer.

##  6\. Testing GPU platforms for LeftoverLocals

Given the diversity of the GPU devices, we wrote several applications to test
for LeftoverLocals, provided in a variety of GPU programming frameworks:

  * •

Vulkan Command Line: A command line application using Vulkan. The kernel is
written in OpenCL and compiled to SPIR-V using
clspv111111https://github.com/google/clspv. It uses a simple Vulkan wrapper
called EasyVK121212https://github.com/ucsc-chpl/easyvk.

  * •

OpenCL Command Line: A command line application that uses the OpenCL
framework.

  * •

Apple App: An Apple app that can be deployed on iOS or Mac OS. It targets the
GPU using Apple’s Metal framework.

  * •

Android App: An Android app that uses Vulkan to target mobile GPUs. The code
uses Vulkan’s C API (through EasyVK again) using JNI. The kernels are the same
as in the Vulkan command line app: they are written in OpenCL and compiled to
SPIR-V using clspv.

Using the above programs, we tested 11 devices spanning seven GPU vendors (and
multiple GPU frameworks in some cases). We observed LeftoverLocals on devices
from three of the vendors (Apple, Qualcomm, and AMD). The amount of memory
leaked depends on the size of the GPU. Larger GPUs contain more physical
memory, and thus, leak more data. For the larger GPUs (e.g., an AMD Radeon RX
7900 XT), we found that we can leak over 5MB per kernel. Table 3 shows the
system information in the cases where we were able able to observe
LeftoverLocals (QC refers to Qualcomm):

Table 3. GPUs and systems that we observed to be impacted by LeftoverLocals. QC refers to Qualcomm.  Device (GPU) |  Framework |  OS/Driver/Build system
---|---|---
Apple iPhone 12 Pro (A14) |  Metal |  iOS 16.6, Xcode 14.3.1 (14E300c)
Apple iPad Air 3rd G (A12) |  Metal |  iOS 16.5.1, Xcode 14.3.1 (14E300c)
Apple MacBook Air (M2) |  Metal |  MacOS 13.4.1, Xcode 14.3.1 (14E300c)
AMD Radeon RX 7900 XT |  Vulkan |  Arch Linux, Mesa 23.1.4
AMD Radeon RX 7900 XT |  OpenCL |  Arch Linux, OpenCL 2.1 AMD-APP.dbg (3570.0)
AMD Ryzen 7 5700G (integrated GPU) |  Vulkan |  Arch Linux, Mesa 23.1.4
AMD RX 6700 XT |  Vulkan |  Windows 11 Pro 22H2, AMD Vulkan 2.0.270
HTC 1+ 11 (QC Snapdragon 8 g2) |  Vulkan |  Android 13, Android Studio (2022.3.1)
HTC 1+ 5T (QC Snapdragon 835) |  Vulkan |  Android 13, Android Studio (2022.3.1)
Table 4. GPUs and systems where the listener returned non–zero values, but we could not confirm that they came from a different process. Arm representatives reviewed the results and confirmed that these values did not come from different processes on their devices. Device (GPU) |  Framework |  OS/Driver/Build system
---|---|---
Galaxy Tab A (Arm Mali G78) |  Vulkan |  Android 13, Android Studio (2022.3.1)
Google Pixel 6 (Arm Mali G71) |  Vulkan |  Android 11, Android Studio (2022.3.1)
Google Pixel 7 (Arm Mali G710) |  Vulkan |  Android 13, Android Studio (2022.3.1)
Table 5. GPUs and systems that we did NOT observe to be impacted by LeftoverLocals. However, later representatives from Google notified us that they observed other devices from Imagination (IM) were impacted, and have subsequently been patched, as documented in the paper introduction.  Device (GPU) |  Framework |  OS/Driver/Build system
---|---|---
Nvidia GeForce RTX 4070 |  Vulkan |  Arch Linux, Mesa 23.1.4
Nvidia GeForce RTX 4070 |  OpenCL |  Arch Linux, OpenCL 3.0 CUDA 12.2.128
Intel NUC (NUC10I5FNK) |  Vulkan |  Ubuntu 22.04, Mesa 20.3.2
Intel NUC (NUC10I5FNK) |  OpenCL |  Ubuntu 22.04, OpenCL 3.0 NEO (22.31.23852)
Motorola M.G (IM PowerVR GE8320) |  Vulkan |  Android 11, Android Studio (2022.3.1)

For some devices, specifically those from Arm, we were not able to observe the
canary value from the writer in the listener, but we did observe non-zero
data. Representatives from Arm reviewed our observations and concluded that
although these values are not zero, they are not from a memory leak. We
outline these devices in Tab. 4.

Additionally, we tested some GPUs from NVIDIA, Intel, and Imagination,
outlined in Tab. 5 (IM refers to Immagination). For these devices, we observed
only zeros in local memory, and thus did not observe LeftoverLocals. It is
unclear if all their devices are not impacted. For example, although we did
not observe the issue on our Imagination device, Google notified us that they
were able to observe it on other Imagination devices.

We prepared a YouTube video that demonstrates the different interfaces and
examples of LocalLeftovers—namely the LLM PoC attack, covert communication
channels, and searching for canary values—on a few different platforms using a
few different applications131313https://www.youtube.com/watch?v=g2A7GvbnItg.
All code is provided in our github repository.

#### Vulnerable environments

An attack program must be co-resident on the same machine and must be
“listening” at the same time that the victim is running a sensitive
application on the GPU. This could occur in many scenarios: for example, if
the attack program is co-resident with the victim on a shared cloud computer
with a GPU. On a mobile device, the attack could be implemented in an app or a
library. Listening can be implemented efficiently, and thus can be done
repeatedly and constantly with almost no obvious performance degradation.

Next, we briefly discuss other environments where GPUs are either deployed or
where an attacker might have access to sensitive information. Although it
appears that some current systems (e.g., WebGPU) are not currently impacted,
the ever-growing prevalence of ML and the diversity of modern GPUs mean that
the next iteration of these systems (or other near-future systems) may be
severely compromised by these types of vulnerabilities.

  * •

Cloud providers: Cloud providers (e.g., AWS and Azure) are unlikely to provide
shared GPU instances, especially if users have dedicated access to the GPU
machine. In other cases, GPUs could be shared using very conservative GPU VM
technology (such as NVIDIA’s vGPU141414https://docs.nvidia.com/grid/13.0/grid-
vgpu-user-guide/index.html or AMD’s
MxGPU151515https://www.amd.com/system/files/documents/gpu-consistency-
security-whitepaper.pdf), which physically partitions the GPU and therefore
prevents users from sharing GPU resources (including local memory). Given
this, many current cloud GPU systems may not currently be vulnerable to
LeftoverLocals; however, we do not have conclusive evidence to determine this
given the general lack of visibility into the specification and implementation
of these systems. We note that we have observed LeftoverLocals on multi-user
Linux servers, as well as on desktop (Windows and Mac) systems through
traditional multi-processing. This includes Docker containers on these
systems.

  * •

Mobile applications: In our experiments and explorations in the mobile domain,
we were able to run concurrent GPU processes (from different apps on iOS or
Android) only in very specific instances. That is, we were not able to run a
GPU process (e.g., from a malicious listener app) in the background while
other apps (e.g., the victim) were run in the foreground. As with our analysis
of cloud providers, we were unable to find clear documentation that explicitly
detailed these constraints, and so we cannot definitively claim whether they
are vulnerable. However, as seen our youtube video, LeftoverLocals can be
exploited either when a malicious listener app is run side-by-side with a
victim app, or if the malicious listener app is quickly swapped from the
background into the foreground from a victim app.

  * •

Remote attacks: We preliminarily investigated the possibility of attacks
originating from websites (e.g., those hosted by a remote attacker). To our
knowledge, web applications do not have the low-level features required to
listen to local memory using GPU graphics frameworks, such as WebGL. We note
that the new WebGPU framework does provide low-level capabilities that allow a
webpage to access local memory. Conservatively, WebGPU initializes and
performs dynamic array bounds checking on local memory (and global memory),
which mitigates this vulnerability. However, these checks cause significant
overhead, as documented in online
discussions161616https://github.com/gpuweb/gpuweb/issues/1202. To test this
further, our code repo contains a simple listener in WebGPU. As expected, we
have only observed zeros in local memory, even on devices that are vulnerable
to LeftoverLocals through other frameworks. However, GPU compilers are known
to be fragile (Lidbury et al., 2015), and it is not difficult to imagine
finding a compiler bug that could somehow bypass these checks (especially
using fuzzing techniques). Our position is that LocalLeftovers should be
addressed at a lower level (e.g., the driver).

#### How GPU vendors can resolve this vulnerability

To defend against LocalLeftovers, GPUs should clear their local memory between
kernel calls. While this could cause some performance overhead, our
experiments show that many GPU vendors (e.g., NVIDIA, Intel) currently appear
to provide this functionality. It even appears that some of this functionality
is provided for impacted GPUs. For example, Mesa drivers for AMD GPUs clears
local memory after a compute kernel
launch171717https://github.com/Mesa3D/mesa/blob/957009978ef6d7121fc0d710d03bc20097d4d46b/src/amd/vulkan/radv_shader.c#L709.
However, this approach has a fundamental flaw that makes it vulnerable to
LeftoverLocals: this memory wipe is done with a separate kernel, thus, the GPU
kernel queue may contain a malicious listener between the computation kernel
and the local memory wipe, allowing the listener to steal memory. Instead, the
computation kernel and the local memory wipe need to occur atomically, i.e.,
without allowing any other kernel to be interleaved between them. Otherwise, a
user may attempt to preemptively defend themselves against LeftoverLocals as
described in the next section.

#### Mitigations

Given the lack of comprehensive patches across impacted GPU vendors,
LeftoverLocals can be defended by modifying the source code of all GPU kernels
that use local memory. As we’ve previously noted, before the kernel ends, the
GPU threads can store 0 to any local memory locations that were used in the
kernel. Given that GPU tasks are typically interleaved at the kernel boundary,
this will prevent another user from being able to read leftover values. We
note that this mitigation may be difficult for many users, especially because
GPU code is often buried deep in complex software stacks (e.g., for ML).
Furthermore, the GPU code may be part of a highly optimized library (e.g., ML
linear algebra routines). In these cases, it is very difficult to identify how
local memory is used, and even more difficult to modify the kernel to zero it
out. It may be possible to augment a compiler to add this functionality,
similar to how WebGPU handles GPU memory accesses. These mitigations do have a
performance overhead that should be taken into account. Another blunt
mitigation involves simply avoiding multi-tenant GPU environments.

##  7\. Impact on LLMs and GPU platforms

#### LLM security

Our PoC attack examines only one application: an interactive open-source LLM
session. However, with a little creativity, attackers could likely target many
GPU applications, including those used within privacy-sensitive domains. Our
motivation stems from the recent increased use and support of open-source
models, often accompanied by claims that their “openness” inherently entails
safety and security through transparency. A recent article in
Nature181818https://www.nature.com/articles/d41586-023-03803-y even alleges
that only open-source generative AI models can “safely” revolutionize health
care, a safety-critical domain. Yet, even if open-source models provide the
opportunity to be rigorously audited and assessed (which they have yet to
be191919https://blog.trailofbits.com/2023/11/15/assessing-the-security-
posture-of-a-widely-used-vision-model-yolov7/), their deployment still hinges
on a closed-source stack (i.e., GPUs). And as demonstrated by LeftoverLocals,
open-source LLMs are particularly susceptible to our vulnerability given our
ability to fingerprint these models to obtain remaining weights as needed.
Indeed, we have already observed announcements regarding the deployment of
open-source models in collaboration with impacted GPU vendors, including
Hugging Face’s collaboration with
AMD202020https://twitter.com/AMD/status/1744831880241750112, Lamini’s
deployment on AMD GPUs212121https://www.lamini.ai/blog/lamini-amd-paving-the-
road-to-gpu-rich-enterprise-llms, and the Qualcomm and Meta partnership for
edge devices 222222https://www.qualcomm.com/news/releases/2023/07/qualcomm-
works-with-meta-to-enable-on-device-ai-applications-usi.

Generally, the introduction of ML poses new attack surfaces that traditional
threat models do not account for, and that can lead to implicit and explicit
access to data, model parameters, or resulting outputs, increasing the overall
attack surface of the system. It is crucial to identify and taxonomize novel
classes of failure modes that directly impact ML models, in addition to novel
threats that can compromise the ML Ops pipeline, as we have demonstrated with
LeftoverLocals. We discuss GPU-specific threat implications in the following
section.

#### GPU providers, applications, and vendors

While many platforms are not currently impacted, we emphasize that the GPU
compute landscape is evolving rapidly. As some examples: there are a growing
number of GPU cloud providers have various policies and available
configurations232323https://cloud-gpus.com/; and GPU programming frameworks,
such as Vulkan and Metal, are well-supported on mainstream platforms, and can
be used in apps without requiring extra privileges. While these developments
are exciting, they increase the threat potential of GPU vulnerabilities, as
LeftoverLocals illustrates. As far as we are aware, there is no unified
security specification for how GPUs are required to handle sensitive data, and
no portable test suite to check if systems are vulnerable to simple memory
leaks, like LeftoverLocals. Thus, GPU compute environments should be
rigorously scrutinized when used for processing any type of sensitive data.

As mentioned throughout this paper, while we focus on LLM applications, GPU
local memory is one of the first tools that a GPU developer uses when
optimizing an application. Although these attacks require analyzing the
victim’s GPU kernel code to identify local memory usage, it is likely possible
to find similar application attacks in other GPU compute domains, such as
image processing and scientific computing. It will be increasingly difficult
for users to detect and defend against these attacks since it’s unlikely they
will know if their application is vulnerable to LeftoverLocals; this would
require knowing the details of the exact GPU kernel code, which are often
hidden away in highly optimized linear algebra libraries (e.g., CLBLAST).
Additionally, an overall lack of specification in up-and-coming GPU platforms
makes it difficult to determine whether the compiler or runtime will use
impacted memory regions without the user knowing. For example, Apple GPUs have
a new caching mechanism, called dynamic
caching242424https://www.digitaltrends.com/computing/apple-dynamic-caching-
explained/, which does not have a clear specification regarding if local
memory regions are being used for other purposes.

##  8\. Coordinated disclosure

Since September 2023, we have been working CERT/CC on a large coordinated
disclosure involving all major GPU vendors, including NVIDIA, Apple, AMD, Arm,
Intel, Qualcomm, and Imagination. Trail of Bits provided vendors a total of
125 days to test their products and provide remediations. The coordination
gradually grew to include software stakeholders, including Google, Microsoft,
and others, which allowed us to understand how LocalLeftovers impacts privacy
requirements and impact at different stages in the ML supply chain. Apple did
not respond or engage with us regarding the disclosure until three days before
the embargo ended.

A high-level timeline of the disclosure is provided below:

  * •

September 8, 2023: Trail of Bits submitted report to the CERT/CC

  * •

September 11, 2023: CERT/CC acknowledged the submission of LeftoverLocals and
began the process of vendor outreach and CVE assignment with a preliminary
disclosure date of December 11, 2023

  * •

September 14, 2023: AMD acknowledged the CERT disclosure

  * •

September 15, 2023: Qualcomm acknowledged the CERT disclosure

  * •

September 22, 2023: The case report was shared with the Khronos group, which
produces the OpenCL and Vulkan specifications and some corresponding
infastructure (such as conformance testing)

  * •

September 29, 2023: NVIDIA acknowledged disclosure and confirmed they were not
affected by the vulnerability

  * •

November 22, 2023: ToB extended release of embargo to January 16, 2024 to
accommodate for vendor requests for further time

  * •

January 11, 2024: We received a notice that Qualcomm provided a patch to their
firmware that addresses this issue only for some of their devices.
Additionally, Google noted that ChromeOS Stable 120 and LTS 114 will be
released on January 16 to include AMD and Qualcomm mitigations.

  * •

January 13, 2024: Apple confirmed that the A17 and M3 series processors
contain fixes to the vulnerability.

  * •

January 14, 2024: Google notified us that they observed that that some
Imagination GPUs are impacted.

  * •

January 16, 2024: Embargo lift and public disclosure of LeftoverLocals

##  9\. Moving forward

Now that GPUs are being used in a wide range of applications, including
privacy sensitive applications, we believe that the wider GPU systems
community (vendors, researchers, developers) must work towards hardening the
GPU system stack and corresponding specifications. This should be accomplished
through robust, holistic specifications that describe both GPU programs’
behavior and how GPU devices integrate with the rest of the system stack
(e.g., the OS or hypervisor). Furthermore, these specifications should be
rigorously tested to account for the diversity of GPU systems and safety
requirements of diverse application domains. Looking forward, a wide variety
of new AI chips are being
developed252525https://www.theinformation.com/articles/the-twelve-startups-
battling-for-a-slice-of-nvidias-pie?utm_source=ti_app and will require
rigorous security analysis.

There are positive developments in this direction. For example, AMD’s ROCm
stack262626https://www.amd.com/en/products/software/rocm.html is open, and
thus available for independent rigorous evaluation, and the Khronos Group has
safety critical specification groups272727https://www.khronos.org/syclsc.
Additionally, cross-vendor programming frameworks, such as Vulkan, have been
incredibly useful for writing portable test suites, as opposed to single-
vendor programming frameworks.

While GPU security and privacy guarantees are scattered and scarce, the Vulkan
specification outlines a reasonable definition of security for GPU platforms
to adhere to—a definition that several platforms clearly violate, as
LocalLeftovers
shows:282828https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#fundamentals-
validusage

> …implementations must ensure that […] an application does not affect the
> integrity of the operating system […]. In particular, any guarantees made by
> an operating system about whether memory from one process can be visible to
> another process or not must not be violated by a Vulkan implementation for
> any memory allocation.

With the dust settling, our position is the following: given the wide
diversity of GPUs and their critical importance in enabling machine learning
applications, these devices, and their ecosystems, are in need of (1) a
detailed threat model that considers the various types of data processed on
GPUs and how this data might be compromised; (2) an exploration of the GPU
execution stack to determine where and how GPU security properties should be
specified and implemented; and (3) significant testing and auditing to fortify
GPU ecosystem, which is the computational foundation of machine learning.

For full transparency, we note that Tyler Sorensen has been an invited member
of the Khronos group (sponsored by Google) since 2019, and participates in the
memory model technical specification group.

## Acknowledgements

We thank Dan Guido, Trent Brunson, Max Ammann, Dominik Czarnota, Kelly
Kaoudis, Jay Little, and Adelin Travers for their insightful comments and
feedback on the vulnerability, PoC, and throughout the disclosure process. We
also thank the Khronos Group for discussing technical specification details
with us, and providing an avenue for us to engage with many vendors. We thank
CERT/CC, specifically Vijay Sarvepalli and Ben Koo, for organizing the
coordinated disclosure, especially considering the potential breadth of the
vulnerability. Finally, thank you to everyone who engaged with us on this
issue. This was a large project and we had discussions with many people who
provided valuable insights and perspectives.

## References

  * (1)
  * Alglave et al. (2015) Jade Alglave, Mark Batty, Alastair F. Donaldson, Ganesh Gopalakrishnan, Jeroen Ketema, Daniel Poetzl, Tyler Sorensen, and John Wickerson. 2015.  GPU Concurrency: Weak Behaviours and Programming Assumptions. In _Architectural Support for Programming Languages and Operating Systems_ _(ASPLOS)_. ACM.  https://doi.org/10.1145/2694344.2694391
  * Lee et al. (2014) Sangho Lee, Youngsok Kim, Jangwoo Kim, and Jong Kim. 2014.  Stealing Webpages Rendered on Your Browser by Exploiting GPU Vulnerabilities. In _2014 IEEE Symposium on Security and Privacy_.  https://doi.org/10.1109/SP.2014.9
  * Lidbury et al. (2015) Christopher Lidbury, Andrei Lascu, Nathan Chong, and Alastair F. Donaldson. 2015.  Many-core compiler fuzzing. In _Programming Language Design and Implementation_ _(PLDI ’15)_. ACM.  https://doi.org/10.1145/2737924.2737986
  * Maurice et al. (2014) Clémentine Maurice, Christoph Neumann, Olivier Heen, and Aurélien Francillon. 2014.  Confidentiality Issues on a GPU in a Virtualized Environment. In _Financial Cryptography and Data Security_ , Nicolas Christin and Reihaneh Safavi-Naini (Eds.). Springer Berlin Heidelberg, Berlin, Heidelberg.
  * Mittal et al. (2018) Sparsh Mittal, S. B. Abhinaya, Manish Reddy, and Irfan Ali. 2018.  A Survey of Techniques for Improving Security of GPUs.  _CoRR_ abs/1804.00114 (2018).  arXiv:1804.00114 http://arxiv.org/abs/1804.00114
  * Pietro et al. (2016) Roberto Di Pietro, Flavio Lombardi, and Antonio Villani. 2016.  CUDA Leaks: A Detailed Hack for CUDA and a (Partial) Fix.  _ACM Transactions on Embedded Computing Systems_ 15, 1 (Jan. 2016), 1–25.  https://doi.org/10.1145/2801153
  * Wang et al. (2024) Y. Wang, R. Paccagnella, Z. Gang, W. Vasquez, D. Kohlbrenner, H. Shacham, and C. Fletcher. 2024.  GPU.zip: On the Side-Channel Implications of Hardware-Based Graphical Data Compression. In _2024 IEEE Symposium on Security and Privacy (SP)_. IEEE Computer Society, 84–84.  https://doi.org/10.1109/SP54263.2024.00084
  * Zhou et al. (2016) Zhe Zhou, Wenrui Diao, Xiangyu Liu, Zhou Li, Kehuan Zhang, and Rui Liu. 2016.  Vulnerable GPU Memory Management: Towards Recovering Raw Data from GPU.  _CoRR_ abs/1605.06610 (2016).  arXiv:1605.06610 http://arxiv.org/abs/1605.06610


# Sparse and Constrained Attention for Neural Machine Translation

Chaitanya Malaviya
Language Technologies Institute
Carnegie Mellon University
<EMAIL_ADDRESS>
&Pedro Ferreira
Instituto Superior Técnico
Universidade de Lisboa
Lisbon, Portugal
pedro.miguel.ferreira
@tecnico.ulisboa.pt
&André F. T. Martins
Unbabel
& Instituto de Telecomunicações
Lisbon, Portugal
andre.martins
@unbabel.com Work done during an internship at Unbabel.

###### Abstract

In NMT, words are sometimes dropped from the source or generated repeatedly in
the translation. We explore novel strategies to address the coverage problem
that change only the attention transformation. Our approach allocates
fertilities to source words, used to bound the attention each word can
receive. We experiment with various sparse and constrained attention
transformations and propose a new one, constrained sparsemax, shown to be
differentiable and sparse. Empirical evaluation is provided in three languages
pairs.

##  1 Introduction

Neural machine translation (NMT) emerged in the last few years as a very
successful paradigm (Sutskever et al., 2014; Bahdanau et al., 2014; Gehring et
al., 2017; Vaswani et al., 2017). While NMT is generally more fluent than
previous statistical systems, adequacy is still a major concern Koehn and
Knowles (2017): common mistakes include dropping source words and repeating
words in the generated translation.

Previous work has attempted to mitigate this problem in various ways. Wu et
al. (2016) incorporate coverage and length penalties during beam search—a
simple yet limited solution, since it only affects the scores of translation
hypotheses that are already in the beam. Other approaches involve
architectural changes: providing coverage vectors to track the attention
history Mi et al. (2016); Tu et al. (2016), using gating architectures and
adaptive attention to control the amount of source context provided Tu et al.
(2017a); Li and Zhu (2017), or adding a reconstruction loss Tu et al. (2017b).
Feng et al. (2016) also use the notion of fertility implicitly in their
proposed model. Their “fertility conditioned decoder” uses a coverage vector
and an “extract gate” which are incorporated in the decoding recurrent unit,
increasing the number of parameters.

In this paper, we propose a different solution that does not change the
overall architecture, but only the attention transformation. Namely, we
replace the traditional softmax by other recently proposed transformations
that either promote attention sparsity (Martins and Astudillo, 2016) or upper
bound the amount of attention a word can receive (Martins and Kreutzer, 2017).
The bounds are determined by the fertility values of the source words. While
these transformations have given encouraging results in various NLP problems,
they have never been applied to NMT, to the best of our knowledge.
Furthermore, we combine these two ideas and propose a novel attention
transformation, constrained sparsemax, which produces _both_ sparse and
bounded attention weights, yielding a compact and interpretable set of
alignments. While being in-between soft and hard alignments (Figure 2), the
constrained sparsemax transformation is end-to-end differentiable, hence
amenable for training with gradient backpropagation.

To sum up, our contributions are as follows:111Our software code is available
at the OpenNMT fork www.github.com/Unbabel/OpenNMT-py/tree/dev and the running
scripts at www.github.com/Unbabel/
sparse_constrained_attention.

  * •

We formulate constrained sparsemax and derive efficient linear and sublinear-
time algorithms for running forward and backward propagation. This
transformation has two levels of sparsity: over time steps, and over the
attended words at each step.

  * •

We provide a detailed empirical comparison of various attention
transformations, including softmax (Bahdanau et al., 2014), sparsemax (Martins
and Astudillo, 2016), constrained softmax (Martins and Kreutzer, 2017), and
our newly proposed constrained sparsemax. We provide error analysis including
two new metrics targeted at detecting coverage problems.

##  2 Preliminaries

Our underlying model architecture is a standard attentional encoder-decoder
(Bahdanau et al., 2014). Let $x:=x_{1:J}$ and $y:=y_{1:T}$ denote the source
and target sentences, respectively. We use a Bi-LSTM encoder to represent the
source words as a matrix
$\bm{H}:=[\bm{h}_{1},\ldots,\bm{h}_{J}]\in\mathbb{R}^{2D\times J}$. The
conditional probability of the target sentence is given as

$\textstyle p(y\,\,|\,\,x):=\prod_{t=1}^{T}p(y_{t}\,\,|\,\,y_{1:(t-1)},x),$
(1)

where $p(y_{t}\,\,|\,\,y_{1:(t-1)},x)$ is computed by a softmax output layer
that receives a decoder state $\bm{s}_{t}$ as input. This state is updated by
an auto-regressive LSTM,
$\bm{s}_{t}=\operatorname*{\mathsf{RNN}}(\operatorname*{\mathsf{embed}}(y_{t-1}),\bm{s}_{t-1},\bm{c}_{t})$,
where $\bm{c}_{t}$ is an input context vector. This vector is computed as
$\bm{c}_{t}:=\bm{H}\bm{\alpha}_{t}$, where $\bm{\alpha}_{t}$ is a probability
distribution that represents the attention over the source words, commonly
obtained as

$\bm{\alpha}_{t}=\operatorname*{\mathsf{softmax}}(\bm{z}_{t}),$ (2)

where $\bm{z}_{t}\in\mathbb{R}^{J}$ is a vector of scores. We follow Luong et
al. (2015) and define $z_{t,j}:=\bm{s}_{t-1}^{\top}\bm{W}\bm{h}_{j}$ as a
bilinear transformation of encoder and decoder states, where $\bm{W}$ is a
model parameter.222This is the default implementation in the OpenNMT package.
In preliminary experiments, feedforward attention (Bahdanau et al., 2014) did
not show improvements.

##  3 Sparse and Constrained Attention

In this work, we consider alternatives to Eq. 2. Since the softmax is strictly
positive, it forces all words in the source to receive _some_ probability mass
in the resulting attention distribution, which can be wasteful. Moreover, it
may happen that the decoder attends repeatedly to the same source words across
time steps, causing repetitions in the generated translation, as Tu et al.
(2016) observed.

With this in mind, we replace Eq. 2 by
$\bm{\alpha}_{t}=\rho(\bm{z}_{t},\bm{u}_{t})$, where $\rho$ is a
transformation that may depend both on the scores
$\bm{z}_{t}\in\mathbb{R}^{J}$ and on upper bounds
$\bm{u}_{t}\in\mathbb{R}^{J}$ that limit the amount of attention that each
word can receive. We consider three alternatives to softmax, described next.

Figure 1: Illustration of the different attention transformations for a toy
example with three source words. We show the attention values on the
probability simplex. In the first row we assume scores
$\bm{z}=(1.2,0.8,-0.2)$, and in the second and third rows
$\bm{z}=(0.7,0.9,0.1)$ and $\bm{z}=(-0.2,0.2,0.9)$, respectively. For
constrained softmax/sparsemax, we set unit fertilities to every word; for each
row the upper bounds (represented as green dashed lines) are set as the
difference between these fertilities and the cumulative attention each word
has received. The last row illustrates the cumulative attention for the three
words after all rounds.

#### Sparsemax.

The sparsemax transformation (Martins and Astudillo, 2016) is defined as:

$\displaystyle\operatorname*{\mathsf{sparsemax}}(\bm{z})$ $\displaystyle:=$
$\displaystyle\operatorname*{\mathsf{arg\,min}}_{\bm{\alpha}\in\Delta^{J}}\|\bm{\alpha}-\bm{z}\|^{2},$
(3)

where $\Delta^{J}:=\\{\bm{\alpha}\in\mathbb{R}^{J}\,\,|\,\,\
\bm{\alpha}\geq\mathbf{0},\sum_{j}\alpha_{j}=1\\}$. In words, it is the
Euclidean projection of the scores $\bm{z}$ onto the probability simplex.
These projections tend to hit the boundary of the simplex, yielding a sparse
probability distribution. This allows the decoder to attend only to a few
words in the source, assigning zero probability mass to all other words.
Martins and Astudillo (2016) have shown that the sparsemax can be evaluated in
$O(J)$ time (same asymptotic cost as softmax) and gradient backpropagation
takes sublinear time (faster than softmax), by exploiting the sparsity of the
solution.

#### Constrained softmax.

The constrained softmax transformation was recently proposed by Martins and
Kreutzer (2017) in the context of easy-first sequence tagging, being defined
as follows:

$\displaystyle\operatorname*{\mathsf{csoftmax}}(\bm{z};\bm{u})$
$\displaystyle\\!\\!:=\\!\\!$
$\displaystyle\operatorname*{\mathsf{arg\,min}}_{\bm{\alpha}\in\Delta^{J}}\mathsf{KL}(\bm{\alpha}\|\operatorname*{\mathsf{softmax}}(\bm{z}))$
(4) $\displaystyle\qquad\text{s.t.}\quad\bm{\alpha}\leq\bm{u},$

where $\bm{u}$ is a vector of upper bounds, and $\mathsf{KL}(.\|.)$ is the
Kullback-Leibler divergence. In other words, it returns the distribution
closest to $\operatorname*{\mathsf{softmax}}(\bm{z})$ whose attention
probabilities are bounded by $\bm{u}$. Martins and Kreutzer (2017) have shown
that this transformation can be evaluated in $O(J\log J)$ time and its
gradients backpropagated in $O(J)$ time.

To use this transformation in the attention mechanism, we make use of the idea
of fertility (Brown et al., 1993). Namely, let
$\bm{\beta}_{t-1}:=\sum_{\tau=1}^{t-1}\bm{\alpha}_{\tau}$ denote the
cumulative attention that each source word has received up to time step $t$,
and let $\bm{f}:=(f_{j})_{j=1}^{J}$ be a vector containing fertility upper
bounds for each source word. The attention at step $t$ is computed as

$\bm{\alpha}_{t}=\operatorname*{\mathsf{csoftmax}}(\bm{z}_{t},\bm{f}-\bm{\beta}_{t-1}).$
(5)

Intuitively, each source word $j$ gets a credit of $f_{j}$ units of attention,
which are consumed along the decoding process. If all the credit is exhausted,
it receives zero attention from then on. Unlike the sparsemax transformation,
which places sparse attention over the source words, the constrained softmax
leads to sparsity over time steps.

#### Constrained sparsemax.

In this work, we propose a novel transformation which shares the two
properties above: it provides _both_ sparse and bounded probabilities. It is
defined as:

$\displaystyle\operatorname*{\mathsf{csparsemax}}(\bm{z};\bm{u})$
$\displaystyle\\!\\!:=\\!\\!$
$\displaystyle\operatorname*{\mathsf{arg\,min}}_{\bm{\alpha}\in\Delta^{J}}\|\bm{\alpha}-\bm{z}\|^{2}$
(6) $\displaystyle\qquad\text{s.t.}\quad\bm{\alpha}\leq\bm{u}.$

The following result, whose detailed proof we include as supplementary
material (Appendix A), is key for enabling the use of the constrained
sparsemax transformation in neural networks.

###### Proposition 1

Let $\bm{\alpha}^{\star}=\operatorname*{\mathsf{csparsemax}}(\bm{z};\bm{u})$
be the solution of Eq. 6, and define the sets
$\mathcal{A}=\\{j\in[J]\,\,|\,\,0<\alpha_{j}^{\star}<u_{j}\\}$,
$\mathcal{A}_{L}=\\{j\in[J]\,\,|\,\,\alpha_{j}^{\star}=0\\}$, and
$\mathcal{A}_{R}=\\{j\in[J]\,\,|\,\,\alpha_{j}^{\star}=u_{j}\\}$. Then:

  * •

Forward propagation. $\bm{\alpha}^{\star}$ can be computed in $O(J)$ time with
the algorithm of Pardalos and Kovoor (1990) (Alg. 1 in Appendix A). The
solution takes the form
$\alpha_{j}^{\star}=\max\\{0,\min\\{u_{j},z_{j}-\tau\\}\\}$, where $\tau$ is a
normalization constant.

  * •

Gradient backpropagation. Backpropagation takes _sublinear_ time
$O(|\mathcal{A}|+|\mathcal{A_{R}}|)$. Let $L(\bm{\theta})$ be a loss function,
$\mathrm{d}\bm{\alpha}=\nabla_{\bm{\alpha}}L(\bm{\theta})$ be the output
gradient, and $\mathrm{d}\bm{z}=\nabla_{\bm{z}}L(\bm{\theta})$ and
$\mathrm{d}\bm{u}=\nabla_{\bm{u}}L(\bm{\theta})$ be the input gradients. Then,
we have:

$\displaystyle\mathrm{d}z_{j}$ $\displaystyle=$
$\displaystyle\mathds{1}(j\in\mathcal{A})(\mathrm{d}\alpha_{j}-m)$ (7)
$\displaystyle\mathrm{d}u_{j}$ $\displaystyle=$
$\displaystyle\mathds{1}(j\in\mathcal{A}_{R})(\mathrm{d}\alpha_{j}-m),$ (8)

where $m=\frac{1}{|\mathcal{A}|}\sum_{j\in\mathcal{A}}\mathrm{d}\alpha_{j}$.

##  4 Fertility Bounds

We experiment with three ways of setting the fertility of the source words:
constant, guided, and predicted. With constant, we set the fertilities of all
source words to a fixed integer value $f$. With guided, we train a word
aligner based on IBM Model 2 (we used fast_align in our experiments, Dyer et
al. (2013)) and, for each word in the vocabulary, we set the fertilities to
the maximal observed value in the training data (or 1 if no alignment was
observed). With the predicted strategy, we train a separate fertility
predictor model using a bi-LSTM tagger.333A similar strategy was recently used
by Gu et al. (2018) as a component of their non-autoregressive NMT model. At
training time, we provide as supervision the fertility estimated by
fast_align. Since our model works with fertility _upper bounds_ and the word
aligner may miss some word pairs, we found it beneficial to add a constant to
this number (1 in our experiments). At test time, we use the expected
fertilities according to our model.

Figure 2: Attention maps for $\operatorname*{\mathsf{softmax}}$ and $\operatorname*{\mathsf{csparsemax}}$ for two De-En sentence pairs (white means zero attention). Repeated words are highlighted. The reference translations are _“This is Moore’s law over the last hundred years”_ and _“I am going to go ahead and select government.”_ | De-En | Ja-En | Ro-En
---|---|---|---
| bleu | meteor | rep | drop | bleu | meteor | rep | drop | bleu | meteor | rep | drop
$\operatorname*{\mathsf{softmax}}$ | 29.51 | 31.43 | 3.37 | 5.89 | 20.36 | 23.83 | 13.48 | 23.30 | 29.67 | 32.05 | 2.45 | 5.59
$\operatorname*{\mathsf{softmax}}$ \+ CovPenalty | 29.69 | 31.53 | 3.47 | 5.74 | 20.70 | 24.12 | 14.12 | 22.79 | 29.81 | 32.15 | 2.48 | 5.49
$\operatorname*{\mathsf{softmax}}$ \+ CovVector | 29.63 | 31.54 | 2.93 | 5.65 | 21.53 | 24.50 | 11.07 | 22.18 | 30.08 | 32.22 | 2.42 | 5.47
$\operatorname*{\mathsf{sparsemax}}$ | 29.73 | 31.54 | 3.18 | 5.90 | 21.28 | 24.25 | 13.09 | 22.40 | 29.97 | 32.12 | 2.19 | 5.60
$\operatorname*{\mathsf{sparsemax}}$ \+ CovPenalty | 29.83 | 31.60 | 3.24 | 5.79 | 21.64 | 24.49 | 13.36 | 21.91 | 30.07 | 32.20 | 2.20 | 5.47
$\operatorname*{\mathsf{sparsemax}}$ \+ CovVector | 29.22 | 31.18 | 3.13 | 6.15 | 21.35 | 24.74 | 10.11 | 21.25 | 29.30 | 31.84 | 2.18 | 5.87
$\operatorname*{\mathsf{csoftmax}}$ ($c=0.2$) | 29.39 | 31.33 | 3.29 | 5.86 | 20.71 | 24.00 | 12.38 | 22.73 | 29.39 | 31.83 | 2.37 | 5.64
$\operatorname*{\mathsf{csparsemax}}$ ($c=0.2$) | 29.85 | 31.76 | 2.67 | 5.23 | 21.31 | 24.51 | 11.40 | 21.59 | 29.77 | 32.10 | 1.98 | 5.44
Table 1:  BLEU, METEOR, REP and DROP scores on the test sets for different attention transformations. | bleu | meteor
---|---|---
constant, $f=2$ | 29.66 | 31.60
constant, $f=3$ | 29.64 | 31.56
guided, | 29.56 | 31.45
predicted, $c=0$ | 29.78 | 31.60
predicted, $c=0.2$ | 29.85 | 31.76
Table 2:  Impact of various fertility strategies for the
$\operatorname*{\mathsf{csparsemax}}$ attention model (De-En).

#### Sink token.

We append an additional $<$sink$>$ token to the end of the source sentence, to
which we assign unbounded fertility ($f_{J+1}=\infty$). The token is akin to
the null alignment in IBM models. The reason we add this token is the
following: without the sink token, the length of the generated target sentence
can never exceed $\sum_{j}f_{j}$ words if we use constrained
softmax/sparsemax. At training time this may be problematic, since the target
length is fixed and the problems in Eqs. 4–6 can become infeasible. By adding
the sink token we guarantee $\sum_{j}f_{j}=\infty$, eliminating the problem.

#### Exhaustion strategies.

To avoid missing source words, we implemented a simple strategy to encourage
more attention to words with larger credit: we redefine the pre-attention word
scores as $\bm{z}_{t}^{\prime}=\bm{z}_{t}+c\bm{u}_{t}$, where $c$ is a
constant ($c=0.2$ in our experiments). This increases the score of words which
have not yet exhausted their fertility (we may regard it as a “soft” lower
bound in Eqs. 4–6).

##  5 Experiments

We evaluated our attention transformations on three language pairs. We focused
on small datasets, as they are the most affected by coverage mistakes. We use
the IWSLT 2014 corpus for De-En, the KFTT corpus for Ja-En (Neubig, 2011), and
the WMT 2016 dataset for Ro-En. The training sets have 153,326, 329,882, and
560,767 parallel sentences, respectively. Our reason to prefer smaller
datasets is that this regime is what brings more adequacy issues and demands
more structural biases, hence it is a good test bed for our methods. We
tokenized the data using the Moses scripts and preprocessed it with subword
units Sennrich et al. (2016) with a joint vocabulary and 32k merge operations.
Our implementation was done on a fork of the OpenNMT-py toolkit Klein et al.
(2017) with the default parameters 444We used a 2-layer LSTM, embedding and
hidden size of 500, dropout 0.3, and the SGD optimizer for 13 epochs.. We used
a validation set to tune hyperparameters introduced by our model. Even though
our attention implementations are CPU-based using NumPy (unlike the rest of
the computation which is done on the GPU), we did not observe any noticeable
slowdown using multiple devices.

As baselines, we use softmax attention, as well as two recently proposed
coverage models:

  * •

CovPenalty (Wu et al., 2016, §7). At test time, the hypotheses in the beam are
rescored with a global score that includes a length and a coverage
penalty.555Since our sparse attention can become $0$ for some words, we
extended the original coverage penalty by adding another parameter $\epsilon$,
set to $0.1$:
$\mathsf{cp}(x;y):=\beta\sum_{j=1}^{J}\log\max\\{\epsilon,\min\\{1,\sum_{t=1}^{|y|}\alpha_{jt}\\}\\}$.
We tuned $\alpha$ and $\beta$ with grid search on $\\{0.2k\\}_{k=0}^{5}$, as
in Wu et al. (2016).

  * •

CovVector (Tu et al., 2016). At training and test time, coverage vectors
$\bm{\beta}$ and additional parameters $\bm{v}$ are used to condition the next
attention step. We adapted this to our bilinear attention by defining
$z_{t,j}=\bm{s}_{t-1}^{\top}(\bm{W}\bm{h}_{j}+\bm{v}{\beta}_{t-1,j})$.

We also experimented combining the strategies above with the sparsemax
transformation.

As evaluation metrics, we report tokenized BLEU, METEOR (Denkowski and Lavie
(2014), as well as two new metrics that we describe next to account for over
and under-translation.666Both evaluation metrics are included in our software
package at www.github.com/Unbabel/
sparse_constrained_attention.

#### REP-score:

a new metric to count repetitions. Formally, given an $n$-gram $s\in V^{n}$,
let $t(s)$ and $r(s)$ be the its frequency in the model translation and
reference. We first compute a sentence-level score

$\displaystyle\sigma(t,r)$ $\displaystyle=$
$\displaystyle\textstyle\lambda_{1}\sum_{s\in V^{n},\,\,t(s)\geq
2}\max\\{0,t(s)-r(s)\\}$ $\displaystyle+$
$\displaystyle\textstyle\lambda_{2}\sum_{w\in V}\max\\{0,t(ww)-r(ww)\\}.$

The REP-score is then given by summing $\sigma(t,r)$ over sentences,
normalizing by the number of words on the reference corpus, and multiplying by
100. We used $n=2$, $\lambda_{1}=1$ and $\lambda_{2}=2$.

#### DROP-score:

a new metric that accounts for possibly dropped words. To compute it, we first
compute two sets of word alignments: from source to reference translation, and
from source to the predicted translation. In our experiments, the alignments
were obtained with fast_align Dyer et al. (2013), trained on the training
partition of the data. Then, the DROP-score computes the percentage of source
words that aligned with some word from the reference translation, but not with
any word from the predicted translation.

Table 1 shows the results. We can see that on average, the sparse models
($\operatorname*{\mathsf{csparsemax}}$ as well as
$\operatorname*{\mathsf{sparsemax}}$ combined with coverage models) have
higher scores on both BLEU and METEOR. Generally, they also obtain better REP
and DROP scores than $\operatorname*{\mathsf{csoftmax}}$ and
$\operatorname*{\mathsf{softmax}}$, which suggests that sparse attention
alleviates the problem of coverage to some extent.

To compare different fertility strategies, we ran experiments on the De-En for
the $\operatorname*{\mathsf{csparsemax}}$ transformation (Table 2). We see
that the Predicted strategy outperforms the others both in terms of BLEU and
METEOR, albeit slightly.

Figure 2 shows examples of sentences for which the
$\operatorname*{\mathsf{csparsemax}}$ fixed repetitions, along with the
corresponding attention maps. We see that in the case of
$\operatorname*{\mathsf{softmax}}$ repetitions, the decoder attends repeatedly
to the same portion of the source sentence (the expression “letzten hundert”
in the first sentence and “regierung” in the second sentence). Not only did
$\operatorname*{\mathsf{csparsemax}}$ avoid repetitions, but it also yielded a
sparse set of alignments, as expected. Appendix B provides more examples of
translations from all models in discussion.

##  6 Conclusions

We proposed a new approach to address the coverage problem in NMT, by
replacing the softmax attentional transformation by sparse and constrained
alternatives: sparsemax, constrained softmax, and the newly proposed
constrained sparsemax. For the latter, we derived efficient forward and
backward propagation algorithms. By incorporating a model for fertility
prediction, our attention transformations led to sparse alignments, avoiding
repeated words in the translation.

## Acknowledgments

We thank the Unbabel AI Research team for numerous discussions, and the three
anonymous reviewers for their insightful comments. This work was supported by
the European Research Council (ERC StG DeepSPIN 758969) and by the Fundação
para a Ciência e Tecnologia through contracts UID/EEA/50008/2013, PTDC/EEI-
SII/7092/2014 (LearnBig), and CMUPERI/TIC/0046/2014 (GoLocal).

## References

  * Almeida and Martins (2013) Miguel B. Almeida and André F. T. Martins. 2013.  Fast and Robust Compressive Summarization with Dual Decomposition and Multi-Task Learning.  In Proc. of the Annual Meeting of the Association for Computational Linguistics.
  * Bahdanau et al. (2014) Dzmitry Bahdanau, Kyunghyun Cho, and Yoshua Bengio. 2014.  Neural machine translation by jointly learning to align and translate.  arXiv preprint arXiv:1409.0473 .
  * Blum et al. (1973) Manuel Blum, Robert W Floyd, Vaughan Pratt, Ronald L Rivest, and Robert E Tarjan. 1973.  Time bounds for selection.  Journal of Computer and System Sciences 7(4):448–461.
  * Brown et al. (1993) Peter F Brown, Vincent J Della Pietra, Stephen A Della Pietra, and Robert L Mercer. 1993.  The mathematics of statistical machine translation: Parameter estimation.  Computational linguistics 19(2):263–311.
  * Denkowski and Lavie (2014) Michael Denkowski and Alon Lavie. 2014.  Meteor universal: Language specific translation evaluation for any target language.  In Proceedings of the EACL 2014 Workshop on Statistical Machine Translation.
  * Dyer et al. (2013) Chris Dyer, Victor Chahuneau, and Noah A Smith. 2013.  A simple, fast, and effective reparameterization of ibm model 2.  Association for Computational Linguistics.
  * Feng et al. (2016) Shi Feng, Shujie Liu, Nan Yang, Mu Li, Ming Zhou, and Kenny Q. Zhu. 2016.  Improving attention modeling with implicit distortion and fertility for machine translation.  In Proceedings of COLING 2016, the 26th International Conference on Computational Linguistics: Technical Papers. The COLING 2016 Organizing Committee, Osaka, Japan, pages 3082–3092.
  * Gehring et al. (2017) Jonas Gehring, Michael Auli, David Grangier, and Yann Dauphin. 2017.  A convolutional encoder model for neural machine translation.  In Proceedings of the 55th Annual Meeting of the Association for Computational Linguistics (Volume 1: Long Papers). volume 1, pages 123–135.
  * Gu et al. (2018) Jiatao Gu, James Bradbury, Caiming Xiong, Victor OK Li, and Richard Socher. 2018\.  Non-autoregressive neural machine translation.  In Proc. of International Conference on Learning Representations.
  * Klein et al. (2017) Guillaume Klein, Yoon Kim, Yuntian Deng, Jean Senellart, and Alexander M. Rush. 2017\.  Opennmt: Open-source toolkit for neural machine translation.  In Proceedings of the 55th Annual Meeting of the Association for Computational Linguistics, ACL 2017, Vancouver, Canada, July 30 - August 4, System Demonstrations. pages 67–72.
  * Koehn and Knowles (2017) Philipp Koehn and Rebecca Knowles. 2017.  Six challenges for neural machine translation.  In Proceedings of the First Workshop on Neural Machine Translation. Association for Computational Linguistics, Vancouver, pages 28–39.
  * Li and Zhu (2017) Junhui Li and Muhua Zhu. 2017.  Learning when to attend for neural machine translation.  arXiv preprint arXiv:1705.11160 .
  * Luong et al. (2015) Thang Luong, Hieu Pham, and Christopher D Manning. 2015.  Effective approaches to attention-based neural machine translation.  In Proceedings of the 2015 Conference on Empirical Methods in Natural Language Processing. pages 1412–1421.
  * Martins and Astudillo (2016) Andre Martins and Ramon Astudillo. 2016.  From softmax to sparsemax: A sparse model of attention and multi-label classification.  In International Conference on Machine Learning. pages 1614–1623.
  * Martins and Kreutzer (2017) André FT Martins and Julia Kreutzer. 2017.  Learning what’s easy: Fully differentiable neural easy-first taggers.  In Proceedings of the 2017 Conference on Empirical Methods in Natural Language Processing. pages 349–362.
  * Mi et al. (2016) Haitao Mi, Baskaran Sankaran, Zhiguo Wang, and Abe Ittycheriah. 2016.  Coverage embedding models for neural machine translation.  In Proceedings of the 2016 Conference on Empirical Methods in Natural Language Processing, EMNLP 2016, Austin, Texas, USA, November 1-4, 2016. pages 955–960.
  * Neubig (2011) Graham Neubig. 2011.  The Kyoto free translation task.  http://www.phontron.com/kftt.
  * Pardalos and Kovoor (1990) Panos M. Pardalos and Naina Kovoor. 1990.  An algorithm for a singly constrained class of quadratic programs subject to upper and lower bounds.  Mathematical Programming 46(1):321–328.
  * Sennrich et al. (2016) Rico Sennrich, Barry Haddow, and Alexandra Birch. 2016.  Neural machine translation of rare words with subword units.  In Proceedings of the 54th Annual Meeting of the Association for Computational Linguistics, ACL 2016, August 7-12, 2016, Berlin, Germany, Volume 1: Long Papers.
  * Sutskever et al. (2014) Ilya Sutskever, Oriol Vinyals, and Quoc V Le. 2014.  Sequence to sequence learning with neural networks.  In Advances in neural information processing systems. pages 3104–3112.
  * Tu et al. (2017a) Zhaopeng Tu, Yang Liu, Zhengdong Lu, Xiaohua Liu, and Hang Li. 2017a.  Context gates for neural machine translation .
  * Tu et al. (2017b) Zhaopeng Tu, Yang Liu, Lifeng Shang, Xiaohua Liu, and Hang Li. 2017b.  Neural machine translation with reconstruction.  In AAAI. pages 3097–3103.
  * Tu et al. (2016) Zhaopeng Tu, Zhengdong Lu, Yang Liu, Xiaohua Liu, and Hang Li. 2016.  Modeling coverage for neural machine translation.  In Proceedings of the 54th Annual Meeting of the Association for Computational Linguistics.
  * Vaswani et al. (2017) Ashish Vaswani, Noam Shazeer, Niki Parmar, Jakob Uszkoreit, Llion Jones, Aidan N Gomez, Łukasz Kaiser, and Illia Polosukhin. 2017.  Attention is all you need.  In Advances in Neural Information Processing Systems. pages 6000–6010.
  * Wu et al. (2016) Yonghui Wu, Mike Schuster, Zhifeng Chen, Quoc V Le, Mohammad Norouzi, Wolfgang Macherey, Maxim Krikun, Yuan Cao, Qin Gao, Klaus Macherey, et al. 2016.  Google’s neural machine translation system: Bridging the gap between human and machine translation.  arXiv preprint arXiv:1609.08144 .

##  Appendix A Proof of Proposition 1

We provide here a detailed proof of Proposition 1.

###  A.1 Forward Propagation

The optimization problem can be written as

$\displaystyle\operatorname*{\mathsf{csparsemax}}(\bm{z},\bm{u})=\operatorname*{\mathsf{arg\,min}}$
$\displaystyle\frac{1}{2}\|\bm{\alpha}\|^{2}-{\bm{z}}^{\top}{\bm{\alpha}}$
s.t.
$\displaystyle\left\\{\begin{array}[]{l}{\mathbf{1}}^{\top}{\bm{\alpha}}=1\\\
\mathbf{0}\leq\bm{\alpha}\leq\bm{u}.\end{array}\right.$ (10)

The Lagrangian function is:

$\displaystyle\mathcal{L}(\bm{\alpha},\tau,\bm{\mu},\bm{\nu})=-\frac{1}{2}\|\bm{\alpha}\|^{2}-{\bm{z}}^{\top}{\bm{\alpha}}+\tau({\mathbf{1}}^{\top}{\bm{\alpha}}-1)-{\bm{\mu}}^{\top}{\bm{\alpha}}+{\bm{\nu}}^{\top}{(\bm{\alpha}-\bm{u})}.$
(11)

To obtain the solution, we invoke the Karush-Kuhn-Tucker conditions. From the
stationarity condition, we have
$\mathbf{0}=\bm{\alpha}-\bm{z}+\tau\mathbf{1}-\bm{\mu}+\bm{\nu}$, which due to
the primal feasibility condition implies that the solution is of the form:

$\bm{\alpha}=\bm{z}-\tau\mathbf{1}+\bm{\mu}-\bm{\nu}.$ (12)

From the complementarity slackness condition, we have that
$0<\alpha_{j}<u_{j}$ implies that $\mu_{j}=\nu_{j}=0$ and therefore
$\alpha_{j}=z_{j}-\tau$. On the other hand, $\mu_{j}>0$ implies
$\alpha_{j}=0$, and $\nu_{j}>0$ implies $\alpha_{j}=u_{j}$. Hence the solution
can be written as $\alpha_{j}=\max\\{0,\min\\{u_{j},z_{j}-\tau\\}$, where
$\tau$ is determined such that the distribution normalizes:

$\tau=\frac{\sum_{j\in\mathcal{A}}z_{j}+\sum_{j\in\mathcal{A}_{R}}u_{j}-1}{|\mathcal{A}|},$
(13)

with $\mathcal{A}=\\{j\in[J]\,\,|\,\,0<\alpha_{j}<u_{j}\\}$ and
$\mathcal{A}_{R}=\\{j\in[J]\,\,|\,\,\alpha_{j}=u_{j}\\}$. Note that $\tau$
depends itself on the set $\mathcal{A}$, a function of the solution. In §A.3,
we describe an algorithm that searches the value of $\tau$ efficiently.

###  A.2 Gradient Backpropagation

We now turn to the problem of backpropagating the gradients through the
constrained sparsemax transformation. For that, we need to compute its
Jacobian matrix, i.e., the derivatives $\frac{\partial\alpha_{i}}{\partial
z_{j}}$ and $\frac{\partial\alpha_{i}}{\partial u_{j}}$ for $i,j\in[J]$. Let
us first express $\bm{\alpha}$ as

$\alpha_{i}=\left\\{\begin{array}[]{ll}0,&i\in\mathcal{A}_{L},\\\
z_{i}-\tau,&i\in\mathcal{A},\\\ u_{i},&i\in\mathcal{A}_{R},\end{array}\right.$
(14)

with $\tau$ as in Eq. 13. Note that we have ${\partial\tau}/{\partial
z_{j}}=\mathds{1}(j\in\mathcal{A})/|\mathcal{A}|$ and
${\partial\tau}/{\partial
u_{j}}=\mathds{1}(j\in\mathcal{A}_{R})/|\mathcal{A}|$. Thus, we have the
following:

$\frac{\partial\alpha_{i}}{\partial
z_{j}}=\left\\{\begin{array}[]{ll}1-1/|\mathcal{A}|,&\text{if
$j\in\mathcal{A}$ and $i=j$}\\\ -1/|\mathcal{A}|,&\text{if $i,j\in\mathcal{A}$
and $i\neq j$}\\\ 0,&\text{otherwise,}\end{array}\right.$ (15)

and

$\frac{\partial\alpha_{i}}{\partial
u_{j}}=\left\\{\begin{array}[]{ll}1,&\text{if $j\in\mathcal{A}_{R}$ and
$i=j$}\\\ -1/|\mathcal{A}|,&\text{if $j\in\mathcal{A}_{R}$ and
$i\in\mathcal{A}$}\\\ 0,&\text{otherwise.}\end{array}\right.$ (16)

Finally, we obtain:

$\displaystyle\mathrm{d}z_{j}$ $\displaystyle=$
$\displaystyle\sum_{i}\frac{\partial\alpha_{i}}{\partial
z_{j}}\mathrm{d}\alpha_{i}$ (17) $\displaystyle=$
$\displaystyle\mathds{1}(j\in\mathcal{A})\left(\mathrm{d}\alpha_{j}-\frac{\sum_{i\in\mathcal{A}}\mathrm{d}\alpha_{i}}{|\mathcal{A}|}\right)$
$\displaystyle=$
$\displaystyle\mathds{1}(j\in\mathcal{A})(\mathrm{d}\alpha_{j}-m),$

and

$\displaystyle\mathrm{d}u_{j}$ $\displaystyle=$
$\displaystyle\sum_{i}\frac{\partial\alpha_{i}}{\partial
u_{j}}\mathrm{d}\alpha_{i}$ (18) $\displaystyle=$
$\displaystyle\mathds{1}(j\in\mathcal{A}_{R})\left(\mathrm{d}\alpha_{j}-\frac{\sum_{i\in\mathcal{A}}\mathrm{d}\alpha_{i}}{|\mathcal{A}|}\right)$
$\displaystyle=$
$\displaystyle\mathds{1}(j\in\mathcal{A}_{R})(\mathrm{d}\alpha_{j}-m),$

where $m=\frac{1}{|\mathcal{A}|}\sum_{j\in\mathcal{A}}\mathrm{d}\alpha_{j}$.

Algorithm 1 Pardalos and Kovoor’s Algorithm

1: input: $\bm{a},\bm{b},\bm{c},d$

2: Initialize working set $\mathcal{W}\leftarrow\\{1,\ldots,J\\}$

3: Initialize set of split points:
$\mathcal{P}\leftarrow\\{a_{j},b_{j}\\}_{j=1}^{J}\cup\\{\pm\infty\\}$

4: Initialize $\tau_{\mathrm{L}}\leftarrow-\infty$,
$\tau_{\mathrm{R}}\leftarrow\infty$, $s_{\mathrm{tight}}\leftarrow 0$,
$\xi\leftarrow 0$.

5: while $\mathcal{W}\neq\varnothing$ do

6: Compute $\tau\leftarrow\mathrm{Median}(\mathcal{P})$

7: Set $s\leftarrow
s_{\mathrm{tight}}+\sum_{j\in\mathcal{W}\,|\,b_{i}<\tau}c_{j}b_{j}+\sum_{j\in\mathcal{W}\,|\,a_{j}>\tau}c_{j}a_{j}+(\xi+\sum_{j\in\mathcal{W}\,|\,a_{j}\leq\tau\leq
b_{j}}c_{j})\tau$

8: If $s\leq d$, set $\tau_{\mathrm{L}}\leftarrow\tau$; if $s\geq d$, set
$\tau_{\mathrm{R}}\leftarrow\tau$

9: Reduce set of split points:
$\mathcal{P}\leftarrow\mathcal{P}\cap[\tau_{\mathrm{L}},\tau_{\mathrm{R}}]$

10: Update tight-sum: $s_{\mathrm{tight}}\leftarrow
s_{\mathrm{tight}}+\sum_{j\in\mathcal{W}\,|\,b_{i}<\tau_{L}}c_{j}b_{j}+\sum_{j\in\mathcal{W}\,|\,a_{j}>\tau_{R}}c_{j}a_{j}$

11: Update slack-sum:
$\xi\leftarrow\xi+\sum_{j\in\mathcal{W}\,|\,a_{j}\leq\tau_{L}\wedge
b_{j}\geq\tau_{R}}c_{j}$

12: Update working set:
$\mathcal{W}\leftarrow\\{j\in\mathcal{W}\,|\,\tau_{L}<a_{j}<\tau_{R}\vee\tau_{L}<b_{j}<\tau_{R}\\}$

13: end while

14: Define $y^{*}\leftarrow(d-s_{\mathrm{tight}})/\xi$

15: Set $x_{j}^{\star}=\max\\{a_{j},\min\\{b_{j},y\\}\\},\,\,\forall j\in[J]$

16: output: $\bm{x}^{\star}$.

###  A.3 Linear-Time Evaluation

Finally, we present an algorithm to solve the problem in Eq. 6 in linear time.

Pardalos and Kovoor (1990) describe an algorithm, reproduced here as Algorithm
1, for solving a class of singly-constrained convex quadratic problems, which
can be written in the form above (where each $c_{j}\geq 0$):

$\displaystyle\mathrm{minimize}\,\,$
$\displaystyle\sum_{j=1}^{J}c_{j}x_{j}^{2}$ $\displaystyle\mathrm{s.t.}\,\,$
$\displaystyle\textstyle\sum_{j=1}^{J}c_{j}x_{j}=d,$ $\displaystyle\textstyle
a_{j}\leq x_{j}\leq b_{j},\quad j=1,\ldots,J.$ (19)

The solution of the problem in Eq. A.3 is of the form
$x_{j}^{\star}=\max\\{a_{j},\min\\{b_{j},y\\}\\}$, where $y\in[a_{j},b_{j}]$
is a constant. The algorithm searches the value of this constant (which is
similar to $\tau$ in our problem), which lies in a particular interval of
split-points (line 3), iteratively shrinking this interval. The algorithm
requires computing medians as a subroutine, which can be done in linear time
(Blum et al., 1973). The overall complexity in $O(J)$ (Pardalos and Kovoor,
1990). The same algorithm has been used in NLP by Almeida and Martins (2013)
for a budgeted summarization problem.

To show that this algorithm applies to the problem of evaluating
$\operatorname*{\mathsf{csparsemax}}$, it suffices to show that our problem in
Eq. 6 can be rewritten in the form of Eq. A.3. This is indeed the case, if we
set:

$\displaystyle x_{j}$ $\displaystyle=$ $\displaystyle\frac{p_{j}-z_{j}}{2}$
(20) $\displaystyle a_{j}$ $\displaystyle=$ $\displaystyle-z_{j}/2$ (21)
$\displaystyle b_{j}$ $\displaystyle=$ $\displaystyle(u_{j}-z_{j})/2$ (22)
$\displaystyle c_{j}$ $\displaystyle=$ $\displaystyle 1$ (23) $\displaystyle
d$ $\displaystyle=$ $\displaystyle\frac{1-\sum_{j=1}^{J}z_{j}}{2}.$ (24)

##  Appendix B Examples of Translations

We show some examples of translations obtained for the German-English language
pair with different systems. Blue highlights the parts of the reference that
are correct and red highlights the corresponding problematic parts of
translations, including repetitions, dropped words or mistranslations.

input | überlassen sie das ruhig uns .
---|---
reference |  leave that up to us .
$\operatorname*{\mathsf{softmax}}$ |  give us a silence .
$\operatorname*{\mathsf{sparsemax}}$ | leave that to us .
$\operatorname*{\mathsf{csoftmax}}$ |  let’s leave that .
$\operatorname*{\mathsf{csparsemax}}$ | leave it to us .
input | so ungefähr , sie wissen schon .
---|---
reference |  like that , you know .
$\operatorname*{\mathsf{softmax}}$ |  so , you know , you know .
$\operatorname*{\mathsf{sparsemax}}$ |  so , you know , you know .
$\operatorname*{\mathsf{csoftmax}}$ |  so , you know , you know .
$\operatorname*{\mathsf{csparsemax}}$ | like that , you know .
input | und wir benutzen dieses wort mit solcher verachtung .
---|---
reference | and we say that word with such contempt .
$\operatorname*{\mathsf{softmax}}$ | and we use this word with such contempt contempt .
$\operatorname*{\mathsf{sparsemax}}$ | and we use this word with such contempt .
$\operatorname*{\mathsf{csoftmax}}$ | and we use this word with like this .
$\operatorname*{\mathsf{csparsemax}}$ | and we use this word with such contempt .
input | wir sehen das dazu , dass phosphor wirklich kritisch ist .
---|---
reference | we can see that phosphorus is really critical .
$\operatorname*{\mathsf{softmax}}$ | we see that that phosphorus is really critical .
$\operatorname*{\mathsf{sparsemax}}$ | we see that that phosphorus really is critical .
$\operatorname*{\mathsf{csoftmax}}$ | we see that that phosphorus is really critical .
$\operatorname*{\mathsf{csparsemax}}$ | we see that phosphorus is really critical .
input | also müssen sie auch nicht auf klassische musik verzichten , weil sie kein instrument spielen .
---|---
reference | so you don’t need to abstain from listening to classical music because you don’t play an instrument .
$\operatorname*{\mathsf{softmax}}$ | so you don’t have to rely on classical music because you don’t have an instrument .
$\operatorname*{\mathsf{sparsemax}}$ | so they don’t have to kill classical music because they don’t play an instrument .
$\operatorname*{\mathsf{csoftmax}}$ | so they don’t have to rely on classical music , because they don’t play an instrument .
$\operatorname*{\mathsf{csparsemax}}$ | so you don’t have to get rid of classical music , because you don’t play an instrument .

input | je mehr ich aber darüber nachdachte , desto mehr kam ich zu der ansicht , das der fisch etwas weiß .
---|---
reference | the more i thought about it , however , the more i came to the view that this fish knows something .
$\operatorname*{\mathsf{softmax}}$ | the more i thought about it , the more i got to the point of the fish .
$\operatorname*{\mathsf{sparsemax}}$ | the more i thought about it , the more i got to the point of view of the fish .
$\operatorname*{\mathsf{csoftmax}}$ | but the more i thought about it , the more i came to mind , the fish .
$\operatorname*{\mathsf{csparsemax}}$ | the more i thought about it , the more i came to the point that the fish knows .
input |  all diese menschen lehren uns , dass es noch andere existenzmöglichkeiten , andere denkweisen , andere wege zur orientierung auf der erde gibt .
---|---
reference |  all of these peoples teach us that there are other ways of being , other ways of thinking , other ways of orienting yourself in the earth .
$\operatorname*{\mathsf{softmax}}$ |  all of these people teach us that there are others , other ways , other ways of guidance to the earth .
$\operatorname*{\mathsf{sparsemax}}$ |  all these people are teaching us that there are other options , other ways , different ways of guidance on earth .
$\operatorname*{\mathsf{csoftmax}}$ |  all of these people teach us that there’s other ways of doing other ways of thinking , other ways of guidance on earth .
$\operatorname*{\mathsf{csparsemax}}$ |  all these people teach us that there are other actors , other ways of thinking , other ways of guidance on earth .
input | in der reichen welt , in der oberen milliarde , könnten wir wohl abstriche machen und weniger nutzen ,
---|---
| aber im durchschnitt wird diese zahl jedes jahr steigen und sich somit
insgesamt mehr als verdoppeln ,
| die zahl der dienste die pro person bereitgestellt werden .
reference | in the rich world , perhaps the top one billion , we probably could cut back and use less , but every year , this number ,
| on average , is going to go up , and so , over all , that will more than
double the services delivered per person .
$\operatorname*{\mathsf{softmax}}$ | in the rich world , in the upper billion , we might be able to do and use less use , but on average , that number
| is going to increase every year and so on , which is the number of services
that are being put in .
$\operatorname*{\mathsf{sparsemax}}$ | in the rich world , in the upper billion , we may be able to do and use less use , but in average , that number
| is going to rise every year , and so much more than double , the number of
services that are being put together .
$\operatorname*{\mathsf{csoftmax}}$ | in the rich world , in the upper billion , we might be able to take off and use less , but in average , this number
| is going to increase every year and so on , and that’s the number of people
who are being put together per person .
$\operatorname*{\mathsf{csparsemax}}$ | in the rich world , in the upper billion , we may be able to turn off and use less , but in average , that number will
| rise every year and so far more than double , the number of services that
are being put into a person .