# BarraCUDA:
Bringing Electromagnetic Side Channel Into Play to Steal the Weights of Neural
Networks from NVIDIA GPUs

Péter Horváth1, Łukasz Chmielewski2, Leo Weissbart1, Lejla Batina1, Yuval
Yarom3  1Radboud University  2Masaryk University  3Ruhr University Bochum

###### Abstract

Over the last decade, applications of neural networks have spread to cover all
aspects of life. A large number of companies base their businesses on building
products that use neural networks for tasks such as face recognition, machine
translation, and autonomous cars. They are being used in safety and security-
critical applications like high definition maps and medical wristbands, or in
globally used products like Google Translate and ChatGPT. Much of the
intellectual property underpinning these products is encoded in the exact
configuration of the neural networks. Consequently, protecting these is of
utmost priority to businesses. At the same time, many of these products need
to operate under a strong threat model, in which the adversary has unfettered
physical control of the product.

Past work has demonstrated that with physical access, attackers can reverse
engineer neural networks that run on scalar microcontrollers, like ARM Cortex
M3. However, for performance reasons, neural networks are often implemented on
highly-parallel general purpose graphics processing units (GPGPUs), and so
far, attacks on these have only recovered course-grained information on the
structure of the neural network, but failed to retrieve the detailed weights
and biases.

In this work, we present BarraCUDA, a novel attack on GPGPUs that can
completely extract the parameters of neural networks. BarraCUDA uses
correlation electromagnetic analysis to recover the weights and biases in the
convolutional layers of neural networks. We use BarraCUDA to attack the
popular NVIDIA Jetson Nano device, demonstrating successful reverse
engineering of neural networks in a highly parallel and noisy environment.

##  1 Introduction

The field of machine learning has seen an explosive increase in interest and
use over the last decade. In particular, deep learning has proven to be a
versatile technique that provides state-of-the-art performance for many real-
world application. The use of deep learning proved useful for a broad range of
domains including playing chess [50], object detection [35], image
classification [29, 51, 34, 17, 23], audio processing [43], forecasting [31,
45, 47, 48] and natural language processing [40]. Thus, applications of deep
learning are changing many areas of our lives and have become indispensable to
our everyday life.

Deep learning typically utilizes artificial neural networks, consisting of
multiple layers of (simulated) neurons. In a nutshell, a neuron takes a number
of inputs and computes a non-linear function of the weighted sum of the
inputs. A common example of such a function is a rectified linear unit (ReLU),
which basically returns its input if it is positive and zero otherwise. When
designing a deep learning solution for a problem, the designer first chooses
the network architecture, which specifies the layers of neurons, including
their sizes and types, as well as how the neurons are connected, i.e. which
neurons outputs are connected to which inputs. The designer then trains the
network, selecting the weights used for each weighted sum as well as bias
values that are added to the sums before the computation of the non-linear
function.

Training a network for any non-trivial example is a resource intensive
process. There is a need to curate a specialized dataset of correctly labeled
samples that can be used for the training, and the training process often
requires days and even weeks of computation on specialized high-performance
hardware, such as large quantities of graphical processing units (GPUs).
Moreover, the design of the right network architecture for a given purpose
requires specialized expertise. Consequently, the design and parameters of
trained models are in many cases considered as trade secrets, which vendors
try to protect against undesired disclosure.

Side-channel attacks, which exploit information leakage via unintended effects
of performing computation, are a major threat to the confidentiality of deep
learning designs. Such attacks against neural network implementations on CPUs
have been demonstrated using both power analysis [13] and microarchitectural
attacks [54], among others. However, to achieve a better performance neural
networks are often implemented on GPUs. So far, attacks on such
implementations only recover the network architecture, but not the parameters
[16, 37]. Moreover, these attacks only target open-source neural networks,
implemented using PyTorch [41] and TensorFlow [12], while we target the
commonly used NVIDIA closed-source TensorRT framework [10].

Therefore, this work focuses on the following research questions:

Are implementations of neural networks on GPU vulnerable to parameter
extraction using side-channel analysis?

### Our Contribution

In this work we answer the question in the affirmative. We analyze NVIDIA’s
closed-source neural network framework TensorRT [10] executing on an
industrial-strength Jetson Nano device [9]. While TensorRT is a closed-source
framework, the produced CUDA binaries can be disassembled.

We demonstrate a side-channel attack that monitors the electromagnetic (EM)
emanations emitted from the device. The attack allows full extraction of the
parameters of the neural network running on the device, overcoming the
measurement noise and the uncertainty inherent in closed platforms.

The attack uses correlation electromagnetic analysis (CEMA) to recover the
weights and biases of neural networks. For the attack, we collect and process
15 million traces over a period of 10 days. We then show how to analyze the
traces to recover weights and biases used in the target network.

The results of attack demonstrate that GPU implementations of deep learning
are not protected against side-channel attacks, even when the source code is
not available.

Disclosure. We notified NVIDIA of the found vulnerabilities on September 22nd,
2023. We shortly received a response on 24/09/2023 that they are looking into
our findings and as a result of mutual agreement the findings would be not
released until December 13rd, 2023. On November 22nd, 2023 NVIDIA notified us
that they reviewed our disclosure and classified the vulnerabilities as a low
severity issue while also recommending the users to follow guidelines to
prevent physical access to secure systems to avoid information leakage.

Organization. The rest of this paper is organized as follows: After providing
the necessary background on side-channel attacks and GPU architecture (Section
2), we describe our weight extraction attack. in Section 3. Finally, Section 4
covers the limitations, possible extensions, countermeasures, and related
work.

##  2 Background

Here we introduce the concepts of side channels and related techniques that
were used for the attacks and provide background on NVIDIA’S CUDA programming
model and GPU architecture.

###  2.1 Side-Channel Analysis

Side-Channel Analysis (SCA) exploits the physical leakages of electronic
devices to extract secret information [28, 27]. Leakage can occur through
various channels, including power consumption, electromagnetic, timing,
optical or sound, and various types of secret information can leak. In
academic settings, side-channel attacks were first introduced in the 90’s
targeting constrained cryptographic devices [28, 27] and pose ever since a
constant threat to the security of various embedded systems. In this work, we
exploit the EM side channel.

####  2.1.1 Electromagnetic emanation

The electromagnetic (EM) emanations from a computing device correlate with the
code and the data that the device processes. This correlation has been used to
break cryptographic implementations [44, 30], reverse engineer neural networks
[13, 16] and eavesdrop on display units [20, 24, 36].

In Simple EM Analysis (SEMA), the attacker samples the emitted signal during
the execution of the target program, creating a _trace_ of measurements. The
attacker then analyzes the trace to identify components that correspond to the
secret information.

In Correlation EM Analysis (CEMA), the attacker builds a model of the leakage
based on the data the program processes. The attacker then executes the
program multiple times with different inputs, and looks for correlation
between the predicted leakage based on the model and the observed leakage at
some points of interest (POIs) in the collected traces. Two common models are
the Hamming weight (HW) model, which predicts that the leakage is linear with
the number of set bits in the data (i.e. its Hamming weight), and the Hamming
distance (HD) model, which predicts that the leakage is linear with the number
of bits that need to be changed between consecutive data values.

####  2.1.2 Leakage evaluation

Multiple methods have been proposed to evaluate the leakage of embedded
devices, such as intermediate correlation, Test Vector Leakage Assessment
(TVLA) [49] and ${\chi}^{2}$-test [39]. In this paper, we use intermediate
correlation and TVLA to place our EM probe at a location that leaks the most
information and detect where each intermediate value leaks in the collected
measurements in the parameter extraction attack presented in Section 3.

###  2.2 CUDA programming model

To leverage the parallelism offered by GPUs, NVIDIA exposes the CUDA
programming model [5] to developers. In this model, multiple abstraction
levels exist and each level has different implications with respect to the GPU
hardware. The lowest level of abstraction is the thread, which executes a CUDA
function defined by the developer. The number of threads executing the CUDA
function in parallel is specified at the time of invoking the
function.111Functions in CUDA are also called _kernels_. We use the term
function to avoid confusion with kernels in CNNs. Subsequently, multiple
threads can be grouped together into a single block of threads. Threads in a
block have a per block on-chip shared memory region where they can exchange
data with other threads in the block. Blocks of threads form a grid of thread
blocks. Each block in a grid executes independently from the other blocks, but
all blocks in the grid share the same off-chip global memory region.

Figure 1: CUDA programming interface and GPU hardware implementation.

###  2.3 GPU Streaming Multiprocessor

Figure 2: GPU Streaming Multiprocessor architecture.

When a CUDA function is invoked, the parallel threads execute on the Streaming
Multiprocessor (SM) of the GPU. A GPU can consist of one or more SMs to
further improve parallelism. Our target, the Jetson Nano, features a Tegra X1
System-On-Chip that consists of one SM of the Maxwell architecture [11] and
CUDA Compute Capability 5.3 [1] from NVIDIA. As illustrated in Figure 1, when
blocks of threads are scheduled onto a particular SM, the threads in the
blocks are divided into groups of 32 threads, also called warps. Warp is the
lowest level of execution unit and is of great concern for SCA. Every warp is
assigned to a particular Processing Unit (PU) where every PU has dedicated
resources to schedule and execute warps in the SM. Figure 2 shows the building
blocks of the Maxwell SM and the PUs in the SM.222Graphics related components,
such as the PolyMorph engine, are omitted in our figure as these are not
relevant for our target application. The warp scheduler in a PU is responsible
for scheduling and issuing instructions for warps that are ready. Every PU has
a set of 32 CUDA cores for arithmetic instructions as well as load/store units
(LD/ST) and special function units (SPU) for memory operations and
transcendental functions, respectively. In addition, a pair of PUs share a
texture/L1 cache but all PUs have access to the same shared memory on the
chip. More importantly for side-channel attacks, each PU also has a dedicated
register file. Since the details of assigning individual warps to PUs are not
publicly disclosed information, multiple register files could mean that a
CPA/CEMA attack is significantly hindered. If a particular warp, one that
calculates our targeted intermediate variable, is assigned to different PUs in
subsequent executions, then our targeted intermediate variable is saved into
multiple, different register files which potentially results in different
power/EM signatures. Therefore, mounting side-channel attacks that rely on the
HD leakage model might become difficult on the GPU, but we present a
successful attack using the HW as well as the HD leakage model.

##  3 Weight and bias extraction attack

In this section, we detail our attack which recovers the weights and biases of
a CNN with 2 convolutional layers, each containing one kernel. First, we
describe the threat model and explore the NN implementations from NVIDIA on
the assembly level. Second, we describe our experiment setup and the process
of establishing points of interest (PoI) where the implementations leak
information. Lastly, we apply CEMA on intermediate values in the convolution
operation to extract the weights and bias of each layer. To the best of our
knowledge, this is the first time that the weights and biases of NNs are
extracted from implementations running on an NVIDIA GPU. In addition, the
experiments are carried out using half-precision calculations in the
convolutional layers.

###  3.1 Threat model

In this scenario, we consider an adversary who has access to the target
device, can observe random inputs fed to the device, and knows the
architecture of the NN that runs on the device. In addition, we also assume
that the adversary can collect electromagnetic side-channel measurements. We
assume that the architecture of the target NN is already extracted as the
architecture of a NN needs to be known for the NN parameters to be extracted
using CEMA.

###  3.2 Exploring TensorRT implementations

The approach of NVIDIA’s TensorRT framework is “one size fits many”, that is,
one implementation can be used for layers with different kernel sizes, input
sizes, batch sizes etc. TensorRT provides information [8] about the optimized
model and this information includes the names of the CUDA function
implementations that are used for the model as well as a hexadecimal
undocumented value called Tacticvalue.333While there is no documentation about
Tacticvalue, it might impact the implementation’s execution since it might be
used in predicated instructions.  Using this information, it is possible to
extract the assembly device code that is executed by each thread on the GPU.
One can use cuobjdump [4] from the CUDA toolkit to extract all the GPU
assembly code that can be found in the different CUDA shared library files
(CuBLAS, CuDNN, TensorRT). Eventually, we found the CUDA function
corresponding to the implementation of a convolutional layer with ReLU
activation in the libnvinfer.so file by the name that can be retrieved using
the TensorRT API.

In our parameter extraction attack, we reverse engineer the weights of a NN
with two convolutional layers each with a kernel containing 5 weights. In the
shared library files, 4 implementations can be found for a convolutional layer
with half-precision calculations and ReLU activation, named:

$maxwell\\_fp16x2\\_hcudnn\\_fp16x2\\_128x\\_\textbf{k}\\_relu\\_\textbf{t}\\_nn\\_v1,$

where

$k\in\\{32,64,128\\},$ $t\in\\{small,medium,large,interior\\}.$

Based on repeated testing, the $small,medium$ and $interior$ implementations
can all be an implementation that the TensorRT API chooses for each
convolutional layer in our 2-layer model configuration. This is possible if
the timing differences are small in the implementations. Thus, it could be
that the same implementation is used for both layers or a combination of two.

Furthermore, based on testing and using the information provided by the API
[8], these are also the implementations that are chosen by the API for the
large-scale neural networks such as MobileNet [25].

The names of implementations also include $fp16x2$ which refers to the half2
[6] data type in CUDA. When half2 is used, a register is packed with 2 half-
precision floating point values and the GPU can execute a floating point
instruction on each half in parallel, doubling the throughput of the GPU, as
opposed to using single-precision floating point values. This also means
additional noise for SCA in addition to the parallel threads executing in the
GPU.

###  3.3 Convolutional layer structure

Although there are 4 possible implementations for a convolutional layer with
ReLU activation, these implementations are very similar in their structure and
they consist of the following blocks:

  1. 1.

block of initialization instructions,

  2. 2.

block of convolution operations, and

  3. 3.

block of bias addition and ReLU calculation.

1.

Init block. The first main block consists of instructions to set up the CUDA
function. In this block, 64 accumulator registers are initialized which are
later used to calculate the sum of multiplications in the convolution
operation.

2.

Convolutional block. The second block is where the convolution operations are
carried out. This block consists of vectorized load and HFMA2 instructions. In
addition, it contains predicated load instructions from global memory. Since
registers are packed with two 16-bit floating point values, it is also crucial
to know exactly how the input data is loaded into the registers. If this
information is available, it sheds light on how the convolution operation is
carried out in the implementations. According to the API, the input is
reformatted so that it is laid out in “Two wide channel vectorized row major
FP16 format”. This means that each register holding inputs is loaded with two
half-precision floating point inputs from two different input channels.
Therefore, this also suggests that each register holds weights from different
channels in a kernel. In the case of our one dimensional convolution, we set
the number of input channels to 1, as it is common for one-dimensional
convolutional use cases. In other cases, such as RGB image inputs, the number
of input channels would be 3 so each register would hold input values from two
different input channels. With multiple input channels processed in parallel,
our CEMA attack would then not require more time to get all the weights
because the channels in the kernels are independent from each other.
Additionally, the code in the second block indicates that each accumulator
register is used in 8 HFMA2 instruction in this block, unless a predicated
jump instruction is executed at the end of the code block which jumps back to
the start of this block. This suggests that the execution of the jump
instruction likely depends on the kernel size in the layer.

3.

ReLU block. The third block, repeated four times, first calculates the sum of
the lower and higher half of each register. This suggests that the channel
combinations of the convolutional operation, after the sum of multiplications,
are executed in these blocks. Then, half2 floating point additions and HSET
instructions are executed in the block suggesting the calculation of the bias
addition and the ReLU output.

Overall, the four implementations are very similar. For our particular two-
layer model configuration, more than 50% of the times the chosen
implementation by the API is the $relu\\_small$ implementation for both
layers, so we choose to carry out the parameter extraction attack with this
implementation. However, we also analyzed the $relu\\_medium$ and
$relu\\_interior$ implementations in our leakage detection process in Section
3.6.

Figure 3: Location of the Langer EM probe. The probe tip is located between
two capacitors. Figure 4: Raw trace of the whole operation on the GPU. The 2
layers are clearly separated in the traces. Additionally, the CUDA device-to-
host memory copy is also clearly visisble in the end of the trace. Figure 5:
Result of TVLA for the multiplication with the first weight in the kernel in
convolutional block. The top figure depicts an elastically aligned trace while
the bottom figure shows the corresponding absolute t-test statistics at each
time sample. There are multiple leaking spots with $|t|>4.5$ but the overall
leakage is not high. Figure 6: Result of fixed vs random bias TVLA. The top
figure depicts an example trace while the bottom figure shows the
corresponding absolute t-test statistics at each time sample. There is clear
leakage scattered across multiple clock cycles with $|t|>>4.5$. Figure 7:
Segmented raw trace of first layer. The first highlighted block contains the
initialization instructions. The second highlighted block is where the
convolution operation is carried out. The third highlighted block corresponds
to the bias addition and ReLU activation calculation.

###  3.4 Experiment setup

In our setup, the Jetson Nano’s GPU cores operate at the highest possible
clock frequency at 921 MHz. In order to collect electromagnetic traces, we use
the Lecroy 8404M-MS oscilloscope at a sampling rate of 10GS/s with the Langer
MFA-R 0.275 near-field probe [7]. Each trace contains measurements of the
electromagnetic emanations of the GPU when the target neural network executes
inference on random inputs. By scanning the Tegra SoC and the nearby
capacitors with the EM probe, we found multiple interesting locations that
leak information about the parameters of neural network. Eventually, we
observed the best signal between two capacitors, as shown in Figure 3. In
terms of leakage magnitude using TVLA, locations above the chip and between
the capacitors are very similar. However, the emitted signal between the
capacitors provides less noisier traces which makes trace acquisition and
alignment easier. To trigger the oscilloscope to collect measurements when the
inference of the NN starts, we use the rising edge of inference signal as
trigger.

###  3.5 Trace preprocessing

Figure 4 shows an example trace of the inference operation of a CNN with two
layers. The two layers are clearly separated and display similar patterns.
However, the first layer takes considerably more time to execute. Since the
input sizes of the layers are 500 and 496, respectively, this difference alone
might not explain the observed difference in execution times. However, the GPU
has various caches which might explain this behaviour. More importantly, we
observed that different patterns can repeat in the second layer from trace to
trace. Whereas the patterns are very similar in the first layer for all
collected traces, the same does not hold for the second layer. For the second
layer, only about 40% of the collected traces display completely similar
patterns. We tested this behaviour with the other implementations as well and
this issue only comes up when the same implementation is used for each layer.
For example, the combination of the $relu\\_small$ and $relu\\_interior$
implementations for the first and second layers, respectively, result in
uniform patterns for both layers. In contrast, if the $relu\\_small$
implementation is used for both layers, then it results in non-uniform
patterns for the second layer. We also tested this behaviour with networks
with more than two convolutional layers, and as soon as the same
implementation is used in two different layers the non-uniform behaviour is
encountered. Regardless, as mentioned in Section 3.2, we carry out the
parameter extraction with the $relu\\_small$ as the used implementation for
each layers. Since accurate alignment with static alignment [38] 444Static
align employs a standard pattern-based approach: we select a part of a trace
as a reference, and compute correlation for each offset within a chosen range
for each of the traces. We then shift each trace by the respective offset that
maximizes the correlation. is not possible using all of the traces for the
second layer, 60% of the collected traces are not used in the CEMA attack for
the second layer.

In addition, despite collecting measurements with minimum misalignment, the
traces are hard to fully synchronize as they contain lot of jitter. To combat
this problem, we use elastic alignment [52] in the leakage detection process.
Elastic alignment allows us to align the traces at every time point, however,
it is also a technique that requires tuning the correct parameters and it is
computationally expensive to do so. Therefore, it is used only to align traces
to detect leaking points but the CEMA attack is carried out on the raw traces
which are aligned using static alignment.

(a) Key rank vs number of traces of the first weight in the first layer with
value of -1.22.

(b) Correlation vs number of traces of the first weight in the first layer
with value of -1.22.

(c) Key rank vs number of traces of the fifth weight in the first layer with
value of 0.3564.

(d) Correlation vs number of traces of the fifth weight in the first layer
with value of 0.3564.

Figure 8: Key ranks and correlations of the first and fifth weights in the
first layer.

###  3.6 Leakage detection

Detecting where intermediate values might leak is a necessary step before
applying CEMA. In order to do so, we apply random vs fixed TVLA [21] and
intermediate correlation to detect PoIs in both layers. We also used leakage
detection process to decide the best location for the probe in our setup.
Furthermore, we performed the leakage detection on the $relu\\_small$,
$relu\\_medium$ and $relu\\_interior$ implementations and the analysis
presented in this section applies to all three.

The TensorRT framework supports refitting networks on-the-fly with different
weights but a new CUDA context [2] has to be created every time the network is
refitted with new weights as per the documentation. Normally, an application
is not required to create a new CUDA context for every inference operation,
but it is required for TVLA as we refit the weights or the bias for every
inference operation.

To the detect PoIs, we applied TVLA 5 times, which is equal to the kernel
size, each time with one weight in the convolutional kernel being fixed vs
random and the rest of the weights 0. The corresponding first input, with
which the weight is multiplied, is fixed and non-zero while the rest of the
inputs are set to 0. This setup allows us to detect leakages corresponding to
the partial sums in the convolutional operation. In this case, the
unexploitable false positives can be loading of the weights. The presented
approach is more desired, i.e. leads to less false positives, then applying
TVLA with random vs fixed inputs as the same input is loaded at different
times, potentially leading to more false positives.

In addition, to detect leakage corresponding to the bias, we apply fixed vs.
random bias TVLA. Similarly to the random vs. fixed weight TVLA, the false
positives in this test can be the loading of the bias. To help filtering these
false positives, we also apply correlation to intermediate values. However, we
also observed several ghost peaks [14, 15] with the correlation method which
cannot be exploited in the CEMA attack.

Figure 5 shows the results of TVLA for the first weight in the kernel in the
first layer with 45k traces. We repeat this process for the rest of the
weights and the results clearly show that leakages of the convolution
operation are in the second highlighted part of Figure 7. Figure 6 shows the
results of TVLA for the bias of the kernel in the first layer with 37k traces.
There is a much clearer leakage present for the bias than for the weights in
the convolution operation.

After establishing the PoIs with the help of elastic alignment, we use static
alignment on the raw traces at these points as static alignment produced
higher individual TVLA peaks as well as correlation for the attack.

After the TVLA and correlations tests, it is possible to segment the raw
traces according the instruction blocks mentioned in Section 3.2. Figure 7
shows a raw trace where the first highlighted section corresponds to the
initialization block. The second highlighted block corresponds to the sum of
multiplications in the convolution operation. The third highlighted segment
corresponds to the calculation of the bias addition and ReLU output, repeated
4 times for different sets of registers. Note, that these instruction blocks
are also separated by synchronization instructions which are also visible in
the trace as the amplitude of the EM signal drops close to 0 between the
blocks. The same segmentation applies for the second layer.

(a) Key rank vs number of traces of the second weight in the second layer with
value of -0.5137.

(b) Correlation vs number of traces of the second weight in the second layer
with value of -0.5137.

(c) Key rank vs number of traces of the third weight in the second layer with
value of -0.6406.

(d) Correlation vs number of traces of the third weight in the second layer
with value of -0.6406.

Figure 9: Key rank and correlation for the third and fourth weights in the
kernel in the second layer.

###  3.7 Weight and bias extraction

If an adversary is able to recover a neural network’s architecture as well as
the parameters of a network, then the adversary completely recovered the
model. In this attack, we show how to extract the weights and the biases of a
CNN with 2 layers using CEMA. Observe that this attack is general and can be
extended to arbitrary number of convolutional layers with arbitrary
dimensions: since the implementations in Section 3.2 are parallelized on a
kernel and channel level, CEMA also can be parallelized to recover all the
weights in a layer if an adversary has multiple machines available.

The input and batch size of our CNN are set to 500 and 1, respectively, while
the weights and biases are initialized randomly. With an input size of 500 and
stride of 1, the number of convolutions executed by a layer is $input\\_size-
kernel\\_size+1$. Consequently, the first and second layers execute 496 and
492 convolutions, respectively. This means there are lots of convolutional
intermediates to target in each layer to recover the weights and bias of a
kernel. However, our attack only requires knowing just a fraction of the
inputs, and the size of this fraction is equivalent to the kernel size.
Therefore, we only use the first 5 inputs to each layer to recover the weights
and biases in the layers.555However, in order to calculate the first 5 inputs
to the second layer, the first 9 inputs to the first layer has to be known.

The layers in our example consist of one kernel each containing 5 unknown
weights and 1 unknown bias. The convolutional layers are followed by ReLU
activation. Altogether, 12 unknowns have to be recovered from both layers.

Formally, each layer computes

$\mathit{ReLU}(W*M+b)=\mathit{ReLU}(\sum_{i=1}^{5}w_{i}\cdot m_{i}+b),$

for a piece of input where $W$ and $M$ are vectors containing the kernel
weights and layer inputs, respectively, and $b$ is the bias that is added
after the sum of multiplications. In this case, there are multiple
intermediate values that depend on the secret weights. First, every weight
$w_{i}$ is multiplied by an input $m_{i}$ so the results of the individual
multiplications can be targeted. Second, the results of multiplications are
summed. Afterward, the bias $b$ is added to the accumulated results. Lastly,
the ReLU activation is applied after the bias addition. All of these
intermediate values depend on the unknown weights or the unknown bias. In the
attack, we have to recover of the parameters of the layers sequentially, i.e.
we have to recover the weights and bias in the first layer to be able to
attack the second layer.

####  3.7.1 Weight extraction

In this experiment, we target the intermediate results of the convolution
between the weights and the inputs, as we could observe results with this
leakage. We present results with two different leakage models as using the HW
leakage model on the partial sum intermediates and the HD between two
subsequent intermediate allows us to extract the weights.

Following Batina et al. [13], we restrict the search space of the weights to
the interval $[-5,5]$. With 16-bit floats, there are 35 330 possible
candidates in this range. We use the HW leakage model on our chosen
intermediate value, the partial sums, to map hypothetical EM consumption
values to real measurements $T$ and calculate the Pearson’s correlation per
sample point between them:

$\rho(HW(\sum_{j=1}^{i}w_{j}\cdot m_{j}),T)\quad i=1..5.$

We calculate the correlation with each weight candidate, selecting the
candidate that shows the highest correlation.

Figure 8 shows the key rankings and correlations of the correct weight value
vs the number of traces in the first layer for the first and fifth weights,
respectively. As shown in Figure 88(a), at around the 3,2 million mark, the
candidate that ranks first is $-1.212$ for the first weight, so it is already
a value that is close to the correct candidate. At this mark, the correct
candidate ranks 106th. Overall, the key rank decreases quickly but in order to
get the first rank for the correct candidate, more than 5 million traces are
required. This might suggest that the exponent of the intermediate values leak
mostly but not the mantissa, and it requires significantly more traces to get
the correct candidate. We leave analyzing this phenomenon for future work.

Figure 88(b) shows the correlation of the correct weight candidate and the
incorrect weight candidates vs the number of traces for the first weight. As
expected, it shows a similar trend to the key rank figure, as the correlation
of the correct candidate gets close to the top at the 3 million traces mark.
However, the correlation of the correct weight candidate only reaches the
maximum after 5 million traces with respect to the incorrect candidates.
Figure 88(c) shows the ranking of the correct weight value vs the number of
traces in the first layer for the fifth weight. Similarly to the first weight,
the key rank drops quickly, but it takes a lot more traces to for the key rank
to converge zero. Similarly, Figure 88(d) shows the correlation of the weight
candidates vs the number of traces for the fifth weight in the first layer.
The correct weight candidates already approaches key rank 0 at 8 million mark
while needing more traces to fully converge. Every weight in the kernel leaks
similarly and targeting the Hamming weight of partial sums can recover the
correct weight candidate for all the weight in the kernel. After recovering
the weights and bias of the kernel in the first layer, we can calculate the
inputs to the second layer to recover the weights and the bias of the kernel
in second kernel. Figure 9 shows the key ranking and correlations of the
correct weight value vs the number of traces in the second layer for the
second and third weights, respectively. Similarly to the weights in the first
layer, the key ranks for both weights drop quickly, but reaching a key rank of
0 requires 800 000 and 1.2 million traces for the second and third weights,
respectively.

####  3.7.2 Bias extraction

There are two operations in the layer that are dependent on the bias value.
The first operation is the calculation of the sum of convolution and the bias
$W*M+b$. The second one is the ReLU output, $ReLU(W*M+b)$. We observed leakage
with the Hamming weight of the sum of convolution and the bias, as well as
with the output of the ReLU activation. Formally, $HW(W*M+b)$ and
$HW(ReLU(W*M+b))$ both show correlation with the EM measurements. Therefore,
we calculate the Pearson correlation for every bias candidate with these
leakage models to the real measurements.

(a) Key rank vs number of traces of the bias in the first layer with the ReLU
intermediate.

(b) Correlation vs number of traces of the bias in the first layer with the
ReLU intermediate.

Figure 10: Key ranks and correlations of the bias with value of 0.856 in the
first layer with ReLU intermediate.

(a) Key rank vs number of traces of the bias in the second layer with the
$W*M+b$ intermediate.

(b) Correlation vs number of traces of the bias in the second layer with the
$W*M+b$ intermediate.

Figure 11: Key ranks and correlations of the bias with value of 0.345 in the
second layer with $W*M+b$ intermediate.

Figure 10 shows the key ranking and correlation of the CEMA attack on the bias
in the first layer with the $ReLU(W*M+b)$ intermediates. The key rank drops
rapidly, but the convergence to key rank 0 takes more than 10 million traces.
Similarly, Figure 10 shows the key rankings and correlation of the CEMA attack
on the bias in the second layer with the $W*M+b$ intermediate. The key rank
also drops quickly for the bias in the second layer, but the convergence to
key rank 0 takes more than 10 million traces. In general, the biases seem to
require more traces to converge to key rank 0 than the weights.

##  4 Discussion

###  4.1 Limitations

####  4.1.1 Parameter extraction

In this work, we demonstrated the parameter extraction of a NN with two
convolutional layers. However, there are multiple aspects that could impact
the CEMA attack:

  * Kernel size.

The kernel size does not impact the CEMA attack in terms of number of traces
needed, but it would take more time to recover all the weights, as our attack
targets the partial sums in the convolution where we recover the weights one-
by-one. In other words, the time it takes to recover all the weights in a
kernel scales linearly with the kernel size.

  * Batch size.

In our experiments, the batch size is set to 1. However, a larger batch size
might not impact the CEMA attack because the same operations could be carried
out in parallel with different inputs. This means that two or more
intermediate values are calculated at the same time that depend on the same
weight or bias. Therefore, the CEMA attack would be applied using the
combination of multiple intermediate values.

  * Concurrent applications.

The GPU is able to handle and schedule concurrently CUDA functions from
multiple applications. This could likely introduce noise in the attacks.
However, this highly depends on the required resources for each CUDA function.
If a CUDA function takes up most of the GPU resources (e.g. shared memory,
registers, etc.), then a different CUDA function will only be scheduled after
all the thread blocks in the previous CUDA function finished execution [3].

###  4.2 Single-precision implementations

The implementations for a convolutional layer with ReLU activation and single-
precision calculations are very similar to half-precision implementations
analyzed in this work. In fact, the single-precision implementations are named
the following:

$maxwell\\_scudnn\\_128x\\_\textbf{k}\\_relu\\_\textbf{t}\\_nn\\_v1$

where

$k\in\\{32,64,128\\},$ $t\in\\{small,medium,large,interior\\}.$

More importantly, the structures of these implementations are also identical
to their half-precision counterparts with the exception of using single-
precision instructions instead of half-precision instructions. Consequently,
our CEMA attack is also applicable for single-precision implementation.
However, targeting 32-bit floats significantly increases the search space even
if it is limited to the $[-5,5]$ interval.

###  4.3 Mitigation

Traditional ways to contain electromagnetic emanation, such as proper
shielding or introducing noise to decrease the Signal-to-Noise ratio, could
alleviate the problem [38]. Specifically against parameter extraction, one of
the possible countermeasures, which is also mentioned in the CSI-NN paper
[13], is shuffling [53] the order of multiplications in the layers, which can
make it significantly harder for an adversary to recover the weights.
Additionally, masking [18, 42], also mentioned in the CSI-NN paper, would
provide a way to break the relationship between the side-channel measurements
and the processed data. However, this comes at the price of execution speed,
which might not be desired in real-time systems.

###  4.4 Comparison with related work

To the best of our knowledge, no previous work has been able to extract the
parameters of neural networks on GPU using physical side-channel.

Author | Platform | Side channel
---|---|---
BarraCUDA (this work) | GPU | EM
Batina, et al. (2019) [13] | microcontroller | EM
Dubet, et al. (2020) [19] | FPGA | Power
Yoshida, et al. (2020) [57] | FPGA | Power
Regazzoni, et al. (2020) [46] | FPGA | EM
Yli-Mäyry, et al. (2021) [55] | FPGA | EM
Li, et al. (2022) [33] | FPGA | Power
Joud, et al. (2022) [26] | microcontroller | EM
Gongye et al. (2023) [22] | FPGA | EM
TABLE I: Comparison with related work.

Previous works have demonstrated parameter extraction on microcontrollers and
FPGAs using power or EM side channel as shown in Table I, but not on GPUs. In
addition, these attacks were performed on neural networks with binary
parameters [19, 46, 56], 8-bit parameters [13, 57, 32, 22] or 32-bit
parameters [13, 26]. In our work, we demonstrate parameter extraction on
16-bit parameters with discussion on extending it to 32-bit parameters.
Additionally, the approach presented in this work is scalable because it does
not rely on chosen inputs as in [22]. Chosen inputs are not scalable to large
neural networks as crafting chosen inputs becomes harder the deeper the target
NN is.

##  5 Conclusions

In this work, we analyze the NVIDIA Jetson Nano GPU, a commonly chosen
platform for real-world neural network implementations, for resilience against
side-channel attacks that aim to extract the weights of the target NN. First,
we establish where the GPU leaks information about the parameters of the
target DNN. Subsequently, we demonstrate the extraction of weights and biases
of a CNN consisting of two convolutional layers using CEMA. Overall, the
neural network implementations of NVIDIA’s TensorRT framework are vulnerable
to parameter extraction using electromagnetic side-channel attack despite the
networks running in a highly parallel and noisy environment. It remains an
open problem to protect their implementations in security or privacy-sensitive
applications.

## Acknowledgments

This work was supported by an ARC Discovery Early Career Researcher Award
DE200101577; an ARC Discovery Project number DP210102670; and the Deutsche
Forschungsgemeinschaft (DFG, German Research Foundation) under Germany’s
Excellence Strategy - EXC 2092 CASA - 390781972. Additionally, this work
received funding in the framework of the NWA Cybersecurity Call with project
name PROACT with project number NWA.1215.18.014, which is (partly) financed by
the Netherlands Organisation for Scientific Research (NWO). Moreover, Łukasz
Chmielewski was partially supported by Ai-SecTools (VJ02010010) project.

## References

  * cuc [a] “Cuda compute capabilities,” https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#compute-capabilities, accessed: 2022-09-30.
  * cuc [b] “Cuda context,” https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#context, accessed: 2022-09-30.
  * [3] https://developer.download.nvidia.com/CUDA/training/StreamsAndConcurrencyWebinar.pdf, accessed: 2022-11-30.
  * [4] “cuobjdump,” https://docs.nvidia.com/cuda/cuda-binary-utilities/#usage, accessed: 2022-09-30.
  * [5] “Cuda programming model,” https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#programming-model, accessed: 2022-09-30.
  * [6] “Cuda half2 data type,” https://docs.nvidia.com/cuda/cuda-math-api/struct____half2.html#struct____half2, accessed: 2022-09-30.
  * [7] https://www.langer-emv.de/en/product/mfa-active-1mhz-up-to-6-ghz/32/mfa-r-0-2-75-near-field-micro-probe-1-mhz-up-to-1-ghz/854, accessed: 2022-01-25.
  * [8] “Tensorrt model inspection,” https://docs.nvidia.com/deeplearning/tensorrt/archives/tensorrt-821/api/c_api/classnvinfer1_1_1_i_engine_inspector.html, accessed: 2022-09-30.
  * [9] “Nvidia jetson nano,” https://developer.nvidia.com/embedded/jetson-nano-developer-kit, accessed: 2022-09-30.
  * [10] “Nvidia tensorrt,” https://developer.nvidia.com/tensorrt, accessed: 2022-09-30.
  * tx [1] “Tegra x1 system-on-chip,” http://international.download.nvidia.com/pdf/tegra/Tegra-X1-whitepaper-v1.0.pdf, accessed: 2022-09-30.
  * Abadi et al. [2015] M. Abadi, A. Agarwal, P. Barham, E. Brevdo, Z. Chen, C. Citro, G. S. Corrado, A. Davis, J. Dean, M. Devin, S. Ghemawat, I. Goodfellow, A. Harp, G. Irving, M. Isard, Y. Jia, R. Jozefowicz, L. Kaiser, M. Kudlur, J. Levenberg, D. Mané, R. Monga, S. Moore, D. Murray, C. Olah, M. Schuster, J. Shlens, B. Steiner, I. Sutskever, K. Talwar, P. Tucker, V. Vanhoucke, V. Vasudevan, F. Viégas, O. Vinyals, P. Warden, M. Wattenberg, M. Wicke, Y. Yu, and X. Zheng, “TensorFlow: Large-scale machine learning on heterogeneous systems,” 2015, software available from tensorflow.org. [Online]. Available: https://www.tensorflow.org/
  * Batina et al. [2019] L. Batina, S. Bhasin, D. Jap, and S. Picek, “CSI–NN: Reverse engineering of neural network architectures through electromagnetic side channel,” in _28th USENIX Security Symposium USENIX Security 19)_ , 2019, pp. 515–532.
  * Brier et al. [2004] E. Brier, C. Clavier, and F. Olivier, “Correlation power analysis with a leakage model,” in _International workshop on cryptographic hardware and embedded systems_. Springer, 2004, pp. 16–29.
  * Canovas and Clédière [2005] C. Canovas and J. Clédière, “What do s-boxes say in differential side channel attacks?” _Cryptology ePrint Archive_ , 2005.
  * Chmielewski and Weissbart [2021] Ł. Chmielewski and L. Weissbart, “On reverse engineering neural network implementation on GPU,” in _International Conference on Applied Cryptography and Network Security_. Springer, 2021, pp. 96–113.
  * Chollet [2017] F. Chollet, “Xception: Deep learning with depthwise separable convolutions,” in _Proceedings of the IEEE conference on computer vision and pattern recognition_ , 2017, pp. 1251–1258.
  * Coron and Goubin [2000] J.-S. Coron and L. Goubin, “On boolean and arithmetic masking against differential power analysis,” in _Cryptographic Hardware and Embedded Systems—CHES 2000: Second International Workshop Worcester, MA, USA, August 17–18, 2000 Proceedings 2_. Springer, 2000, pp. 231–237.
  * Dubey et al. [2020] A. Dubey, R. Cammarota, and A. Aysu, “Maskednet: The first hardware inference engine aiming power side-channel protection,” in _IEEE International Symposium on Hardware Oriented Security and Trust_ , 2020, pp. 197–208.
  * Elibol et al. [2012] F. Elibol, U. Sarac, and I. Erer, “Realistic eavesdropping attacks on computer displays with low-cost and mobile receiver system,” in _2012 Proceedings of the 20th European Signal Processing Conference (EUSIPCO)_. IEEE, 2012, pp. 1767–1771.
  * Gilbert Goodwill et al. [2011] B. J. Gilbert Goodwill, J. Jaffe, P. Rohatgi _et al._ , “A testing methodology for side-channel resistance validation,” in _NIST non-invasive attack testing workshop_ , vol. 7, 2011, pp. 115–136.
  * Gongye et al. [2024] C. Gongye, Y. Luo, X. Xu, and Y. Fei, “Side-channel-assisted reverse-engineering of encrypted dnn hardware accelerator ip and attack surface exploration,” in _IEEE Symposium on Security and Privacy_ , 2024\.
  * He et al. [2016] K. He, X. Zhang, S. Ren, and J. Sun, “Deep residual learning for image recognition,” in _Proceedings of the IEEE conference on computer vision and pattern recognition_ , 2016, pp. 770–778.
  * Hongxin et al. [2009] Z. Hongxin, H. Yuewang, W. Jianxin, L. Yinghua, and Z. Jinling, “Recognition of electro-magnetic leakage information from computer radiation with SVM,” _Computers & Security_, vol. 28, no. 1-2, pp. 72–76, 2009.
  * Howard et al. [2017] A. G. Howard, M. Zhu, B. Chen, D. Kalenichenko, W. Wang, T. Weyand, M. Andreetto, and H. Adam, “Mobilenets: Efficient convolutional neural networks for mobile vision applications,” _arXiv preprint arXiv:1704.04861_ , 2017.
  * Joud et al. [2022] R. Joud, P.-A. Moëllic, S. Pontié, and J.-B. Rigaud, “A practical introduction to side-channel extraction of deep neural network parameters,” in _International Conference on Smart Card Research and Advanced Applications_. Springer, 2022, pp. 45–65.
  * Kocher et al. [1999] P. Kocher, J. Jaffe, and B. Jun, “Differential power analysis,” in _Annual international cryptology conference_. Springer, 1999, pp. 388–397.
  * Kocher [1996] P. C. Kocher, “Timing attacks on implementations of Diffie-Hellman, RSA, DSS, and other systems,” in _Annual International Cryptology Conference_. Springer, 1996, pp. 104–113.
  * Krizhevsky et al. [2012] A. Krizhevsky, I. Sutskever, and G. E. Hinton, “Imagenet classification with deep convolutional neural networks,” _Advances in neural information processing systems_ , vol. 25, pp. 1097–1105, 2012.
  * Kuhn and Anderson [1998] M. G. Kuhn and R. J. Anderson, “Soft tempest: Hidden data transmission using electromagnetic emanations,” in _International Workshop on Information Hiding_. Springer, 1998, pp. 124–142.
  * Laptev et al. [2017] N. Laptev, J. Yosinski, L. E. Li, and S. Smyl, “Time-series extreme event forecasting with neural networks at uber,” in _International conference on machine learning_ , vol. 34. sn, 2017, pp. 1–5.
  * Li et al. [2022b] G. Li, M. Tiwari, and M. Orshansky, “Power-based attacks on spatial DNN accelerators,” _ACM J. Emerg. Technol. Comput. Syst._ , vol. 18, no. 3, pp. 58:1–58:18, 2022. [Online]. Available: https://doi.org/10.1145/3491219
  * Li et al. [2022a] ——, “Power-based attacks on spatial dnn accelerators,” _ACM Journal on Emerging Technologies in Computing Systems_ , vol. 18, no. 3, pp. 1–18, 2022\.
  * Lin et al. [2013] M. Lin, Q. Chen, and S. Yan, “Network in network,” _arXiv preprint arXiv:1312.4400_ , 2013.
  * Liu et al. [2020a] L. Liu, W. Ouyang, X. Wang, P. Fieguth, J. Chen, X. Liu, and M. Pietikäinen, “Deep learning for generic object detection: A survey,” _International journal of computer vision_ , vol. 128, no. 2, pp. 261–318, 2020.
  * Liu et al. [2020b] Z. Liu, N. Samwel, L. Weissbart, Z. Zhao, D. Lauret, L. Batina, and M. Larson, “Screen gleaning: A screen reading tempest attack on mobile devices exploiting an electromagnetic side channel,” _arXiv preprint arXiv:2011.09877_ , 2020.
  * Maia et al. [2022] H. T. Maia, C. Xiao, D. Li, E. Grinspun, and C. Zheng, “Can one hear the shape of a neural network?: Snooping the GPU via magnetic side channel,” in _31st USENIX Security Symposium, USENIX Security 2022, Boston, MA, USA, August 10-12, 2022_ , K. R. B. Butler and K. Thomas, Eds. USENIX Association, 2022, pp. 4383–4400. [Online]. Available: https://www.usenix.org/conference/usenixsecurity22/presentation/maia
  * Mangard et al. [2008] S. Mangard, E. Oswald, and T. Popp, _Power analysis attacks: Revealing the secrets of smart cards_. Springer Science & Business Media, 2008, vol. 31.
  * Moradi et al. [2018] A. Moradi, B. Richter, T. Schneider, and F.-X. Standaert, “Leakage detection with the x2-test,” _IACR Transactions on Cryptographic Hardware and Embedded Systems_ , pp. 209–237, 2018.
  * Otter et al. [2020] D. W. Otter, J. R. Medina, and J. K. Kalita, “A survey of the usages of deep learning for natural language processing,” _IEEE transactions on neural networks and learning systems_ , vol. 32, no. 2, pp. 604–624, 2020.
  * Paszke et al. [2019] A. Paszke, S. Gross, F. Massa, A. Lerer, J. Bradbury, G. Chanan, T. Killeen, Z. Lin, N. Gimelshein, L. Antiga, A. Desmaison, A. Kopf, E. Yang, Z. DeVito, M. Raison, A. Tejani, S. Chilamkurthy, B. Steiner, L. Fang, J. Bai, and S. Chintala, “Pytorch: An imperative style, high-performance deep learning library,” in _Advances in Neural Information Processing Systems 32_. Curran Associates, Inc., 2019, pp. 8024–8035. [Online]. Available: http://papers.neurips.cc/paper/9015-pytorch-an-imperative-style-high-performance-deep-learning-library.pdf
  * Prouff and Rivain [2013] E. Prouff and M. Rivain, “Masking against side-channel attacks: A formal security proof,” in _Advances in Cryptology–EUROCRYPT 2013: 32nd Annual International Conference on the Theory and Applications of Cryptographic Techniques, Athens, Greece, May 26-30, 2013. Proceedings 32_. Springer, 2013, pp. 142–159.
  * Purwins et al. [2019] H. Purwins, B. Li, T. Virtanen, J. Schlüter, S.-Y. Chang, and T. Sainath, “Deep learning for audio signal processing,” _IEEE Journal of Selected Topics in Signal Processing_ , vol. 13, no. 2, pp. 206–219, 2019.
  * Quisquater and Samyde [2001] J.-J. Quisquater and D. Samyde, “ElectroMagnetic Analysis (EMA): Measures and Counter-Measures for Smard Cards,” in _Smart Card Programming and Security (E-smart 2001)_ , ser. Lecture Notes in Computer Science, I. Attali and T. P. Jensen, Eds., vol. 2140. Springer-Verlag, 2001, pp. 200–210.
  * Rangapuram et al. [2018] S. S. Rangapuram, M. W. Seeger, J. Gasthaus, L. Stella, Y. Wang, and T. Januschowski, “Deep state space models for time series forecasting,” _Advances in neural information processing systems_ , vol. 31, 2018.
  * Regazzoni et al. [2020] F. Regazzoni, S. Bhasin, A. A. Pour, I. Alshaer, F. Aydin, A. Aysu, V. Beroulle, G. Di Natale, P. Franzon, D. Hely _et al._ , “Machine learning and hardware security: Challenges and opportunities,” in _International Conference on Computer-Aided Design_ , 2020, pp. 1–6.
  * Sagheer and Kotb [2019] A. Sagheer and M. Kotb, “Time series forecasting of petroleum production using deep lstm recurrent networks,” _Neurocomputing_ , vol. 323, pp. 203–213, 2019.
  * Salinas et al. [2020] D. Salinas, V. Flunkert, J. Gasthaus, and T. Januschowski, “Deepar: Probabilistic forecasting with autoregressive recurrent networks,” _International Journal of Forecasting_ , vol. 36, no. 3, pp. 1181–1191, 2020\.
  * Schneider and Moradi [2015] T. Schneider and A. Moradi, “Leakage assessment methodology,” in _International Workshop on Cryptographic Hardware and Embedded Systems_. Springer, 2015, pp. 495–513.
  * Silver et al. [2017] D. Silver, T. Hubert, J. Schrittwieser, I. Antonoglou, M. Lai, A. Guez, M. Lanctot, L. Sifre, D. Kumaran, T. Graepel _et al._ , “Mastering chess and shogi by self-play with a general reinforcement learning algorithm,” _arXiv preprint arXiv:1712.01815_ , 2017.
  * Simonyan and Zisserman [2014] K. Simonyan and A. Zisserman, “Very deep convolutional networks for large-scale image recognition,” _arXiv preprint arXiv:1409.1556_ , 2014.
  * van Woudenberg et al. [2011] J. G. van Woudenberg, M. F. Witteman, and B. Bakker, “Improving differential power analysis by elastic alignment,” in _Topics in Cryptology–CT-RSA 2011: The Cryptographers’ Track at the RSA Conference 2011, San Francisco, CA, USA, February 14-18, 2011. Proceedings_. Springer, 2011, pp. 104–119.
  * Veyrat-Charvillon et al. [2012] N. Veyrat-Charvillon, M. Medwed, S. Kerckhof, and F.-X. Standaert, “Shuffling against side-channel attacks: A comprehensive study with cautionary note,” in _Advances in Cryptology–ASIACRYPT 2012: 18th International Conference on the Theory and Application of Cryptology and Information Security, Beijing, China, December 2-6, 2012. Proceedings 18_. Springer, 2012, pp. 740–757.
  * Yan et al. [2020] M. Yan, C. W. Fletcher, and J. Torrellas, “Cache telepathy: Leveraging shared resource attacks to learn DNN architectures,” in _USENIX Security_ , 2020, pp. 2003–2020.
  * Yli-Mäyry et al. [2021] V. Yli-Mäyry, A. Ito, N. Homma, S. Bhasin, and D. Jap, “Extraction of binarized neural network architecture and secret parameters using side-channel information,” in _IEEE International Symposium on Circuits and Systems, ISCAS 2021, Daegu, South Korea, May 22-28, 2021_. IEEE, 2021, pp. 1–5. [Online]. Available: https://doi.org/10.1109/ISCAS51556.2021.9401626
  * Yli-Mäyry et al. [2021] V. Yli-Mäyry, A. Ito, N. Homma, S. Bhasin, and D. Jap, “Extraction of binarized neural network architecture and secret parameters using side-channel information,” in _IEEE International Symposium on Circuits and Systems (ISCAS)_ , 2021, pp. 1–5.
  * Yoshida et al. [2020] K. Yoshida, T. Kubota, S. Okura, M. Shiozaki, and T. Fujino, “Model reverse-engineering attack using correlation power analysis against systolic array based neural network accelerator,” in _IEEE International Symposium on Circuits and Systems (ISCAS)_ , 2020, pp. 1–5.


# Optical dynamic nuclear polarization of 13C spins in diamond at a low field
with multi-tone microwave irradiation

Vladimir Vladimirovish Kavtanyuk  Quantum Magnetic Imaging Team, Korea
Research Institute of Standards and Science, Daejeon 34113, Republic of Korea
Hyun Joon Lee  Radio & Satellite Research Division, Electronics and
Telecommunications Research Institute, Daejeon 34129, Republic of Korea
Sangwon Oh  Quantum Magnetic Imaging Team, Korea Research Institute of
Standards and Science, Daejeon 34113, Republic of Korea  Keunhong Jeong  Korea
Military Academy, Seoul 01805, Republic of Korea  Jeong Hyun Shim
<EMAIL_ADDRESS>Quantum Magnetic Imaging Team, Korea Research Institute of
Standards and Science, Daejeon 34113, Republic of Korea  Department of Medical
Physics, University of Science and Technology, Daejeon 34113, Republic of
Korea

###### Abstract

Most of dynamic nuclear polarization (DNP) has been requiring helium
cryogenics and strong magnetic fields for a high degree of polarization. In
this work, we instead demonstrate an optical hyperpolarization of naturally
abundant 13C nuclei in a diamond crystal at a low magnetic field and an
ambient temperature. It exploits continuous irradiations of pump laser for
polarizing electron spins of nitrogen vacancy centers and microwave for
transferring the induced polarization to 13C nuclear spins. Triplet structures
corresponding to 14N hyperfine splitting were clearly observed in the spectrum
of 13C polarization. The powers of microwave irradiation and pump laser were
optimized. By simultaneously irradiating three microwave frequencies matching
to the peaks of the triplet, we achieved a 13C bulk polarization of 0.113
$\%$, leading to an enhancement of about a factor of 90,000 over the thermal
polarization at 17.6 mT. We believe that the multi-tone irradiation can be
universally adopted to further enhance the 13C polarization at a low magnetic
field.

††preprint: APS/123-QED

##  I INTRODUCTION

Dynamic nuclear polarization (DNP) has been a technological breakthrough,
which can significantly boost the signal-to-noise ratio in nuclear magnetic
resonance (NMR) and magnetic resonance imaging (MRI) applications [1, 2, 3, 4,
5, 6, 7]. Enhanced signal enables extracting substantial amount of information
at molecular levels for a wide-range of chemical, biological and physical
processes [8, 9, 10, 11, 12, 13, 14, 15, 16]. Recently, optical dynamic
nuclear polarization, in which electronic polarization of negatively-charged
nitrogen-vacancy (NV) centers is transferred to bulk 13C nuclear spins (Fig.
1(a)), has been demonstrated [17, 18, 19, 20, 21, 22, 23, 24, 25]. Besides
continuous-wave and simultaneous irradiations of microwave and pump laser, a
variety of NV-based optical DNP techniques have been introduced including
pulsed irradiation of microwave (MW) fields satisfying Hartman-Hahn condition
[26], the cross polarization at the level anti-crossing in the ground state
[17] and the excited state [18]. Frequency sweeping [21, 22, 23, 24] was
exploited to hyperpolarize diamond powders or nanodiamonds that can
potentially be used for molecule-targeted in-vivo imaging with the advantage
of a long spin life time [27, 28, 29, 30, 31].

Owing to the optically induced polarization, NV-based optical DNP does not
need cryogenic apparatus and can operate in low magnetic fields of tens of
millitesla. The magnetic fields in this range are easily accessible in
experiments, however certainly not favorable for efficient hyperpolarization.
Firstly, both a weak thermal nuclear polarization and a short T1 relaxation
time are disadvantageous. In addition, the positive and negative nuclear
polarization spectra, being the hallmark of solid-state dynamic nuclear
polarization, are separated by the order of nuclear Zeeman splitting[32].
Thus, in low fields the partial overlap of positive and negative nuclear
polarizations may occur, and this will eventually reduce the net nuclear
polarization. Nevertheless, such nuclear polarization spectrum, which reflects
the weak 13C Zeeman splitting in low magnetic fields, has not been observed
yet in optical DNP studies of diamonds. If 13C resonance frequency is less
than 14N hyperfine splitting (2.16 MHz) of NV spin (Fig. 1(b)), the nuclear
polarization spectrum would be duplicated three times. Not surprisingly such
triplet structure has not been reported either.

Figure 1:  (a) The schematic illustration of optical DNP process in diamond.
The optically induced polarization in NV spins will be transferred to bulk 13C
nuclear spins. (b) The part of energy level structure associated with NV, 14N,
and distant 13C spins is given. At a low field, 13C nuclear Zeeman energy
$\gamma_{c}B$ is weaker than the 14N hyperfine splitting (2.16 MHz) in NV
spins. The arrows indicates the possible transitions induced by MW
irradiations used in optical DNP.

In the present study, we conducted NV-based optical DNP of bulk 13C nuclear
spins in diamond at a low field and an ambient temperature. Applied field was
17.6 mT, which corresponds to 13C resonance frequency of approximately 0.2
MHz. MW and laser were continuously irradiated, and their optimal power levels
were characterized. In 13C nuclear polarization spectrum, we clearly observed
the triplet structure showing the 14N hyperfine splitting. As anticipated,
nitrogen nuclear spin does not participate in the polarization transfer
process except for in determining the resonance frequency of NV spins through
the hyperfine interaction. By exploiting such independence, we irradiated
triple microwave frequencies simultaneously for further boosting, and achieved
a 13C nuclear polarization of 0.113 $\%$, which is about 1.7 times higher
compared to DNP by single MW frequency irradiation.

##  II Experimental methods

We used a HPHT-grown diamond crystal with naturally abundant 13C nuclear
spins. NV centers with a concentration of 1.25 ppm are fabricated via electron
irradiation and thermal annealing process. We developed an optical DNP system,
which includes low-field region for optical pumping, high-field region (6 T)
for 13C NMR, and a rapid shuttling device for displacing the hyperpolarized
diamond crystal to the high-field region in 2 s. From 13C NMR signals at 6 T,
enhancement factors and absolute polarization levels were estimated. 13C NMR
and the timing controls were conducted with a commercial NMR console. (For
more detailed explanations, see the appendix section.)

##  III  Experimental Results

###  III.1 13C polarization spectrum

We first investigated 13C polarization spectrum. Varying the frequency of MW
irradiated for optical DNP, we recorded the 13C NMR signal as represented by
the blue lines in Fig. 2. The green curves illustrate optically-detected
magnetic resonance (ODMR) lines of NV center, including the transitions of ms
= 0 $\rightarrow$ ms = -1 (Fig. 2(a)) and ms = 0 $\rightarrow$ ms = +1 (Fig.
2(b)). All these measurements were performed through NV spins of
$\langle$111$\rangle$ orientation, which is parallel with the direction of
magnetic field. Each data point was averaged by 10 measurements. The recording
sequence of 13C NMR started with saturating 13C spins with a series of 90
pulses, which destroyed a residual 13C polarization and ensured non-
polarization. Triplet structures were clearly observed for both NV spin
transitions. The 14N (I = 1) of the NV center splits the NV spin levels each
into three hyperfine sublevels with an energy splitting of 2.16 MHz, and this
explains the triplet shown in Fig. 2. Although ODMR does not exhibit the 14N
hyperfine splitting, it is visible in the 13C polarization spectra. Notably,
the signs of the 13C polarization are identical for both NV spin transitions.

Figure 2:  Normalized values of integrated Furrier Transform (FT) signals of
the real part plotted as a function of the MW frequency (blue line) and ODMR
signal with normalized intensity (green line) at 17.6 mT. Each data point for
each MW frequency was measured with the same experimental conditions. (a) -1
state and (b) +1 state of NV spins.

The results in Fig. 2 are obtained at a magnetic field of 17.6 mT with a
diamond containing natural abundance of 13C nuclear spins. A similar
measurement was performed in Ref.[19], where a magnetic field of 18 mT was
applied. However, the 13C polarization spectrum showed no triplet structure.
Since a diamond of 10$\%$ enriched 13C was used in Ref.[19], strong hyperfine
interactions between NV and proximate 13C spins may dominate over the
hyperfine interaction with the 14N spin. In Ref.[25], a higher magnetic field
of 473 mT was applied, corresponding to a 13C resonance frequency of 5 MHz. In
this case, the overlap of three spectra with a separation of 2.16 MHz wipes
the triplet structure, resulting in a single curve.

Figure 3:  (a) Illustration of MW irradiation influences on nuclear
polarization of 13C. (b) Normalized integrals of FT signals from
hyperpolarized 13C vs MW power.

###  III.2  Optimal microwave power

To improve the 13C polarization, we tried to optimize experimental parameters,
such as MW and laser powers. Particularly at a low magnetic field, a strong MW
field may not be beneficial. As the separation between positive and negative
peaks in the nuclear polarization spectrum is proportional to the nuclear
Zeeman splitting [32], the strong MW that is on-resonant with one polarization
can also induce the transition of the opposite polarization, which is off-
resonant as illustrated in Fig. 3 (a). Then, the opposite polarization will be
imposed and eventually lead to a lower net nuclear polarization. Figure 3 (b)
shows the 13C nuclear polarization as a function of MW power. The polarization
rises initially but decays afterwards as the MW power increases. From the
curve in Fig. 3, we can determine the optimal MW power, near 10 W. The solid
line is the fitted curve with two exponents for a guide to eyes. A similar
result has been reported in Ref.[19], in which it is explained by a transition
from selective regime to $\Lambda$ regime due to stronger MW fields.

###  III.3  Optimal laser power density

We investigated how the 13C polarization varies with the pump laser power.
Figure 4 (a) shows that the maximum 13C polarization signal was measured at a
laser power density of 30 mW/mm2. After the maximum, the polarization was
gradually reduced as laser power increased. Initially, we speculated that the
13C polarization decrease due to the rise of diamond temperature. The
temperature of the diamond crystal was measured from the position of zero-
field splitting. From the ODMR spectrum, the peak positions of the two
transitions, ms = 0 $\rightarrow$ ms = -1 and ms = 0 $\rightarrow$ ms = +1,
can be obtained. Then, their mean results in the zero-field splitting with the
conversion factor, $\frac{dD}{dT}\cong-74$ kHz/K[33]. As shown in Fig. 4 (b),
the temperature rise induced by the pump laser was linearly proportional to
the laser power density. At the optimal near 30 mW/mm2, the diamond
temperature was raised more than 100 K. Temperature rise will reduce the
thermal polarization of 13C nuclear spins. One may attribute the decrease of
the 13C polarization to such laser-induced heating.

Figure 4: (a) Normalized integrals of FT signals from hyperpolarized 13C vs
laser power density. (b) Laser power vs diamond temperature. (c) Laser power
vs intensity of diamond’s photoluminescence. (d) The curve of the 13C
polarization expected by using the results in Fig. 4 (b) and (c).

We found out, however, that the laser-induced heating is not enough to explain
the 13C polarization decrease. It’s because the NV electronic spin
polarization increases with the laser power. Figure 4 (c) shows the variation
of the fluorescence intensity from NV spins in the diamond. Apparently, the NV
spin polarization was far from the saturation and still increasing linearly.
The enhancement factor is proportional to the NV spin polarization
($P_{\mathrm{NV}}$). Then, 13C polarization ($P_{\mathrm{Hyper}}$) obtained by
the optical DNP becomes proportional to the NV spin polarization
($P_{\mathrm{NV}}$) times the 13C thermal polarization
($P_{\mathrm{Thermal}}$) under the pump laser. Given a pump laser density
$\sigma_{\mathrm{Laser}}$, $P_{\mathrm{NV}}$ and $P_{\mathrm{Thermal}}$ can be
expressed as below,

$\displaystyle P_{\mathrm{NV}}=\beta\sigma_{\mathrm{Laser}},$ $\displaystyle
P_{\mathrm{Thermal}}=\frac{c}{300+\alpha\sigma_{\mathrm{Laser}}}$ (1)

, in which $\alpha$ and $\beta$ are the slopes of the curves in Fig. 4 (b) and
(c), respectively. ($c$ is a constant.) The resulting curve of
$P_{\mathrm{Hyper}}$ is shown in Fig. 4 (d), which is certainly inconsistent
with the result in Fig. 4 (a).

###  III.4  13C hyperpolarization by multi-tone MW frequencies

Since multiple peaks were observed in the 13C polarization spectrum (Fig. 2),
we can enhance further the 13C polarization by simultaneous irradiation of
multiple MW frequencies corresponding to those peaks. To estimate the 13C
polarization obtained through the optical DNP, we compared the 13C NMR spectra
with that of thermal polarization at 6 T (300 K). The thermal polarization of
13C in the diamond was measured for 4 days, with 400 averages and 20 min
between each scan. The hyperpolarization of 13C in the diamond was performed
by optical DNP using the transition from ms = 0 to ms = -1 shown in Fig. 2
(a). The diamond was irradiated for 120 s by the pump laser beam with a power
density of 25 mW/mm2 and by the MW field with a power of 11 W. Due to the
laser irradiation the diamond was heated up to about 400 K.

Figure 5 shows the results of the multiple MW irradiation. The number of the
MW frequencies have been increased from one (f1) to three (f1, f2, f3). The
three MW frequencies are indicated by the arrows in Fig. 2 (a). Clearly, the
13C polarization increases with the number of MW frequency. It was 0.068 %
with f1 only, while 0.113 % with all three frequencies. The enhancement ratio
is 1.7, which is less than the ideal value of about 2.5 predicted by the
summed intensities in the Fig. 2 (a). This may be due to the effect of off-
resonant excitations described above. The polarization of 0.113 % corresponds
to the enhancement of 90,000 times over the in situ thermal polarization at
17.6 mT and the diamond temperature as the calculation shown in the appendix.

Figure 5:  Comparison between thermally polarized 13C at 6 T and
hyperpolarized 13C nuclei obtained by DNP method for -1 state of aligned NV-
at 17.6 mT. All the measurements were performed on the same diamond at the
room temperature. The pink line represents FT signal obtained from thermally
polarized 13C nuclei with 400 averages. The red line represents FT signal
obtained by MW excitation at the first hyperfine splitting corresponding to
0.068$\%$ of 13C hyperpolarization. The green line represents FT signal
obtained by MW excitation at the first and the second hyperfine splitting
corresponding to 0.094$\%$ of 13C hyperpolarization. The blue line represents
FT signal obtained by MW excitation at the first, the second and the third
hyperfine splitting corresponding to 0.113$\%$ of 13C hyperpolarization. The
combined polarization was expected to be higher, we are not sure the reason of
the lower polarization value.

##  IV Conclusion

We have demonstrated the optical DNP results obtained with a diamond crystal
of natural abundance of 13C nuclei at room temperature and a low magnetic
field of 17.6 mT. We observed a triplet structure in 13C nuclear polarization
spectrum, in which the splitting of 2.16 MHz attributes to 14N hyperfine
interaction in NV spin. The bulk 13C nuclear spin polarization reached 0.113%.
To obtain such 13C hyperpolarization, we have found optimal values of laser
power density and MW power. In addition, we adopted multi-tone MW excitation.
Simultaneous irradiation of three MW frequencies, matched to the peak
frequencies in the 13C nuclear spectrum, results in an improvement of 1.7
times. We believe this multi-tone irradiation scheme can be universally
exploited to further enhance the 13C polarization at a low magnetic field.

###### Acknowledgements.

This work was supported by a grant (GP2021-0010) from Korea Research Institute
of Standards and Science, Institute of Information and communications
Technology Planning $\&$ Evaluation (IITP) grants funded by the Korea
government (MSIT) (No.2019-000296, No.2019-0-00007), the Ministry of Science
and ICT in Korea via KBSI (Grant No. C123000).

##  Appendix A Diamond sample

All the experimental data, provided in this paper, have been measured from
only one industrial diamond of 42 mg. It was synthesized by the high-pressure
high-temperature method and was electron irradiated at 1 MeV followed by
annealing at about 800 oC. It contains about 50 ppm of P1 centers, 1.25 ppm of
NV centers and the natural abundance of 13C. The spin concentrations in the
diamond were estimated by performing electron spin resonance (ESR). The peak
intensities in ESR spectra were compared to that from a diamond crystal having
a known P1 concentration. The thickness of the diamond is about 1 mm and its
form can be considered as the equilateral triangle with side equal about 4 mm.
The diamond is $\langle$111$\rangle$ surface-oriented, meaning that if a
magnetic field is aligned along one of the NV center orientations, the other
three orientations would have the angle of 109.5o away from the field
direction.

##  Appendix B Estimation of 13C nuclear hyperpolarization and enhancement
factor

The hyperpolarization of 13C in the diamond obtained at $B_{\mathrm{EM}}$ =
17.6 mT has been estimated based on the thermal polarization of the same
diamond obtained at $B_{SM}$ = 6 T. So, PHyper \- the hyperpolarization of 13C
has been calculated by the equation:

$P_{\mathrm{Hyper}}=\frac{S_{\mathrm{Hyper}}}{S_{\mathrm{Thermal}}}*P_{\mathrm{Thermal}}$
(2)

where $S_{\mathrm{Hyper}}$ is the integral of FT signal measured from a
hyperpolarized 13C nuclei in the diamond, $S_{\mathrm{Thermal}}$ is the
integral of FT signal measured from thermally polarized 13C nuclei in the
diamond and $P_{Thermal}$ = $\hbar$$\gamma$$B_{0}/2kT_{R}$ is the value of
thermal polarization at the room temperature ($T_{R}$ = 297 K).

The enhancement factor ($\varepsilon$) of 13C hyperpolarization is about
90,000 over thermal equilibrium. It has been calculated by this equation:

$\varepsilon=\frac{S_{\mathrm{Hyper}}}{S_{\mathrm{Thermal}}}*\frac{B_{\mathrm{SM}}}{B_{\mathrm{EM}}}*\frac{T_{L}}{T_{R}}$
(3)

where $T_{L}$ is diamond temperature under laser irradiation during DNP
process (Fig. 4, b).

##  Appendix C Experimental Setup

Our experimental setup is sophisticated system capable of hyperpolarizing 13C
in diamonds using DNP method and measuring polarization using NMR method. This
section will thoroughly describe the experimental setup (Fig. 6) and a purpose
of each hardware.

Figure 6: Full experimental setup. The right part shows schematic of our
entire experimental setup. The left part shows small items of the setup in
details.

###  C.1 Superconducting magnet and NMR hardware

The NMR signals, presented in this paper, have been measured at the magnetic
field of a superconducting magnet from OXFORD. It has a diameter of 70 cm and
a height of 97 cm. The magnet produces a homogeneous magnetic field of about 6
T in its central region. There is a cylindrical cavity with a diameter of
about 5.5 cm in the center of the magnet through its entire height. A handmade
saddle coil is installed inside of the cavity in the homogeneous region for
measuring polarization from samples. The coil is connected to a NMR probe
(Fig. 6), which is a LC circuit made of parallel and serial connected variable
capacitors. By manipulating the capacitance, we can match the impendence and
set a required resonance frequency for the saddle coil ($\sim$64.237 MHz for
13C). The probe is subsequently connected to an active-transcoupler and to a
commercial spectrometer LapNMR from Tecmag. The purpose of the transcoupler is
to manage transmitted signals from the spectrometer and received NMR signals
from the saddle coil. The spectrometer is controlled by TNMR software, which
allows us to set all required parameters for generating proper RF signals and
acquire NMR signals. Transmitted RF signals from LapNMR is increased by CPC
9T100M power amplifier. NMR signals captured by the saddle coil is increased
by two +35 dB preamplifiers from NMR service.

###  C.2  MW hardware

One of the aspect in successful hyperpolarization of 13C in a diamond is
application of MW field with enough power. In our experiment, MW signals are
produced by APSIN12G and amplified by Sungsan TS2458BBP0 with a maximum power
output of 50 W. A handmade Helmholtz coil with a diameter of 13 mm is used to
generate MW fields (Fig. 6). Due to impedance mismatch between the coil and
other 50 Ohm hardware, most of applied power to the coil gets reflected. So,
an actual power which can pass through the MW coil is less than 1$\%$ from the
applied power. For this reason PE83CR1004 circulator with a 50 W termination
is used to protect the amplifier from the high reflection power. The MW coil
is installed inside of an electrical Helmholtz magnet (Fig. 6) powered by a
high precision DC power supply 2280S-32-6. The electrical magnet is installed
right below the superconducting magnet and used for DNP measurements, since it
provides a good homogeneity and the easy way to vary the magnetic field.

###  C.3 Laser with optics

Our laser produces vertically polarized 532 nm green light with maximum power
output of 4 W and beam diameter of 2 mm. The delivery of the beam light from
the laser to the diamond is performed by two separated optical setups (Fig.
7). The left setup locates inside of a shielding box where we perform DNP
measurements on the diamond. The box shields from electrical and magnetic
fields. The right setup locates outside of the box and includes the laser.

Figure 7:  Schematic of the laser optical setup.

The isolator is used to protect from the reflected light and to increase the
beam size to a diameter of 5 mm to match the diamond size. The long pass
filter is used to exclude lights except the red light emitting by the diamond.
The lens, with focal length of 25.4 mm, is used to focus the emitted
fluorescence, from diamond, to the photodetector’s sensor with a diameter of
about 2 mm.

##  Appendix D Experimental procedure

Before each DNP experiment, ODMR measurement was performed to find an exact
positions of ms = $\pm$1 substates at a specific magnetic field. The diamond
was irradiated with a wide range of MW sweep (up to 1.5 GHz) and with 532 nm
laser light. The emitted red light from a diamond then was captured by a
photo-detector. A data from the detector could be seen on a computer using a
Labview program, where we could define exact MW frequencies corresponding to
the ms = $\pm$1 substates and estimate a magnetic field. In this work we have
only showed DNP measurements performed on NV centers aligned with the
direction of the magnetic field. It has been problematic to align the diamond
precisely, so that all 3 non-aligned NV orientations would have the same angle
of 109.5o away from the direction of the magnetic field. So, we plan to
improve our setup later and perform DNP through the non-aligned NV
orientations, which will help us to increase the maximum hyperpolarization
value by three times. All DNP experiments in this work have been performed at
the magnetic field of about 17.6 mT, generated by the electrical magnet, which
corresponds to a MW frequency of 2.370$\pm$0.005GHz for ms = -1 substate and
3.357$\pm$0.005GHz for ms = +1 substate of NV centers aligned with the
magnetic field. 13C nuclear spins in the diamond was hyperpolarized for 100+
sec at constant MW irradiation and at constant laser irradiation with
wavelength of 532 nm. Once 13C nuclei were hyperpolarized, the diamond was
shuttled 80 cm up within 1 sec, by a high precision motion controller Newport
XPS-RL, to the superconducting magnet, where the polarization was detected by
the saddle coil. Then the signal, from the coil, was sent by LapNMR at the
resonance frequency of 64.237MHz. The NMR measurement took 5 ms and then
diamond was sent back to the previous position, within 2 sec, for additional
polarization measurements. For a better SNR, the hyperpolarized 13C NMR signal
was always averaged 10 times.

## References

  * Ardenkjær-Larsen  _et al._ [2003] J. H. Ardenkjær-Larsen, B. Fridlund, A. Gram, G. Hansson, L. Hansson, M. H. Lerche, R. Servin, M. Thaning, and K. Golman, Proc Natl Acad Sci USA 100, 10158 (2003).
  * Joo  _et al._ [2006] C.-G. Joo, K.-N. Hu, J. A. Bryant, and R. G. Griffin, J. Am. Chem. Soc. 128, 9428–9432 (2006).
  * Thankamony  _et al._ [2017] A. S. L. Thankamony, J. J.Wittmann, M. Kaushik, and B. Corzilius, Prog. in Nuc. Magn. Res. Spec. 102-103, 120–195 (2017).
  * Rankin  _et al._ [2019] A. G. M. Rankin, J. Trebosc, F. Pourpoint, J.-P. Amoureux, and O. Lafon, Solid State Nuc. Magn. Res. 101, 116–143 (2019).
  * Bucher  _et al._ [2020] D. B. Bucher, D. R. Glenn, H. Park, M. D. Lukin, and R. L. Walsworth, Phys. Rev. X 10, 021053 (2020).
  * Krummenacker  _et al._ [2012] J. G. Krummenacker, V. P. Denysenkov, M. Terekhov, L. M. Schreiber, and T. F. Prisner, J. of Mag. Res. 215, 94–99 (2012).
  * Denysenkov  _et al._ [2017] V. Denysenkov, M. Terekhov, R. Maeder, S. Fischer, S. Zangos, T. Vogl, and T. F. Prisner, Scientific Reports 7, 44010 (2017).
  * Golman, Zandt, and Thaning [2006] K. Golman, R. I. T. Zandt, and M. Thaning, Proc. Natl. Acad. Sci. 103, 11270–11275 (2006).
  * Barnes  _et al._ [2008] A. B. Barnes, G. D. Paepe, P. C. A. van der Wel, K.-N. Hu, C.-G. Joo, V. S. Bajaj, M. L. Mak-Jurkauskas, J. R. Sirigiri, J. Herzfeld, R. J. Temkin, and R. G. Griffin, App. Magn. Res. 34, 237–263 (2008).
  * Rossini  _et al._ [2012] A. J. Rossini, A. Zagdoun, F. Hegner, M. Schwarzwalder, D. Gajan, C. Coperet, A. Lesage, and L. Emsley, J. Am. Chem. Soc. 134, 16899–16908 (2012).
  * Puebla  _et al._ [2013] J. Puebla, E. A. Chekhovich, M. Hopkinson, P. Senellart, A. Lemaitre, M. S. Skolnick, and A. I. Tartakovskii, Phys. Rev. B 88 (4), 9 (2013).
  * Masion  _et al._ [2017] A. Masion, A. Alexandre, F. Ziarelli, S. Viel, and G. M. Santos, Scientific Reports 7, 3430 (2017).
  * Flori  _et al._ [2018] A. Flori, G. Giovannetti, M. F. Santarelli, G. D. Aquaro, D. D. Marchi, S. Burchielli, F. Frijia, V. Positano, L. Landini, and L. Menichetti, Spectr. Acta Part A: Molec. and Biomolec. Spectr. 199, 153–160 (2018).
  * Morze and Merritt [2019] C. V. Morze and M. E. Merritt, NMR in Biomedicine 32, e3937 (2019).
  * Park  _et al._ [2021] S. Park, H. Rintaro, S. K. Kim, and I. Park, Metabolites 11, 504 (2021).
  * Abhyankar and Szalai [2021] N. Abhyankar and V. Szalai, The J. of Phys. Chem. B 125, 5171–5190 (2021).
  * Fischer  _et al._ [2013] R. Fischer, C. O. Bretschneider, P. London, D. Budker, D. Gershoni, and L. Frydman, Phys. Rev. Lett. 111, 057601 (2013).
  * Wang  _et al._ [2013] H.-J. Wang, C. S. Shin, C. E. Avalos, S. J. Seltzer, D. Budker, A. Pines, and V. S. Bajaj, Nat. Comm. 4, 1940 (2013).
  * Alvarez  _et al._ [2015] G. A. Alvarez, C. O. Bretschneider, R. Fisher, P. London, H. Kanda, S. Onoda, J. Isoya, D. Gershoni, and L. Frydman, Nat. Comm. 6, 8456 (2015).
  * King  _et al._ [2013] J. P. King, K. Jeong, C. C. Vassiliou, C. S. Shin, R. H. Page, C. E. Avalos, H.-J. Wang, and A. Pines, Nat. Comm. 6, 8965 (2013).
  * Ajoy  _et al._ [2018a] A. Ajoy, K. Liu, R. Nazaryan, X. Lv, P. R. Zangara, B. Safvati, G. Wang, D. Arnold, G. Li, A. Lin, P. Raghavan, E. Druga, S. Dhomkar, D. Pagliero, J. A. Reimer, D. Suter, C. A. Meriles, and A. Pines, Sci. Adv. 4, eaar5492 (2018a).
  * Ajoy  _et al._ [2018b] A. Ajoy, R. Nazaryan, K. Liu, X. Lv, B. Safvati, G. Wang, E. Druga, J. A. Reimer, D. Suter, C. Ramanathan, C. A. Meriles, and A. Pines, Proc Natl Acad Sci USA 115, 10576 (2018b).
  * Ajoy  _et al._ [2020] A. Ajoy, R. Nazaryan, E. Druga, K. Liu, A. Aguilar, B. Han, M. Gierth, J. T. Oon, B. Safvati, R. Tsang, J. H. Walton, D. Suter, C. A. Meriles, J. A. Reimer, and A. Pines, Rev. of Sci. Instr. 91, 023106 (2020).
  * Ajoy  _et al._ [2021] A. Ajoy, A. Sarkar, E. Druga, P. Zangara, D. Pagliero, C. A. Meriles, and J. A. Reimer, “Low-field microwave-mediated optical hyperpolarization in optically pumped diamond,” Journal of Magnetic Resonance 331, 107021 (2021).
  * Parker  _et al._ [2019] A. J. Parker, K. Jeong, C. E. Avalos, B. J. M. Hausmann, C. C. Vassiliou, A. Pines, and J. P. King, Phys. Rev. B 100, 041203(R) (2019).
  * Scheuer  _et al._ [2016] J. Scheuer, I. Schwartz, Q. Chen, D. Schulze-Sünninghausen, P. Carl, P. Hofer, A. Retzker, H. Sumiya, J. Isoya, B. Luy, M. B. Plenio, B. Naydenov, and F. Jelezko, New J. Phys. 18, 013040 (2016).
  * Cassidy  _et al._ [2013] M. C. Cassidy, H. R. Chan, B. D. Ross, P. K. Bhattacharya, and C. M. Marcus, Nature Nanotechnology 8, 363–368 (2013).
  * Rej  _et al._ [2015] E. Rej, T. Gaebel, T. Boele, D. E. J. Waddington, and D. J. Reilly, Nature Communications 6, 8459 (2015).
  * Kwiatkowski  _et al._ [2017] G. Kwiatkowski, F. Jähnig, J. Steinhauser, P. Wespi, M. Ernst, and S. Kozerke, Scientific Reports 7, 7946 (2017).
  * Hu, Whiting, and Bhattacharya [2018] J. Hu, N. Whiting, and P. Bhattacharya, J. Phys. Chem. C 122, 10575–10581 (2018).
  * Seo  _et al._ [2018] H. Seo, I. Choi, N. Whiting, J. Hu, Q. S. Luu, S. Pudakalakatti, C. McCowan, Y. Kim, N. Zacharias, S. Lee, P. Bhattacharya, and Y. Lee, ChemPhysChem 19, 2143–2147 (2018).
  * Wenckenbach [2016] T. Wenckenbach,  _Essentials of Dynamic Nuclear Polarization_ (Spindrift, 2016).
  * Acosta  _et al._ [2010] V. M. Acosta, E. Bauch, M. P. Ledbetter, A. Waxman, L.-S. Bouchard, and D. Budker, Phys. Rev. Lett. 104, 070801 (2010).


Dataset

This work utilized the clinical notes of pediatric patients from the UW
hospital system. The patient cohort consists of a random sample from the
general pediatric population to improve generalizability across patient
demographics. The clinical notes span a ten-year period (1/1/2012-12/31/2021)
with 198k distinct notes from 36k distinct patients. Clinical notes are
organized into topical sections that are delineated by specific heading
formats. Patient SDoH can be described throughout the clinical narrative;
however, SDoH are most frequently documented in the social history sections of
the clinical notes. To focus the annotation on SDoH-dense portions of the
clinical notes, we applied a rule-based approach to identify topical section
headings and the social history sections, yielding 11k social history sections
for 8k distinct patients. The social history section text for a patient can be
very similar or identical across notes, so we randomly selected one social
history section per patient, resulting in 8k patients, each with a single
social history section. Finally, we randomly sampled 1,260 out of 8K social
history sections for SDoH annotations.

Annotation scheme We created detailed annotation guidelines for ten SDoH
(referred to here as event types), as listed in Table ann_guideline. The three
substance events, alcohol, drug, and tobacco, are annotated and evaluated
separately, but their performance is reported together due to their relatively
low frequency.

Each event is characterized by a trigger and multiple arguments that describe
the event’s status, type, and status. The trigger is a span with an event-type
label. Each argument attaches to the corresponding trigger and is assigned a
multi-class label, referred to here as a subtype label †† arguments and
subtype labels can be considered as attribute names and attribute values. We
chose this naming convention following the previous N2C2 SDoH challenge
lybarger20232022. , representing normalized SDoH concepts (such as Status -
past, current) that are more suitable for downstream clinical applications.
Because the most important clinical information is usually stored in a
structured format in EHRs, the normalized SDoH concepts as labels can be
directly added to other structured information to create a more comprehensive
patient representation. Arguments can be categorized into required and
optional. The required arguments define the most important attributes of the
event. A trigger can only be annotated if all required arguments can be
resolved.

The annotation scheme and event type distribution are specified in Table
ann_guideline. SDoH information was annotated using the BRAT rapid annotation
tool stenetorp-etal-2012-brat. Figure ann_example is an example describing the
patient’s living arrangement and caregivers’ employment.

[width=]figures/sdohann_ex.png justification=raggedright,singlelinecheck=false

Figure : An Annotation example: the triggers are in boldface. The box above a
trigger shows the event type, arguments and subtype labels.

[width=] figures/Models_PedSDOH.png
justification=raggedright,singlelinecheck=false

Figure : Our one-step (T5-Event) and two-step (T5-2sQA) extraction models.
T5-Event extracts all SDoH events, including triggers and arguments, in one
query. T5-2sQA extracts triggers and arguments in separate queries, where Step
Two includes the predicted triggers from Step One.

###  IE evaluation

We follow the previous N2C2 SDoH challenge lybarger20232022 evaluation
criteria. We evaluate the trigger and argument extraction performance for each
event. Two triggers are considered equivalent if they have the same event type
and overlapping spans. The trigger extraction is framed as a named entity
recognition task, and the precision, recall, and F1 are calculated. Two
arguments are considered equivalent if they are attached to equivalent
triggers and have the same argument type and subtype labels, and are evaluated
using precision, recall, and F1.

Table : Annotation scheme and event statistics for PedSHAC, where * indicates
optional arguments. The train, validation, and test sets contain 894, 121, and
245 notes, respectively. The IAA micro-averaged F1 (%) is calculated on the
last round of double annotation, consisting of 90 notes. The IAA F1 micro
averages on triggers, arguments, and triggers plus arguments are 85.1, 80.0,
and 81.9, respectively.

###  Annotator agreement

Six medical students at UW annotated SDoH events in our dataset. We first
performed two practice rounds to train the annotators and refine the
annotation guidelines, with 5 and 10 notes, respectively. After the practice
rounds, each note was annotated by two annotators (double annotation), with a
third annotator adjudicating disagreements. The Inter-Annotator Agreement
(IAA) is evaluated using the criteria in Section . We doubly annotated 360
notes through 4 rounds (90 notes per round) and then singly annotated the
remaining 885 notes. PedSHAC has an IAA micro average of 85.1 F1 across all
triggers and 80.0 F1 across all arguments in the last double-annotation round
with 90 notes. Low IAAs are from infrequently occurring events such as Food
Insecurity and Mental Health, and the annotation group carefully discussed
every disagreement. PedSHAC is split into training, validation, and test sets.
Table  presents the distribution of SDoH for each split along with the IAA for
all event types. The entirety of the test set and the majority of the
validation set are doubly annotated.

##  SDoH Information Extraction

We experimented with various LLM types and learning strategies, including i)
fine-tuning BERT, ii) fine-tuning T5, and iii) in-context learning with GPT-4.
The generative model experimentation with T5 and GPT-4 explored multiple
prompting strategies, including i) single-step text2event (Event), and ii)
two-step question answering (2sQA). Both prompting approaches were explored
with T5 through fine-tuning and GPT-4 through in-context learning. Fine-tuning
BERT (mSpERT): Following prior work in the N2C2 SDoH challenge
lybarger20232022, we use our high-performing, multi-label variation of the
Span-based Entity and Relation Transformer model (mSpERT)
eberts2019SpERT,lybarger2022mSpERT111https://github.com/Lybarger/sdoh_extraction,
as the BERT baseline. mSpERT is a span-based extractor that jointly extracts
entities and relations. In the PedSHAC extraction task, mSpERT assigns
multiple labels to a given span and assumes all predictions for a given span
are associated with the same event. As all PedSHAC arguments share the same
span as the trigger, mSpERT did not generate any relation predictions between
spans. Fine-tuning T5 with single-step text2event Prompting (T5-Event): Recent
work lu-etal-2021-text2event, ma-etal-2023-dice demonstrates that entity and
relation extraction tasks can be reformulated into text2event tasks using
generative encoder-decoder models like T5 t5,chung2022scaling and decoder-only
models like GPT-4 openai2023gpt4. We map each event annotation to a structured
text representation romanowski2023extracting,lu-
etal-2021-text2event222https://github.com/romanows/SDOH-n2c2/. Figure
illustrates our T5-Event approach. Input sequences included the entire social
history section and a model instruction. The target sequence was a sequence of
SDoH events containing trigger type and text span, followed by the required
and optional arguments. The trigger span was repeated with its argument to
associate the arguments with the trigger span. Multiple events in the output
were separated with ‘AND’ for parsing. T5-Event extracts all PedSHAC SDOH
events for a social history section in one step. Fine-tuning T5 with two-step
QA Prompting (T5-2sQA): we utilize a two-step pipeline approach to first
extract triggers spans ma-etal-2023-dice and then resolve subtype labels
through multiple-choice questions ma2023multichoice. Figure  illustrates our
two-step approach. In step 1, the model input is a prompt specifying the
target event type and the social history section text, and the model’s desired
output is a list of trigger spans associated with the target event type. In
step 2, we apply multi-choice QA to resolve the argument subtype labels for
each identified trigger and each argument type relevant to the event type. The
input prompt specifies the argument, the relevant trigger within the note, and
all possible argument subtypes. An additional choice, “none,” is added for
optional arguments, indicating the argument may not be present for that event.
The model output is the selected subtype. GPT-4 with In-context Learning:
Previous research demonstrates LLMs can achieve high performance through in-
context learning agrawal-etal-2022-large. Additionally, some proprietary LLMs,
including GPT-4, cannot currently be fine-tuned. Using prompt-based, in-
context learning, information about the desired task is conveyed through
instructions and few-shot examples. The larger context window of recent LLMs,
including GPT-4 openai2023gpt4, which can accommodate up to 32k tokens, allows
detailed text-based instructions and several response examples to be included
in the prompt. We explored three in-context learning strategies: i) Event and
2sQA – simple instructions without explanation of annotated phenomena. For
GPT-Event, our instruction contained a list of all the event and argument
types and an illustration of the T5-Event output format using a randomly
chosen example note. GPT-2sQA uses the same prompts provided to T5-2sQA, ii)
GPT + guide – 2sQA prompt with a brief description of target trigger/argument
based on a summary of the annotation guideline, iii) GPT + 3-shot – three few-
shot examples, in addition to the GPT + guide prompts. For the +3-shot
setting, we randomly selected three example social history sections from the
train set per GPT query, with some restrictions: (1) for trigger extraction:
the three example notes contained zero, one, and more than one triggers of
specific event type respectively; (2) for required argument extraction, three
randomly selected examples of events with that argument type (positive
examples); and (3) for optional argument extraction such as residence, one
random negative example as an event without that argument, and two random
positive examples, are included from event associated with the argument type.

##  Experimental Paradigm

In fine-tuning, we trained extraction models on the train set, optimized the
hyperparameters on the validation set, and applied the best-performing models
to the withheld test set. In in-context learning, we utilized the annotation
guideline and examples from the train set. We initialized the BERT-based
mSpERT model from Bio+ClinicalBERT alsentzer-etal-2019-publicly. For T5
experimentation, we initialized from Flan-T5-Large (780M) chung2022scaling, an
instruction-tuned T5 variant. For GPT-4 experiments, we used OpenAI’s
GPT-4-32k (version: 2023-03-15-preview) with the chat completion API provided
through our HIPAA-compliant Azure server instance and utilized the ‘role’
preset (‘system’, ‘user’, and ‘assistant’) arguments for providing our
prompts. The system message includes the same instructions as the T5
experiments (except for the subtype options) and the distilled annotation
guideline. The user message includes the note and subtype options for the
argument extraction. We utilize multiple user-assistant input pairs to
simulate the conversation history as in-context learning few-shot examples.