# Locking Machine Learning Models into Hardware

Eleanor Clifford, Cheng Zhang,
Yiren Zhao
Imperial College London
&Adhithya Saravanan††footnotemark: , Harry Langford††footnotemark: ,
Robert Mullins
University of Cambridge &Ilia Shumailov, Jamie Hayes
Google DeepMind
Shared first authorship

###### Abstract

Modern Machine Learning models are expensive IP and business competitiveness
often depends on keeping this IP confidential. This in turn restricts how
these models are deployed – for example it is unclear how to deploy a model
on-device without inevitably leaking the underlying model. At the same time,
confidential computing technologies such as Multi-Party Computation or
Homomorphic encryption remain impractical for wide adoption. In this paper we
take a different approach and investigate feasibility of ML-specific
mechanisms that deter unauthorised model use by restricting the model to only
be usable on specific hardware, making adoption on unauthorised hardware
inconvenient. That way, even if IP is compromised, it cannot be trivially used
without specialised hardware or major model adjustment. In a sense, we seek to
enable cheap locking of machine learning models into specific hardware. We
demonstrate that locking mechanisms are feasible by either targeting
efficiency of model representations, such making models incompatible with
quantisation, or tie the model’s operation on specific characteristics of
hardware, such as number of cycles for arithmetic operations. We demonstrate
that locking comes with negligible work and latency overheads, while
significantly restricting usability of the resultant model on unauthorised
hardware.

##  1 Introduction

The monetary expenditures associated with developing Machine Learning (ML)
models are increasing rapidly with the advent of large generative models.
Models with over a trillion parameters are now being trained on web-scale data
(Brown et al., 2020). These models have become valuable Intellectual Property
(IP) assets, yet ensuring their competitive edge remains uncompromised when
deployed on-device proves challenging. Competitors may reverse engineer the
model’s architecture and parameters, redeploying it on their software and
hardware stack. Concurrently, governance of Machine Learning models is a
concern (Hadfield and Clark, 2023). Especially in safety-critical
applications, it may be necessary to limit model execution to special
authenticated settings. Here, we usually rely on hardware and software
combinations to prevent model use on unverified platforms which may lead to
the potential misuse of the model.

Existing ML governance and IP protection methods can be classified into two
categories: namely policies and centralised serving. Policy-based methods
focus on either access control or licensing. For example, accessing Llama
models requires users to sign a terms of service agreement (Touvron et al.,
2023), and licenses like OpenRail (Ferrandis, 2022) include usage limitations
to prevent misuse of these large language models (LLMs). However, it is
entirely possible for these access-controlled models to leak, as it happened
to LLama2 (Vincent, 2023), and malicious users may not adhere to any existing
licenses, as pointed out by Henderson et al. and Lin et al.. Another approach
involves hosting ML models on centralised servers and providing standard API
access to users. Companies employ this method to safeguard their IP,
implementing safety filters and safeguarding prompts to ensure appropriate
usage. It is worth mentioning that this centralised model serving necessitates
substantial resources to maintain, as all user queries are handled by
centralised computing infrastructure, unlike computations on user devices,
which are typically less checked. Furthermore, these models cannot be used
offline.

Figure 1: A high-level illustration of how Machine Learning Locking functions
work is as follows: the locked model resists efficient, or any, deployment by
adversaries on unauthorised hardware stacks. This resistance occurs because
unauthorised hardware devices inherently lack support for some hardware
operation or are unable to match the hardware properties of the authorised
hardware.

As both policy enforcement and centralised serving fail to address the issue
of deploying whole or parts of ML models on user devices, we adopt a
completely different approach from the aforementioned methods in this study,
illustrated in Figure 1. We explore the feasibility of mechanisms enabling
Machine Learning Locking, whereby a locked ML model resists efficient, or any,
deployment by adversaries on unauthorised hardware stacks. In such scenarios,
should a model be stolen or reverse-engineered, deploying it on unauthorised
platforms would be either impossible or extremely challenging.

  * •

We introduce the concept of machine learning locking, where ML models are
locked to specific hardware stacks, restricting model usage on unauthenticated
hardware.

  * •

We demonstrate the feasibility of a range of software and hardware locking
mechanisms; these methods have negligible overheads when deployed and we
quantify the difficulty of breaking these locks on unauthorised hardware
platforms.

  * •

To the best of our knowledge, ours is the first work to explore hardware
asymmetries as a foundation for governance and to discuss its implications for
AI governance.

Table 1: Taxonomy of locking methods. For reverse engineering, access is given
to the locked model but not the target hardware. * Similar to the cost of
fine-tuning the model. ** The PUF method provides sufficient entropy to use
the cheaper _AES encryption_ transformation.

Category Method Effect Reverse-engineering cost Overhead Soft locking
Sparsity-aware Slowdown or Accuracy drop Moderate* Small Quantisation-aware
Slowdown or Accuracy drop Moderate* Small Hard locking Clock fingerprint
Accuracy drop High Moderate Finite Precision Accuracy drop High Moderate SRAM
PUF Accuracy drop Infeasible Small**

##  2 Background

Software and hardware Traditional software systems often employ measures to
ensure that software executes exclusively on authorised hardware. This is
achieved by integrating hardware identifiers, or fingerprints, into the
software itself, and is similar to the methods presented in this paper. Upon
execution, the software verifies the authenticity of the hardware it is
running on by comparing its fingerprint to the expected value. Such
fingerprints can be, for example, generated with Picasso by using web-browser
agents and HTML canvases for identification (Bursztein et al., 2016), or
DrawnApart, which relies on GPU speed variance (Laor et al., 2022).

Furthermore, comprehensive verification of the software and hardware loading
process is frequently implemented. This involves employing secure boot-
loaders, firmware checks, hardware authentication tokens, and platform modules
to ensure the integrity of the system as a whole. By combining hardware
fingerprinting with a thorough verification process, these systems can
effectively restrict software execution to pre-approved hardware. Note, that
all of the solutions are usually used in conjunction. In this paper we build
examples of such mechanisms which are specifically suited for machine
learning.

Machine Learning deployment Modern machine learning deployment primarily
utilises centralised serving due to challenges in data sharing and the
necessity for specialised hardware. In the realm of on-device inference, the
prevailing approach to access and security management is through policy-based
restrictions (Outchakoucht et al., 2017) and safety fine tuning (Qi et al.,
2023).

Table 2: A comparison of the supported features of existing AI accelerators. *
We consider also the TensorCore supported instruction for NVIDIA GPUs. TEE
stands for the Trusted Execution Environment. ** DOJO utilised a specialised
floating point arithmetic (CFP), where they have a different setup for
exponent and mantissa bit widths.

Device INT4 INT8 FP8 FP16 FP32 TF32 BFLOAT16 Sparsity TEE NVIDIA A100 *
Choquette et al. (2021) ✓ ✓ ✗ ✓ ✓ ✓ ✓ ✓ ✗ NVIDIA H100 * Choquette (2023) ✗ ✓ ✓
✓ ✓ ✓ ✓ ✓ ✓ NVIDIA H200 * NVIDIA  ✗ ✓ ✓ ✓ ✓ ✓ ✓ ✓ ✓ Cerebras WSE 2 Selig
(2022) ✗ ✗ ✗ ✓ ✓ ✗ ✓ ✓ ✗ Tesla DOJO Talpes et al. (2022) ✗ ✓ ✓** ✓** ✓ ✗ ✓ ✗ ✗
Groq Gwennap (2020) ✗ ✓ ✗ ✓ ✓ ✗ ✗ ✗ ✗ Qualcomm AI100 Chatha (2021) ✗ ✓ ✗ ✓ ✓ ✗
✗ ✗ ✗ Google TPU V4i Jouppi et al. (2021) ✗ ✓ ✗ ✓ ✓ ✗ ✓ ✗ ✗ AMD MI300 AMD  ✗ ✓
✓ ✓ ✓ ✓ ✓ ✓ ✗ AMD AIE Alok (2020) ✓ ✓ ✗ ✓ ✗ ✗ ✓ ✗ ✗

###  2.1 Specialised Hardware

AI hardware vendors have designed a wide range of chips featuring a
comprehensive range of hardware intrinsic supports. This often focuses on
hardware arithmetic, optimisations such as sparsity, and security support like
Trusted Execution Environment (TEE). Table 2 presents a compilation of
recently developed AI accelerators along with their respective hardware
intrinsic supports. It details features encompassing various arithmetic
supports, including INT4, INT8, FP8, FP16, FP32, TF32, and BFLOAT16. Table 2
also considers the availability of support for sparse matrix multiplication
(sparsity) and TEE in the devices listed. The device-level hardware intrinsic
asymmetry, detailed in Table 2, provides substantial opportunities for its
exploitation in hardware-locking. While software or compiler optimisations can
imitate circuit design variations, these emulations are inherently less
efficient, often by orders of magnitude, in terms of operation-per-Watt
performance.

Furthermore, Table 2 illustrates how the same hardware support can vary in
implementation. For instance, Tesla’s DOJO adopted distinctive FP8 and FP16
schemes, deviating from the conventional schemes, and termed them CFP.
Notably, computations occur at finite precision, and the execution of the same
operation often varies across hardware platforms. Consequently, even a
standard FP32 convolution operation might yield different numerical results on
different hardware, as detailed by Schlögl et al. (2023). The differences in
hardware implementations and numerical deviations described above can also
serve as sources of asymmetry for exploration in the realm of locking.

##  3 Methodology

###  3.1 Assumptions, Goals, and Definitions

The goal of this paper is to develop locking mechanisms that make it hard to
move a machine learning model from an authorised hardware platform to an
unauthorised (different) hardware platform. That is, to build mechanisms that
deter unauthorised model use by restricting supported hardware. Note that our
locks are not designed to replace existing cryptographic security solutions
e.g. distribution and storage of encrypted weights, hardware security modules,
and root of trust, but are rather developed to complement them for settings
where restriction to specific hardware is preferred. Our methods are similar
in function to standard encryption techniques, but use hardware-specific
characteristics as key material, do not require explicit key management, and
do not require dedicated cryptographic hardware.

In what follows, we explicitly separate two main types of locking: hard and
soft. Soft locking mechanisms refer to mechanisms that do not fully restrict
normal use of the model on non-authorised hardware, but rather make use of the
model less performant or efficient. For example, consider a model that during
inference produces abnormally large amounts of internal data that slows down
inference on normal GPUs, but specialised GPUs filter the produced data to
only keep task-informative data. Hard locking mechanisms refer to mechanisms
that fully restrict use of models on non-authorised hardware, ideally with
formal e.g. cryptographic guarantees. In practice, we envision both locking
types to be used in conjunction. To make such mechanisms usable, we seek to
minimise effort required to instrument the model for deployment. At the same
time, given that both developing proprietary hardware platforms and training
large-scale foundation models are expensive, we either opt to 1 convert models
into representations that are unique to specific platforms e.g. in Table 2 we
show that INT4 representation can be favourable for soft locking, since only
two widely available hardware platforms support it; or 2 condition models on
hardware-specific behaviours e.g. introducing model dependency on latency of
scatter/gather operations.

We explicitly note that none of the mechanisms described in this paper on
their own provide security and could be circumvented by a knowledgeable
attacker with enough effort and appropriate access. Furthermore, they in no
way stop an adversary from performing model extraction, but these would result
in significant adversary costs (Tramèr et al., 2016; Truong et al., 2021;
Shafran et al., 2023). Yet, our locks present a significant challenge for an
attacker with locked model-only access.

###  3.2 Soft locking methods

Optimisations for model inference, such as pruning and quantisation, typically
convert model parameters into sparse or low-precision tensors. These
perturbations in the parameter space can lead to a test-time disparity,
meaning that models, even if derived from the same original model but
quantised with different arithmetics or pruned to varying sparsities, can
misclassify distinct samples. Prior work has leveraged artefacts from
quantisation or pruning to develop stealthy backdoor attacks (Ma et al., 2023;
Hong et al., 2021). Nonetheless, in soft locking, our interest lies in
implementing strategies that enable optimisations exclusively on hardware
platforms with specific intrinsic support, not on others. Unsupported or
unauthorised hardware platforms may still be capable of running the same model
but would suffer from inefficient execution and/or considerable performance
degradation.

Sparsity-aware locking Pruning is a family of methods that transforms models
with dense parameters into sparse ones (LeCun et al., 1989). Pruning reduces
the number of parameters needed to store the model and potentially decreases
the amount of floating-point operations required, if supported by the
hardware. Table 2 reveals that only NVIDIA A100, H100, Cerebras WSE2, and AMD
MI300 possess hardware intrinsic support for sparse tensor acceleration,
indicating the existence of hardware asymmetry in pruning. We propose a simple
fine-tuning scheme to produce models that significantly degrade in performance
when used at an unauthorised sparsity level. The loss for this proposed
manipulation takes the form:
$\mathcal{L}=\mathcal{L}(f_{p}(x,p_{1}),y)+\lambda(\epsilon-\mathcal{L}(f_{p}(x,p_{2}),y))^{2}$.
Here $f_{p}(\cdot)$ represents the pruned version of the original network
$f(\cdot)$, where $x$ is the input and $p_{1}$ and $p_{2}$ are values between
$0$ and $1$ that define the level of pruning in the pruned network. When
$p=0$, we have an unpruned, dense network. $\mathcal{L}$ denotes a valid loss,
for example cross-entropy. $\lambda$ is a hyperparameter that allows tuning of
the relative magnitudes of the original training term and the pruning-
degradation term. $\epsilon$ denotes a target loss value for the pruned model.
In essence, our loss promotes free optimisation for models with a sparsity
$p_{1}$ while aiming to drive models with a sparsity $p_{2}$ toward a
suboptimal point, given that $p_{1}\neq p_{2}$. This results in a sparsity-
aware locking, where the model exhibits higher accuracy at a sparsity level of
$p_{1}$ and significantly worse performance at a sparsity level of $p_{2}$.

Quantisation-aware locking Machine Learning models are often distributed at
lower quantisation to allow deployment on specialist or constrained hardware.
ML locking could therefore be accomplished by limiting which quantisations a
particular model could be run at. We propose using a loss similar to the
pruning-degradation loss defined above and that used by Hong et al. (2021) to
backdoor models, in order to lock the quantisation levels which a model can be
used on. The differences between our aim and Hong et al. (2021) are two-fold.
Firstly, we aim to make transferring across hardware systems more challenging
in order to accomplish ML locking instead of trying to backdoor model
quantisation. Secondly, we expand upon the approach of Hong et al. (2021)
across various arithmetic types, not just within the integer arithmetic
domain, as many chips listed in Table 2 support integer arithmetic but vary in
their support for different floating-point or even custom arithmetics.

###  3.3 Hard locking methods

Our _hard locking_ methods are based on using a fingerprint obtained from a
specific device to transform model parameters in a way that is difficult to
invert without the fingerprint. Note that such signatures can be shared across
devices from the same family, as we show with clock fingerprints, as well as,
specific individual devices as we show with finite precision fingerprints.

First, the properties of the target device model are used to generate a high-
entropy fingerprint unique to the device or the device model, which is hard or
impossible to replicate from other devices or device families. Then, a
parameter transformation function is used to modify the model parameters based
on the fingerprint. The model is only ever stored in its transformed form, and
detransformed on the fly by the authorised device. Without knowledge of the
fingerprint, the transformed model is not functional, and the fingerprinting
method is designed to have high entropy, such that it cannot easily be guessed
or brute-forced without access to the authorised hardware.

Device fingerprinting The fingerprinting method can be anything which produces
a consistent and unguessable output on one device or device model but a
different output on other devices. It should have sufficient entropy to
produce sufficient key material for the encryption-like parameter
transformations. We propose three candidate device fingerprints: the _clock
fingerprint_ , the _finite precision fingerprint_ , and the _PUF fingerprint_.
Different methods can be combined together for additional entropy.

Table 3: Clock fingerprint on various devices Device | GPU | Fingerprint
---|---|---
GTX 1080Ti | GP102 | 72100 to 723ff
RTX 2080Ti | TU102 | 49a59
RTX 3090 | GA102 | 4b85a
RTX A6000 | GA102 | 4b85a

Clock fingerprint The _clock fingerprint_ is generated by counting the number
of clock cycles taken by a CUDA device to repeatedly add numbers together. The
fingerprint is this number represented as five symbols of hexadecimal. This
can be seen in Table 3. We find experimentally that the GP102 produces an
output stochastically between 72100 and 723ff, while the TU102 and GA102
produce a fingerprint deterministically. The entropy of the clock fingerprint
is upper bounded by the number of bits in the output (20). Clock fingerprints
are an example of device-family fingerprint. We show our clock fingerprint
generator in Appendix J.

Finite precision fingerprint The _finite precision fingerprint_ is generated
via numerical errors specific to an arithmetic and precision. In ML systems,
these can also occur due to inference-time microbenchmarking (Schlögl et al.,
2023). Device-specific ML framework implementation choices produce a unique
error which can be used as a fingerprint. The total entropy is determined by
the total number of possible sets of such choices. Schlögl et al. (2023) find
a maximum of four error equivalence classes, or two bits of entropy, for a
single convolution. The total bits of entropy available in a convolution-based
finite precision fingerprint is therefore upper bounded by twice the number of
convolutions performed. In reality, a large number of bits would be difficult
to obtain due to strong correlations between the algorithm choices. At the
same time, we find that floating point operations in general can be used to
generate device-specific errors. We compute a finite precision fingerprint
based on the SHA-256 hash of the error generated in a sequence of linear
layers. These can be seen in Appendix I. We find that the fingerprint is
consistent between different GPUs of the same model on the same system, but
differ between systems. This could be used to further lock models to other
components of the system. We show our precision fingerprint generator in
Appendix K.

PUF fingerprint Physically Unclonable Functions (PUFs) that exist in (or are
explicitly introduced into) hardware can be used to derive a high-entropy
device fingerprint. For example, Van Aubel et al. (2015) finds that the
initial state of shared SRAM memory in NVIDIA GPUs is one example of an un-
advertised PUF, while other constructions are possible (Li et al., 2015;
Forlin et al., 2020). These PUFs could be combined with fuzzy extractors
(Dodis et al., 2004) to achieve a reliable fingerprint. In model-distributed
setups, these PUFs could be chained together across multiple devices.

Parameter transformation In order to achieve _hard locking_ , we must
transform the model parameters in a way that can only be easily inverted with
knowledge of the fingerprint. There are three ideal properties of the
transformation, defined informally as follows:

  1. 1.

Destruction: The performance of a model with transformed parameters or
parameters detransformed with an incorrect fingerprint should be significantly
hampered, ideally equivalent to random guessing. Without destruction, the
model would still be usable without the fingerprint, and hard locking fails.

  2. 2.

Encryption: No information about the original parameters should be obtainable
from the transformed parameters, other than the information required for
_indistinguishability_. If encryption holds, then the attacker must brute-
force all possible fingerprints to determine the correct one to reveal the
secret. If it does not, in some cases cheaper methods such as gradient descent
may be applicable for extracting the original parameters.

  3. 3.

Indistinguishability: Incorrectly detransformed parameters should be
statistically indistinguishable in aggregate from correct parameters. If
indistinguishability holds, then in a brute-force attack the attacker must run
each candidate model on test data and choose the correct one by maximising
test accuracy. This is in general more expensive and error-prone than a
statistical test.

Table 4: Comparison of parameter transformations that can be applied with different fingerprints. Method | Indistinguishablity | Encryption | Destruction
---|---|---|---
AES encryption | ✗ | ✓ | ✓
Parameter shuffling | ✓ | ✗ | ✓
Pre-transformed AES encryption | ✓ | ✓ | ✓

Table 4 considers three parameter transformation methods in terms of these
properties: _AES encryption_ , _Parameter shuffling_ , and _Pre-transformed
AES encryption_. These are described below.

AES Encryption The obvious method to transform the model parameters is with a
classical encryption scheme, this gives rise to the _AES encryption_ method,
in which the parameters of the model are collected together into a bytestream,
which is then AES-encrypted, keyed by the SHA-256 hash of the fingerprint. The
resulting bytestream is then the transformed parameters. This achieves perfect
_encryption_ and _destruction_ , but not _indistinguishability_ , because
incorrectly decrypted parameters will be uniformly distributed, in contrast to
the correct parameters.

In conventional cryptographic schemes, security can be increased by using
higher-entropy keys to make brute-forcing infeasible. Since the entropy of a
fingerprint is fixed and cannot be increased, we focus instead on making each
decryption attempt more expensive, i.e. key-stretching (Kelsey et al., 1997).
We propose achieving this through _indistinguishability_. With
indistinguishability, for each candidate fingerprint the attacker must
evaluate the candidate model’s accuracy, which is computationally expensive.
Indistinguishability of parameters is feasible since the statistics of ML
parameters are much easier to fake than the plaintext of most encryption
schemes (for example, coherent English text, or a valid JPEG file).

In order to achieve _indistinguishability_ and thus make a brute-force attack
more difficult, we developed two further transformation methods: _parameter
shuffling_ , and _pre-transformed AES encryption_.

Parameter shuffling In the _parameter shuffling_ method, the fingerprint is
used as a seed to generate a random permutation of the parameters, which is
then used in place of the original parameters. This achieves near-perfect
_indistinguishability_ and _destruction_ , but does not achieve _encryption_ ,
as some information about the original parameters still exists in the permuted
parameters.

In practice, however, because the search space of permutations (factorial of
the number of parameters) is many orders of magnitude larger than the search
space of keys (two to the power of the fingerprint entropy), it is much more
efficient for an attacker to brute force the key space than attempt gradient
descent in the permutation space, and so the lack of perfect encryption is
unimportant. For example, a very large fingerprint with 256 bits of key
material has a search space of $2^{256}\approx 10^{80}$, but even a small
ResNet18 test model has over $10^{7}$ parameters, equating to a permutation
search space of $\left(10^{7}\right)!\approx 10^{{10}^{8}}$, untenably larger.

Pre-transformed AES encryption The problem with naïve AES transformation is
that incorrectly detransformed parameters are uniformly distributed in byte-
space, while correct parameters have some other distribution, often Gaussian,
making cracking considerably easier.

We can correct for this problem by defining an additional transformation
function, the ‘pre-transformation’ function. This is applied to the parameters
before encryption, such that these parameters are then uniform in byte-space,
and indistinguishable from incorrectly decrypted bytes. This pre-
transformation function is reversed after decryption.

The sole purpose of the pre-transformation function is to achieve
indistinguishability by transforming the uniform distribution of incorrectly
decrypted AES to the same distribution as the correct parameters. It must be
stored unencrypted to prevent distinguishability by statistics on the pre-
transformation. We therefore say that the pre-transformed AES encryption
transformation achieves the full trifecta of _indistinguishability_ ,
_encryption_ , and _destruction_.

A generic distribution can be pre-transformed to be uniform by applying its
cumulative distribution function to it. For example, if $X$, the distribution
of model parameters, is Gaussian with mean $\mu$ and variance $\sigma^{2}$,
then $Y$, the distribution of pre-transformed model parameters, is uniform in
the region (0, 1), where
$Y=\frac{1}{2}\left[1+\textrm{erf}\left(\frac{X-\mu}{\sigma\sqrt{2}}\right)\right].$
If $Y$ is then encoded as integers spanning the full possible integer range,
its distribution will also be uniform in byte space, and thus
indistinguishable from incorrectly decrypted bytes.

We evaluated pre-transforming the parameters of a test ResNet18 model trained
on CIFAR10, first by assuming that it was Gaussian, and secondly by directly
estimating the cumulative distribution function via sampling of the
parameters. This can be seen in Appendix H.

In practice, assuming that the distribution is Gaussian may be problematic. If
there are any outliers, then applying the cumulative distribution function
will take these outliers so close to 0 or 1 that they can no longer be
represented with sufficient precision in floating point, and they will become
exactly 0 or 1. Then, these outliers detransform to positive or negative
infinity, destroying the model. There are sufficiently many outliers that
subsequently casting these outliers to any fixed finite value will destroy the
accuracy of the model. Regardless of any scheme to correct for this,
_indistinguishability_ is broken, because these infinities do not occur with
comparable frequency in incorrectly detransformed data.

To fix this problem, we instead computed the cumulative distribution function
of the parameters empirically, then used this as a look-up table to transform
the distribution into a uniform distribution. For an $n$-bit precision,
computing the pre-transformation requires $\mathcal{O}(2^{n})$ time and space,
but applying the pre-transformation or de-pre-transformation is cheap. This is
computationally tractable for precisions up to 32-bits and inexpensive for
precisions up to 16-bit. The results for FP16 can be seen in Figure 16 in
Appendix H. There exists a trade-off between the pre-transform being exactly
invertible and it exactly matching the target distribution for random data.
Reported results are for an exact pre-transformation.

##  4 Evaluation

###  4.1 Evaluating soft locking

We evaluate the effects of different soft locking schemes as illustrated in
Table 1. When run on unauthorised hardware, soft-locked models will cause
either a slow down in inference, or a drop in accuracy or model performance.
The former occurs because unauthorised hardware may emulate the execution of
authorised hardware at the software level, due to a lack of hardware intrinsic
support, this will result in slowdowns, as we evaluate in Section 4.1.3.
Alternatively, if unauthorised hardware directly executes only what is
natively supported, it would incur an accuracy penalty, as illustrated in
Section 4.1.1, and Section 4.1.2.

####  4.1.1 Performance degradation for sparsity-aware lock

We investigate sparsity-aware locking, for the widely-used
$l_{1}$-unstructured pruning (LeCun et al., 1989; Li et al., 2017) across
pruning levels (0.05, 0.10, 0.25, 0.50, 0.75) for both vision and language
models, across different datasets. For all vision models, we fine-tune a
trained model 111Training and fine-tuning hyper-parameters presented in the
Appendix D. with the loss defined in Section 3.2 for 25 epochs, with $\lambda$
= 1 and $\epsilon$ = 5. The language models are fine-tuned for 3 epochs, as
this was sufficient for effective manipulation.

We fine-tuned the BERT model (Devlin et al., 2019), and specifically the bert-
base-cased-finetuned-{sst2, cola, mrpc} HuggingFace checkpoints on the
respective GLUE tasks (Wang et al., 2018), with the pruning-resistant loss
from Section 3.2. We also evaluate performance on various vision tasks such as
CIFAR10, CIFAR100 (Krizhevsky and Hinton, 2009), and Flowers102 (Nilsback and
Zisserman, 2008), utilising ResNet18, ResNet50 (He et al., 2016), and ViT-B
(Dosovitskiy et al., 2021) for these specific tasks.

Table 5 displays sparsity-aware lock results, we present the following metrics
as illustrated in Figure 2:

  * •

$\mathit{Acc}_{\text{authorised}}^{\text{locked}}$: the accuracy of models
with soft locks running on authorised hardware.

  * •

${\color[rgb]{0.05078125,0.54296875,0.36328125}\definecolor[named]{pgfstrokecolor}{rgb}{0.05078125,0.54296875,0.36328125}\Delta_{\text{orig}}}=\mathit{Acc}_{\text{original}}-\mathit{Acc}_{\text{authorised}}^{\text{locked}}$
measures the impact of soft locking on authorised execution by comparing the
accuracy of the original model (without locks) to that of the locked model
executing on authorised hardware. A small
${\color[rgb]{0.05078125,0.54296875,0.36328125}\definecolor[named]{pgfstrokecolor}{rgb}{0.05078125,0.54296875,0.36328125}\Delta_{\text{orig}}}$
value is desirable.

  * •

${\color[rgb]{0.80078125,0,0}\definecolor[named]{pgfstrokecolor}{rgb}{0.80078125,0,0}\Delta_{\text{lock}}}=\mathit{Acc}_{\text{authorised}}^{\text{locked}}-\mathit{Acc}_{\text{unauthorised}}^{\text{locked}}$
measures the degradation in accuacy when locked models are deployed on
unauthorised devices. A large
${\color[rgb]{0.80078125,0,0}\definecolor[named]{pgfstrokecolor}{rgb}{0.80078125,0,0}\Delta_{\text{lock}}}$
value is desirable.

Figure 2: Applying soft locks to a model with
$\mathit{Acc}_{\text{original}}$,
${\color[rgb]{0.05078125,0.54296875,0.36328125}\definecolor[named]{pgfstrokecolor}{rgb}{0.05078125,0.54296875,0.36328125}\Delta_{\text{orig}}}$
and
${\color[rgb]{0.80078125,0,0}\definecolor[named]{pgfstrokecolor}{rgb}{0.80078125,0,0}\Delta_{\text{lock}}}$
measure together the effectiveness of locking.

Table 5 demonstrates that sparsity-aware lock leaves the accuracy of the
authorised execution largely unaffected, as shown by the small
${\color[rgb]{0.05078125,0.54296875,0.36328125}\definecolor[named]{pgfstrokecolor}{rgb}{0.05078125,0.54296875,0.36328125}\Delta_{\text{orig}}}$
values. Meanwhile, locked models experience significant accuracy degradation
when operating on unauthorised hardware, as indicated by the substantial
${\color[rgb]{0.80078125,0,0}\definecolor[named]{pgfstrokecolor}{rgb}{0.80078125,0,0}\Delta_{\text{lock}}}$
values. We present the accuracy drop from hardware transfer without locking,
$\Delta_{\text{base}}$, in the Appendix (F). It is clear that the significant
accuracy drop from transfer can be attributed to locking (and not merely
pruning), as the degradation in accuracy of the original models,
$\Delta_{\text{base}}$, is generally much lower than that of the locked
models,
${\color[rgb]{0.80078125,0,0}\definecolor[named]{pgfstrokecolor}{rgb}{0.80078125,0,0}\Delta_{\text{lock}}}$.
Certain settings, such as BERT on MRPC, show that the sparsity-aware lock can
trigger a larger degradation in performance when the unauthorised
configuration is more sparse (i.e. greater pruning proportion, $p$). However,
even a low unauthorised pruning proportion of $p=0.05$ is sufficient for
sparsity-aware locks to be effective.

Table 5: Results are presented as
$\mathit{Acc}_{\text{authorised}}^{\text{locked}}({\color[rgb]{0.05078125,0.54296875,0.36328125}\definecolor[named]{pgfstrokecolor}{rgb}{0.05078125,0.54296875,0.36328125}\Delta_{\text{orig}}},{\color[rgb]{0.80078125,0,0}\definecolor[named]{pgfstrokecolor}{rgb}{0.80078125,0,0}\Delta_{\text{lock}}})$.

Pruning Levels Dataset 0.05 0.25 0.50 BERT CoLA 0.80 (-0.03, 0.49) 0.84 (0.00,
0.53) 0.83 (0.00, 0.53) MRPC 0.84 (-0.02, 0.15) 0.86 (0.00, 0.18) 0.86 (0.00,
0.54) SST-2 0.92 (-0.01, 0.43) 0.93 (0.00, 0.43) 0.92 (0.00, 0.43) Resnet18
CIFAR10 0.89 (0.03, 0.67) 0.93 (0.00, 0.83) 0.93 (0.00, 0.83) CIFAR100 0.73
(0.03, 0.71) 0.76 (0.00, 0.75) 0.76 (0.00, 0.75) Flowers102 0.84 (0.04, 0.83)
0.89 (0.00, 0.88) 0.89 (0.00, 0.88) Resnet50 CIFAR10 0.92 (0.01, 0.81) 0.93
(-0.01, 0.83) 0.94 (-0.01, 0.84) CIFAR100 0.69 (0.09, 0.63) 0.78 (0.00, 0.77)
0.78 (0.00, 0.77) Flowers102 0.76 (0.10, 0.74) 0.86 (0.00, 0.84) 0.86 (0.00,
0.84) ViT-B_16-224 CIFAR10 0.99 (0.00, 0.89) 0.99 (0.00, 0.89) 0.99 (0.07,
0.89) CIFAR100 0.92 (-0.02, 0.92) 0.93 (-0.02, 0.92) 0.93 (-0.02, 0.93)
Flowers102 1.00 (0.00, 0.98) 1.00 (0.00, 0.99) 1.00 (0.00, 0.99)

An obvious attack on the sparsity-lock is re-training with a pruning-aware
loss on the unauthorised hardware. We investigate the recovery accuracy in
G.2, and find that at lower levels of sparsity, $p$ (0.05 and 0.25), re-
training does not allow the model to reach the accuracy of the original (i.e.
unlocked) model running on unauthorised hardware. At $p=0.50$, accuracy is
regained after roughly 5 epochs. As such, the locking procedure is highly
effective, as in the best case it causes terminal damage, and otherwise, it
necessitates further training from the unauthorised users, who may not have
access to the training data. We present further investigation in the Appendix
(G, F.3), including into the specificity of locking to the assumed
unauthorised sparsity level, and typical locking profiles.

####  4.1.2 Performance degradation for quantisation-aware lock

We investigate quantisation-aware locking on ResNet models across a set of
authorised, unauthorised arithmetic pairs, using the quantisation-aware locks
described in Section 3.2. The models are trained for 25 epochs, with $\lambda$
= 1 and $\epsilon$ = 5, as with sparsity-aware locking. We note that Hong et
al. (2021) previously used a similar loss in Section 3.2 to learn models that
degrade upon quantisation to a lower precision. We extend this by
investigating the effectiveness of preventing transfer across both precision
and arithmetics and present the results below, using the metrics from Section
4.1.1.

Table 6: Results are presented as
$\mathit{Acc}_{\text{authorised}}^{\text{locked}}({\color[rgb]{0.05078125,0.54296875,0.36328125}\definecolor[named]{pgfstrokecolor}{rgb}{0.05078125,0.54296875,0.36328125}\Delta_{\text{orig}}},{\color[rgb]{0.80078125,0,0}\definecolor[named]{pgfstrokecolor}{rgb}{0.80078125,0,0}\Delta_{\text{lock}}})$.

authorised $\rightarrow$ Unauthorised Resnet18 Resnet50 FP32 $\rightarrow$
8-bit MiniFloat 0.90 (-0.02, 0.65) 0.92 (-0.04, 0.67) Int8 $\rightarrow$ 8-bit
MiniFloat 0.47 (+0.41, 0.06) 0.50 (+0.39, 0.07) FP16 $\rightarrow$ Int8 0.90
(-0.02, 0.62) 0.91 (-0.04, 0.61) 16-bit MiniFloat2 $\rightarrow$ Int8 0.90
(-0.01, 0.72) 0.91 (-0.03, 0.81)

Here, we use FP32 to simulate these formats during fine-tuning to prove the
idea, which is independent of hardware. In both models, the quantization-aware
lock causes performance to degrade to near-random guessing performance when
the quantization is to a lower precisions than the original model, as evident
by both FP32 to 8-bit MiniFloat and 16-bit MiniFloat222Besides FP16, vendors
have various MiniFloat formats, such as Google’s BFloat16 and NVidia’s
TensorFloat. Here we set an exponent width of 5 and an exponent bias of 11. to
Int8 in Table 6. Manipulation across the same precision, but different
arithmetic, was not successful (Int8 to 8-bit MiniFloat), showing there is not
sufficient discrepancy in their representations for this simple manipulation
procedure to yield model degradation. However, the transfer across arithmetics
(16-Bit MiniFloat to Int8) results in a greater degradation, with a lower
difference in precision than transfer across just precision (FP32 to 8-bit
MiniFloat).

####  4.1.3 Emulation cost for soft locking

Table 7: Emulation costs for soft locking. Though the accuracy of soft locked
model can be recovered by emulation on unauthorised hardware, the inference
suffers from inefficient execution, such as a lower throughput and higher
latency. The single matmul is at a size of (2048, 2048), and we use the
OPT-2.7B model at a batch size of 4 on NVIDIA A6000. TOPS denotes Tera
operations per second, and TPS is tokens per second.

Sparsity-aware Quantization-aware Workload Metric Real Emulated Real Emulated
Single Matmul Throughput (TOPS) 49.97 18.85 79.22 22.72 Latency (ms) 0.43 1.14
0.27 0.95 OPT inference Throughput (TPS) 4692.20 2468.31 3505.22 1865.41
Latency (ms) 436.47 829.72 584.27 1097.88

As described in Section 3.2 and shown in Table 1, unauthorised devices without
supported hardware intrinsics may opt for operation emulation, which incurs
additional throughput costs. For a fair comparison, we assess the performance
of both a single matrix multiplication operation (matmul) and the prefill
phase of the entire network with and without software emulation on the same
device, an NVIDIA A6000 GPU. In Table 7, for sparsity-aware locking, the term
“Real” signifies that weight matrices exhibit a sparsity level of $0.995$,
implemented via torch.sparse.mm. Conversely, “Emulated” indicates that the
linear layer employs a full weight matrix alongside a $0.995$-sparsity mask;
during runtime, this mask is applied to the weights, which are then processed
using nn.functional.linear. Regarding quantisation-aware lock, “Real” means
the hardware executes the quantised operation using its intrinsically
supported INT8 GEMM, where “Emulated” is emulating the INT8 operations using
ordinary FP32 operations. We use the MASE flow for both sparsity and
quantisation emulations (Cheng et al., 2023). We consider both a single matrix
multiply as the size of (2048, 2048) and a full OPT-2.7B model inference at a
batch size of 4.

Our results in Section 3.2 suggest that emulation comes at a large cost. For
instance for a single matrix multiply, both sparsity-aware and quantisation-
aware locks induce a 2.65$\times$ and 3.52$\times$ overhead in latency, and a
2.65$\times$ and 3.49$\times$ reduction in throughput. For the inference of
OPT-2.7B, the emulation overhead is around 2$\times$ in both latency and
throughput.

###  4.2 Hard locking

We tested the three transformation functions presented in Section 3.3 on
ResNet18 (He et al., 2016) trained on CIFAR10 (Krizhevsky and Hinton, 2009).
All three successfully achieve _destruction_ , as can be seen in Table 8. We
further tested the cost to fully brute force each. In our tests on this very
small model, the _indistinguishability_ property of the _Parameter shuffling_
and _Pre-transformed AES_ methods added to the cracking cost by just under one
order of magnitude. In a larger and more expensive to evaluate model, we
expect the difference to be several orders of magnitude.

Table 8:  Accuracy of test model (FP16 ResNet18 trained on CIFAR10) with
parameter transformations. 10% is random guessing. "Correctly detransformed"
is detransformed with the same fingerprint that was transformed with, while
"Incorrectly detransformed" is detransformed with any other fingerprint. Cost
refers to experimental cost of model extraction on a toy model (CPU hours for
full brute-force). $b$ is the number of bits of entropy in the fingerprint,
discussed further in Appendix B.

Accuracy Cracking cost (s) Method Original Transformed Correctly detransformed
Incorrectly detransformed AES encryption 95.4% 10% 95.4% 10% $0.3\times 2^{b}$
Parameter shuffling 95.4% 10% 95.4% 10% $2.7\times 2^{b}$ Pre-transformed AES
encryption 95.4% 10% 95.4% 10% $1.3\times 2^{b}$

##  5 Conclusion

In this paper we introduce machine learning locking, a novel paradigm for
safeguarding machine learning models from unauthorised use to address a
growing concern about intellectual property protection and responsible AI use.
We investigate a number of different locking mechanisms, encompassing both
soft locking, which discourages model theft by imposing performance penalties
on unauthorised hardware, and hard locking, which leverages hardware
fingerprints to (cryptographically) bind models to specific platforms. Our
experiments demonstrated the effectiveness of these locking mechanisms, both
in preserving model performance and introducing significant complexity in
removing the locks. By investigating hardware-based locking mechanisms, we
offer a potential solution for safeguarding valuable on-device ML models.

## References

  * [1] AMD Instinct MI300 Series Accelerators.  https://www.amd.com/en/products/accelerators/instinct/mi300.html.  Accessed: 2024-03-03.
  * Alok [2020] G. Alok.  Architecture apocalypse dream architecture for deep learning inference and compute-versal ai core.  _Embedded World_ , 2020.
  * Brown et al. [2020] T. Brown, B. Mann, N. Ryder, M. Subbiah, J. D. Kaplan, P. Dhariwal, A. Neelakantan, P. Shyam, G. Sastry, A. Askell, S. Agarwal, A. Herbert-Voss, G. Krueger, T. Henighan, R. Child, A. Ramesh, D. Ziegler, J. Wu, C. Winter, C. Hesse, M. Chen, E. Sigler, M. Litwin, S. Gray, B. Chess, J. Clark, C. Berner, S. McCandlish, A. Radford, I. Sutskever, and D. Amodei.  Language Models are Few-Shot Learners.  In _Advances in Neural Information Processing Systems (NIPS)_ , pages 1877–1901. Curran Associates, Inc., 2020.  URL https://proceedings.neurips.cc/paper_files/paper/2020/file/1457c0d6bfcb4967418bfb8ac142f64a-Paper.pdf.
  * Bursztein et al. [2016] E. Bursztein, A. Malyshev, T. Pietraszek, and K. Thomas.  Picasso: Lightweight device class fingerprinting for web clients.  In _Proceedings of the 6th Workshop on Security and Privacy in Smartphones and Mobile Devices (SPSM)_ , pages 93–102, 2016.
  * Chatha [2021] K. Chatha.  Qualcomm® Cloud Al 100: 12TOPS/W Scalable, High Performance and Low Latency Deep Learning Inference Accelerator.  In _2021 IEEE Hot Chips 33 Symposium (HCS)_ , pages 1–19. IEEE, 2021\.
  * Cheng et al. [2023] J. Cheng, C. Zhang, Z. Yu, A. Montgomerie-Corcoran, C. Xiao, C.-S. Bouganis, and Y. Zhao.  Fast Prototyping Next-Generation Accelerators for New ML Models using MASE: ML Accelerator System Exploration.  _arXiv preprint arXiv:2307.15517_ , 2023.
  * Choquette [2023] J. Choquette.  NVIDIA Hopper H100 GPU: Scaling Performance.  _IEEE Micro_ , (3):9–17, 2023.
  * Choquette et al. [2021] J. Choquette, W. Gandhi, O. Giroux, N. Stam, and R. Krashinsky.  NVIDIA A100 tensor core GPU: Performance and innovation.  _IEEE Micro_ , (2):29–35, 2021.
  * Devlin et al. [2019] J. Devlin, M. Chang, K. Lee, and K. Toutanova.  BERT: Pre-training of Deep Bidirectional Transformers for Language Understanding.  In _Proceedings of the 2019 Conference of the North American Chapter of the Association for Computational Linguistics: Human Language Technologies (NAACL-HLT)_ , pages 4171–4186. Association for Computational Linguistics, 2019.  URL https://doi.org/10.18653/v1/n19-1423.
  * Dodis et al. [2004] Y. Dodis, L. Reyzin, and A. Smith.  Fuzzy Extractors: How to Generate Strong Keys from Biometrics and Other Noisy Data.  In _Advances in Cryptology - EUROCRYPT 2004_ , pages 523–540. Springer Berlin Heidelberg, 2004.
  * Dosovitskiy et al. [2021] A. Dosovitskiy, L. Beyer, A. Kolesnikov, D. Weissenborn, X. Zhai, T. Unterthiner, M. Dehghani, M. Minderer, G. Heigold, S. Gelly, J. Uszkoreit, and N. Houlsby.  An Image is Worth 16x16 Words: Transformers for Image Recognition at Scale.  In _9th International Conference on Learning Representations (ICLR)_. OpenReview.net, 2021.  URL https://openreview.net/forum?id=YicbFdNTTy.
  * Ferrandis [2022] C. M. Ferrandis.  Openrail: Towards open and responsible AI licensing frameworks, 2022\.  URL https://www.licenses.ai/faq-2.
  * Forlin et al. [2020] B. Forlin, R. Husemann, L. Carro, C. Reinbrecht, S. Hamdioui, and M. Taouil.  G-PUF: An Intrinsic PUF Based on GPU Error Signatures.  In _2020 IEEE European Test Symposium (ETS)_ , pages 1–2, 2020.
  * Gwennap [2020] L. Gwennap.  Groq rocks neural networks.  _Microprocessor Report, Tech. Rep., jan_ , 2020.
  * Hadfield and Clark [2023] G. K. Hadfield and J. Clark.  Regulatory Markets: The Future of AI Governance.  _arXiv preprint arXiv:2304.04914_ , 2023.
  * 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 (CVPR)_ , pages 770–778, 2016.
  * Henderson et al. [2023] P. Henderson, E. Mitchell, C. Manning, D. Jurafsky, and C. Finn.  Self-Destructing Models: Increasing the Costs of Harmful Dual Uses of Foundation Models.  In _Proceedings of the 2023 AAAI/ACM Conference on AI, Ethics, and Society_ , AIES ’23, page 287–296. Association for Computing Machinery, 2023\.  URL https://doi.org/10.1145/3600211.3604690.
  * Hong et al. [2021] S. Hong, M.-A. Panaitescu-Liess, Y. Kaya, and T. Dumitras.  Qu-ANTI-zation: Exploiting Quantization Artifacts for Achieving Adversarial Outcomes.  In _Advances in Neural Information Processing Systems (NIPS)_ , pages 9303–9316. Curran Associates, Inc., 2021.  URL https://proceedings.neurips.cc/paper_files/paper/2021/file/4d8bd3f7351f4fee76ba17594f070ddd-Paper.pdf.
  * Jouppi et al. [2021] N. P. Jouppi, D. H. Yoon, M. Ashcraft, M. Gottscho, T. B. Jablin, G. Kurian, J. Laudon, S. Li, P. Ma, X. Ma, et al.  Ten Lessons from Three Generations Shaped Google’s TPUv4i: Industrial Product.  In _2021 ACM/IEEE 48th Annual International Symposium on Computer Architecture (ISCA)_ , pages 1–14. IEEE, 2021.
  * Kelsey et al. [1997] J. Kelsey, B. Schneier, C. Hall, and D. Wagner.  Secure applications of low-entropy keys.  In _International Workshop on Information Security_ , pages 121–134. Springer, 1997.
  * Krizhevsky and Hinton [2009] A. Krizhevsky and G. Hinton.  Learning Multiple Layers of Features from Tiny Images.  Technical report, University of Toronto, 2009.  URL https://www.cs.toronto.edu/~kriz/learning-features-2009-TR.pdf.
  * Laor et al. [2022] T. Laor, N. Mehanna, A. Durey, V. Dyadyuk, P. Laperdrix, C. Maurice, Y. Oren, R. Rouvoy, W. Rudametkin, and Y. Yarom.  DRAWN APART: A Device Identification Technique based on Remote GPU Fingerprinting.  In _Proceedings 2022 Network and Distributed System Security Symposium (NDSS)_ , NDSS 2022. Internet Society, 2022.  URL http://dx.doi.org/10.14722/ndss.2022.24093.
  * LeCun et al. [1989] Y. LeCun, J. Denker, and S. Solla.  Optimal Brain Damage.  In _Advances in Neural Information Processing Systems (NIPS)_. Morgan-Kaufmann, 1989.  URL https://proceedings.neurips.cc/paper_files/paper/1989/file/6c9882bbac1c7093bd25041881277658-Paper.pdf.
  * Li et al. [2015] F. Li, X. Fu, and B. Luo.  POSTER: A Hardware Fingerprint Using GPU core frequency variations.  In _Proceedings of the 22nd ACM SIGSAC Conference on Computer and Communications Security_ , CCS ’15, page 1650–1652. Association for Computing Machinery, 2015.  URL https://doi.org/10.1145/2810103.2810105.
  * Li et al. [2017] H. Li, A. Kadav, I. Durdanovic, H. Samet, and H. P. Graf.  Pruning Filters for Efficient ConvNets.  In _5th International Conference on Learning Representations (ICLR)_. OpenReview.net, 2017.  URL https://openreview.net/forum?id=rJqFGTslg.
  * Lin et al. [2024] Z. Lin, J. Cui, X. Liao, and X. Wang.  Malla: Demystifying Real-world Large Language Model Integrated Malicious Services.  _arXiv preprint arXiv:2401.03315_ , 2024.
  * Ma et al. [2023] H. Ma, H. Qiu, Y. Gao, Z. Zhang, A. Abuadbba, M. Xue, A. Fu, J. Zhang, S. F. Al-Sarawi, and D. Abbott.  Quantization Backdoors to Deep Learning Commercial Frameworks.  _IEEE Transactions on Dependable and Secure Computing_ , 2023.
  * Nilsback and Zisserman [2008] M.-E. Nilsback and A. Zisserman.  Automated flower classification over a large number of classes.  In _2008 Sixth Indian conference on computer vision, graphics & image processing_, pages 722–729. IEEE, 2008.
  * [29] NVIDIA.  NVIDIA Nvidia Data Center GPU Resource Center.  https://resources.nvidia.com/l/en-us-gpu.  Accessed: 2024-22-05.
  * Outchakoucht et al. [2017] A. Outchakoucht, E.-S. Hamza, and J. P. Leroy.  Dynamic access control policy based on blockchain and machine learning for the internet of things.  _International journal of advanced Computer Science and applications_ , 8(7), 2017.
  * Qi et al. [2023] X. Qi, Y. Zeng, T. Xie, P.-Y. Chen, R. Jia, P. Mittal, and P. Henderson.  Fine-tuning aligned language models compromises safety, even when users do not intend to!  _arXiv preprint arXiv:2310.03693_ , 2023.
  * Schlögl et al. [2023] A. Schlögl, N. Hofer, and R. Böhme.  Causes and Effects of Unanticipated Numerical Deviations in Neural Network Inference Frameworks.  In _Advances in Neural Information Processing Systems (NIPS)_ , pages 56095–56107. Curran Associates, Inc., 2023.  URL https://proceedings.neurips.cc/paper_files/paper/2023/file/af076c3bdbf935b81d808e37c5ede463-Paper-Conference.pdf.
  * Selig [2022] J. Selig.  The cerebras software development kit: A technical overview.  2022\.
  * Shafran et al. [2023] A. Shafran, I. Shumailov, M. A. Erdogdu, and N. Papernot.  Beyond Labeling Oracles: What does it mean to steal ML models?, 2023.
  * Talpes et al. [2022] E. Talpes, D. Williams, and D. D. Sarma.  DOJO: The Microarchitecture of Tesla’s Exa-Scale Computer.  In _2022 IEEE Hot Chips 34 Symposium (HCS)_ , pages 1–28, 2022.
  * Touvron et al. [2023] H. Touvron, L. Martin, K. Stone, P. Albert, A. Almahairi, Y. Babaei, N. Bashlykov, S. Batra, P. Bhargava, S. Bhosale, et al.  Llama 2: Open foundation and fine-tuned chat models.  _arXiv preprint arXiv:2307.09288_ , 2023.
  * Tramèr et al. [2016] F. Tramèr, F. Zhang, A. Juels, M. K. Reiter, and T. Ristenpart.  Stealing Machine Learning Models via Prediction APIs.  In _25th USENIX Security Symposium (USENIX Security 16)_ , pages 601–618. USENIX Association, 2016.  URL https://www.usenix.org/conference/usenixsecurity16/technical-sessions/presentation/tramer.
  * Truong et al. [2021] J.-B. Truong, P. Maini, R. J. Walls, and N. Papernot.  Data-Free Model Extraction.  In _Proceedings of the IEEE/CVF conference on computer vision and pattern recognition (CVPR)_ , pages 4771–4780, 2021.
  * Van Aubel et al. [2015] P. Van Aubel, D. J. Bernstein, and R. Niederhagen.  Investigating SRAM PUFs in large CPUs and GPUs.  In _Security, Privacy, and Applied Cryptography Engineering_ , pages 228–247. Springer International Publishing, 2015.
  * Vincent [2023] J. Vincent.  Meta’s powerful AI language model has leaked online — what happens now?, 2023.  URL https://www.theverge.com/2023/3/8/23629362/meta-ai-language-model-llama-leak-online-misuse.
  * Wang et al. [2018] A. Wang, A. Singh, J. Michael, F. Hill, O. Levy, and S. Bowman.  GLUE: A Multi-Task Benchmark and Analysis Platform for Natural Language Understanding.  In _Proceedings of the 2018 EMNLP Workshop BlackboxNLP: Analyzing and Interpreting Neural Networks for NLP_ , pages 353–355. Association for Computational Linguistics, 2018.  URL https://aclanthology.org/W18-5446.

##  6 Appendix / supplemental material

##  Appendix A Broader Impact

Our research addresses the growing need to protect machine learning models
from misuse. By introducing the concept of machine learning locking, our work
offers a new tool to safeguard machine learning IP and aids responsible ML
development and deployment. Our work also has implications for the governance
of ML. By connecting models to specific hardware, another tool becomes
available to control where and how such models are used. This could be
particularly valuable in safety-critical applications, where ensuring that
models are only executed in authenticated settings is paramount.

##  Appendix B Hard Locking and Entropy estimates

Results are presented in Table 9. Here, we list 20 as a strict upper bound for
the clock fingerprint as it is the number of bits in the output. We expect the
real entropy to still be considerable to provide useful security guarantees,
but cannot estimate it without a large-scale experiment on diverse hardware.

Twice the number of convolutions is a theoretical result based on Schlögl et
al. [2023]’s work which finds four equivalence classes for each convolution
performed. It is an upper bound for a convolution-based finite precision
fingerprint, but convolution is not the only possible strategy for a finite
precision fingerprint, indeed our finite precision fingerprints in Appendix K
and Appendix I are not based on convolutions. The entropy of the SRAM PUF is
an experimental result [Van Aubel et al., 2015], theoretically more than 256
is possible, but in our implementation encryption schemes we assume 256 bits,
so any further entropy is not applicable.

Table 9: Comparison of methods of generating fingerprints from hardware. * estimated Method | Entropy (bits) | Error rate
---|---|---
Clock | 20 (upper bound) | <0.1%
Finite precision | 2$\times\textrm{num. of convolutions}$ (theoretical) | 5%*
SRAM PUF [Van Aubel et al., 2015] | >256 | 5%*

##  Appendix C Base-model Training

##### Sparsity-aware lock

We use open-source SoTA setups for training the base models in soft locking
experiments:

  * •

We use an open-source Github checkpoint333Github repository:
huyvnphan/PyTorch_CIFAR10 as the base model for ResNet18/50 on CIFAR10. For
ResNet18/50 on CIFAR100 and Flowers102, we train the models from scratch using
the open-source scripts444Github repository: weiaicunzai/pytorch-cifar100. We
adopt their hyperparameter settings except that we resize the Flowers102
images to 128$\times$128.

  * •

We download ViT_B_16-224 checkpoint from an open-sourced Github
repository555Github repository: jeonsworld/ViT-pytorch and fine-tune it using
the scripts and hyperparameters provided in the same repository. All images
are resized to 224$\times$224 during training.

  * •

We download bert-base-cased checkpoint from HuggingFace666HuggingFace
checkpoint: google-bert/bert-base-cased, and use the default hyperparameters
provided in the sequence classification training script open-sourced in the
transformers repository777HuggingFace transformer: run_glue.py for sequence
classification.

##### Quantisation-aware lock

We use the same checkpoints and settings for quantisation-aware locking
experiments. The checkpoints were trained for 50 epochs with a quantisation-
aware loss using AdamW with a learning-rate of $1e^{-3}$, a batch size of 256,
and a random seed of $0$.

##  Appendix D Soft-locking Fine-Tuning

We use AdamW with a learning-rate of $1e^{-5}$ for both soft-locking
procedures. ViT-B_16-224 and the ResNet models were locked with batch-sizes of
32 and 256 respectively. Flowers102 images were resized to $128\times 128$ for
the ResNet models, and all datasets were resized to $224\times 224$ for ViT-
B_16-224. All runs were seed with random seed $0$.

##  Appendix E Sparsity-Aware Lock - Extended Table 5

We present below a finer version of 5 below, with two extra sparsity levels,
$p=0.10,0.25$, for completeness.

Table 10: Results are presented as
$\mathit{Acc}_{\text{authorised}}^{\text{locked}}({\color[rgb]{0.05078125,0.54296875,0.36328125}\definecolor[named]{pgfstrokecolor}{rgb}{0.05078125,0.54296875,0.36328125}\Delta_{\text{orig}}},{\color[rgb]{0.80078125,0,0}\definecolor[named]{pgfstrokecolor}{rgb}{0.80078125,0,0}\Delta_{\text{lock}}})$

Pruning Levels Dataset 0.05 0.10 0.25 0.50 0.75 BERT CoLA 0.80 (-0.03, 0.49)
0.83 (-0.01, 0.52) 0.84 (0.00, 0.53) 0.83 (0.00, 0.53) 0.84 (0.00, 0.53) MRPC
0.84 (-0.02, 0.15) 0.84 (-0.02, 0.16) 0.86 (0.00, 0.18) 0.86 (0.00, 0.54) 0.87
(0.00, 0.55) SST-2 0.92 (-0.01, 0.43) 0.92 (0.00, 0.43) 0.93 (0.00, 0.43) 0.92
(0.00, 0.43) 0.92 (0.00, 0.43) Resnet18 CIFAR10 0.89 (0.03, 0.67) 0.92 (0.00,
0.78) 0.93 (0.00, 0.83) 0.93 (0.00, 0.83) 0.93 (0.00, 0.83) CIFAR100 0.73
(0.03, 0.71) 0.75 (0.01, 0.74) 0.76 (0.00, 0.75) 0.76 (0.00, 0.75) 0.76 (0.13,
0.75) Flowers102 0.84 (0.04, 0.83) 0.84 (0.04, 0.84) 0.89 (0.00, 0.88) 0.89
(0.00, 0.88) 0.89 (0.00, 0.88) Resnet50 CIFAR10 0.92 (0.01, 0.81) 0.93 (0.00,
0.83) 0.93 (-0.01, 0.83) 0.94 (-0.01, 0.84) 0.94 (-0.01, 0.84) CIFAR100 0.69
(0.09, 0.63) 0.76 (0.02, 0.75) 0.78 (0.00, 0.77) 0.78 (0.00, 0.77) 0.78 (0.00,
0.77) Flowers102 0.76 (0.10, 0.74) 0.82 (0.04, 0.80) 0.86 (0.00, 0.84) 0.86
(0.00, 0.84) 0.85 (0.00, 0.84) ViT-B_16-224 CIFAR10 0.99 (0.00, 0.89) 0.99
(0.00, 0.89) 0.99 (0.00, 0.89) 0.99 (0.07, 0.89) 0.93 (-0.02, 0.93) CIFAR100
0.92 (-0.02, 0.92) 0.93 (-0.02, 0.92) 0.93 (-0.02, 0.92) 0.93 (-0.02, 0.93)
0.93 (-0.02, 0.93) Flowers102 1.00 (0.00, 0.98) 1.00 (0.00, 0.98) 1.00 (0.00,
0.99) 1.00 (0.00, 0.99) 1.00 (0.00, 1.00)

##  Appendix F Baseline Performance Degradation

###  F.1 Sparsity-aware lock

We present below the degradation in performance from the transfer of the dense
original model, i.e. before locking, to a higher-levels of sparsity, $p$.

Table 11: Results are presented as
$Acc_{\text{authorised}}(\Delta_{\text{base}}=Acc_{\text{authorised}}-Acc_{\text{un\\-authorised}})$

Pruning Levels Dataset 0.05 0.10 0.25 0.50 0.75 BERT MRPC 0.86 (0.00) 0.86
(0.00) 0.86 (0.01) 0.86 (0.18) 0.86 (0.54) SST-2 0.92 (-0.00) 0.92 (-0.00)
0.92 (-0.00) 0.92 (0.03) 0.92 (0.30) CoLA 0.84 (-0.00) 0.84 (-0.00) 0.84
(0.01) 0.84 (0.53) 0.84 (0.53) Resnet18 CIFAR10 0.93 (0.00) 0.93 (0.00) 0.93
(0.00) 0.93 (0.06) 0.93 (0.68) CIFAR100 0.76 (0.00) 0.76 (-0.00) 0.76 (0.01)
0.76 (0.07) 0.76 (0.08) Flowers102 0.89 (0.00) 0.89 (0.01) 0.89 (0.01) 0.89
(0.08) 0.89 (0.68) Resnet50 CIFAR10 0.93 (-0.00) 0.93 (-0.00) 0.93 (0.00) 0.93
(0.06) 0.93 (0.70) CIFAR100 0.78 (-0.00) 0.78 (-0.00) 0.78 (0.01) 0.78 (0.07)
0.78 (0.65) Flowers102 0.86 (0.00) 0.86 (0.00) 0.86 (0.01) 0.86 (0.10) 0.86
(0.66) ViT-B_16-224 CIFAR10 0.99 (-0.00) 0.99 (-0.00) 0.99 (0.00) 0.99 (0.07)
0.99 (0.00) CIFAR100 0.91 (0.00) 0.91 (0.01) 0.91 (0.01) 0.91 (0.30) 0.91
(0.90) Flowers102 1.00 (0.00) 1.00 (0.00) 1.00 (0.00) 1.00 (0.23) 1.00 (0.98)

As seen from Table 11, in most settings, models can consistently be deployed
at a pruning level of up to $p=0.5$, without incurring a notable cost in terms
of accuracy. In certain cases, for example ResNet18 on CIFAR100 and ViT-
B_16-224 on CIFAR10, even a sparsity level of $p=0.75$ does not significantly
affect performance. Table 11 therefore highlights that the significant
degradation in performance seen when deploying locked models on unauthorised
levels of sparsity 5 can attributed to sparsity-aware locking, and not merely
pruning.

###  F.2 Quantisation-aware lock

We present below the degradation in performance from the transfer of the
original model, i.e. before locking, to different arithmetics.

Table 12: Results are presented as $Acc_{\text{authorised}}(\Delta_{\text{base}}=Acc_{\text{authorised}}-Acc_{\text{un-authorised}})$ authorised $\rightarrow$ Unauthorised | Resnet18 | Resnet50
---|---|---
FP32 $\rightarrow$ 8-bit MiniFloat | 0.88 (0.00) | 0.88 (0.01)
Int8 $\rightarrow$ 8-bit MiniFloat | 0.88 (0.00) | 0.88 (0.01)
16-bit MiniFloat $\rightarrow$ Int8 | 0.88 (0.00) | 0.88 (0.01)
FP16 $\rightarrow$ Int8 | 0.88 (0.00) | 0.88 (-0.01)

There are negligible performance drops, if any, from quantising the base model
to an un-authorised arithmetic, as seen by the low $\Delta_{\text{base}}$
values. Note that this is as the base-model was partially trained with a
quantisation-aware loss at Int8. Like with sparsity-aware locking, Table 12
clearly shows that the degradation in accuracy when the locked model is
transferred to unauthorised arithmetic schemes 6 is a result of soft-locking,
and not merely from the differences in model representation across hardware.

###  F.3 Soft-Locking Optimisation Profiles

####  F.3.1 Sparsity-aware Lock

We present below the validation curves of the soft-locking optimisation
procedure for the three vision models for a subset of the datasets, CIFAR10,
CIFAR100, and pruning proportions, $p=\\{0.05,0.25,0.50\\}$.

Figure 3: ResNet18 locking curves Figure 4: ResNet50 locking curves Figure 5:
ViT-B_16-224 locking curves

The optimisation profiles generally converge much more rapidly, across
datasets and models, for higher pruning proportions, $p$. This reinforces the
intuitive result that the optimisation solved is easier for greater sparsity
values, and hence, larger values of $p$. For $p=0.05$, the performance of the
authorised locked model ($Acc_{authorised}^{locked}$) first reduces in
performance, suggesting the optimisation is dominated by the second term of
the manipulation loss function. Across the course of the optimisation, this
accuracy increases to an acceptable level, due to the eventual effect of the
first term of the loss function.

The notable exception to this trend is ViT-B_16-224. We posit that this is due
to its increased size (in number of parameters), which allows for a decoupling
of the opposing optimisation objectives, as the number of parameters at $5\%$
is much greater in ViT-B_16-224 than in the ResNet models. The larger number
of parameters can be used to effectively encode the differential knowledge
between the authorised and unauthorised variants of the model.

###  F.4 Quantisation-aware lock

We present below the validation curves of the quantisation-aware locking
procedure for the ResNet models on CIFAR10, CIFAR100 \- for the authorised,
unauthorised arithmetic pairs presented in 6. Note the failure of the locking
optimisation procedure for MiniFloat8 $\rightarrow$ Int8.

Figure 6: ResNet18 locking curves Figure 7: ResNet50 locking curves

##  Appendix G Appendix

In the following section, we present various miscellaneous findings regarding
the sparsity-aware lock.

###  G.1 Specificity of Sparsity-Aware Locking

We present below the the performance of locked model across a range of
sparsity levels for the ResNet models on CIFAR10.

Figure 8: Validation accuracy across sparsity levels (ResNet18 on CIFAR10)
Figure 9: Validation accuracy across sparsity levels (ResNet50 on CIFAR10)

In both models and all locking levels, $p_{\text{lock}}$, there is a highly
localised drop in accuracy at the sparsity-level that is locked. This is
desirable for setting in which the provider that is locking a proprietary
model is aware of the downstream sparsity-level utilised by unauthorised
users. If this is not known, or there is no single sparsity level that is
unauthorised, models can be conservatively locked at $p_{\text{lock}}$ = 0.05.
As seen in 8 and 9, locking the model at $p_{\text{lock}}$ = 0.05 effectively
renders transfer to any level of sparsity that enables efficient execution on
hardware useless.

###  G.2 Attacking soft-locking

If the unauthorised user has access to data, they may attack the soft-locking
procedures by re-training the models on the unauthorised hardware. For re-
training, it is natural for the user to employ a quantisation- or pruning-
aware loss and minimise the loss of the quantised or the pruned model.

To investigate the effectiveness of re-training, we re-trained a subset of the
locked vision models, namely Resnet18 and ResNet50, on CIFAR100. We present
below the re-training curves of the sparsity-locked models.

Figure 10: Re-training sparsity-locked ResNet18 on CIFAR100 Figure 11: Re-
training sparsity-locked ResNet50 on CIFAR100

In both models, re-training the models locked at $p=0.05$ and $p=0.25$ does
not return them to the accuracy of the unlocked model deployed on unauthorised
hardware. For the models locked at $p=0.50$, the majority of the accuracy is
recovered in roughly 5 epochs.

Below, we also present the training curves resulting from re-training the
sparsity-locked BERT models on the GLUE tasks, with a pruning-aware loss. The
fine-tuning curve of a base BERT model, at the same level of sparsity, with a
pruning-aware loss is presented as a baseline of reference. The recovery of
performance of locked models is comparable to the fine-tuning a sparse BERT
model from scratch. However, we note that GLUE tasks may not provide the
resolution that the previous vision tasks did, as 5 epochs are sufficient for
convergence in this setting.

Figure 12: (Re-)Fine-tuning of Sparsity-locked Models

###  G.3 Effect of Sparsity-Aware Locking on Prune Sum

To better understand the workings of the sparsity-locking scheme we
investigated the evolution of the prune sum, which is the sum of the absolute
values of the weights pruned. For brevity, we present the evolution of prune
sum for only ViT-B_16-224.

Figure 13: Evolution of Prune Sum during Sparsity-Locking (ViT-B_16-224)

Note that the $y$-axis scales of the different settings are vastly different.
Interestingly, this diagnostic illuminates that there are two regimes for the
sparsity-aware locking scheme, dependent on the value of $p$. For locking low-
levels of sparsity and $p$, the locking procedure increases the absolute
magnitudes of the $p$-smallest parameters which, to a first-order, increases
their importance in model inference. At larger values of $p$, increases to the
prune sum are negligible, if at all, suggesting the locking procedure does not
rely on the increasing of magnitudes.

##  Appendix H Parameter transform distributions

$-$0.10.00.10500000Original Parameters0.02.55.07.51e60100000200000Original
Significands0100200021e6Original Exponents$-$0.10.00.10241e6Transformed
Parameters0.02.55.07.51e60100000Transformed
Significands010020002000040000Transformed
Exponents$-$0.10.00.10500000Detransformed
Parameters0.02.55.07.51e60100000200000Detransformed
Significands0100200021e6Detransformed
Exponents$-$0.10.00.10241e6IncorrectlyDetransformed
Parameters0.02.55.07.51e60100000200000IncorrectlyDetransformed
Significands0100200050000IncorrectlyDetransformed Exponents

Figure 14: AES encryption transformation method

$-$0.10.00.10.00.51.01e6Original
Parameters0.02.55.07.51e60100000200000Original
Significands0100200021e6Original Exponents$-$0.10.00.10.00.51.01e6Transformed
Parameters0.02.55.07.51e60100000200000Transformed
Significands0100200021e6Transformed
Exponents$-$0.10.00.10.00.51.01e6Detransformed
Parameters0.02.55.07.51e60100000200000Detransformed
Significands0100200021e6Detransformed
Exponents$-$0.10.00.10.00.51.01e6IncorrectlyDetransformed
Parameters0.02.55.07.51e60100000200000IncorrectlyDetransformed
Significands0100200021e6IncorrectlyDetransformed Exponents

Figure 15: Shuffle transformation method

$-$0.10.00.10241e6Original Parameters050010000100000200000Original
Significands01020300121e6Original Exponents$-$0.10.00.10121e6Transformed
Parameters050010000100000Transformed Significands01020300200000Transformed
Exponents$-$0.10.00.10241e6Detransformed
Parameters050010000100000200000Detransformed
Significands01020300121e6Detransformed
Exponents$-$0.10.00.10241e6IncorrectlyDetransformed
Parameters050010000100000200000IncorrectlyDetransformed
Significands01020300121e6IncorrectlyDetransformed Exponents

Figure 16: Pretransformed AES encryption transformation method, with direct
estimation of distribution, on 16-bit floating point

$-$0.10.00.10.00.51.01e6Original
Parameters0.02.55.07.51e60100000200000Original
Significands0100200021e6Original Exponents$-$0.10.00.10241e6Transformed
Parameters0.02.55.07.51e60100000Transformed
Significands010020002000040000Transformed
Exponents$-$0.10.00.10.00.51.01e6Detransformed
Parameters0.02.55.07.51e60100000200000Detransformed
Significands0100200021e6Detransformed
Exponents$-$0.10.00.10250000500000IncorrectlyDetransformed
Parameters0.02.55.07.51e60100000200000IncorrectlyDetransformed
Significands0100200021e6IncorrectlyDetransformed Exponents

Figure 17: Pretransformed AES encryption transformation method, assuming
distribution to be Gaussian, on 32-bit floating point

$-$0.10.00.10.00.51.01e6Original
Parameters0.02.55.07.51e60100000200000Original
Significands0100200021e6Original Exponents$-$0.10.00.10241e6Transformed
Parameters0.02.55.07.51e60100000Transformed
Significands010020002000040000Transformed
Exponents$-$0.10.00.10.00.51.01e6Detransformed
Parameters0.02.55.07.51e60100000200000Detransformed
Significands0100200021e6Detransformed
Exponents$-$0.10.00.1011e6IncorrectlyDetransformed
Parameters0.02.55.07.51e60100000200000IncorrectlyDetransformed
Significands01002000121e6IncorrectlyDetransformed Exponents

Figure 18: Pretransformed AES encryption transformation method, directly
estimating distribution but assuming significant and exponent to be
independent, on 32-bit floating point

##  Appendix I Finite Precision Fingerprints

Table 13: Finite precision fingerprint on various devices

Device Fingerprint GTX 1080Ti
834a709ba2534ebe3ee1397fd4f7bd288b2acc1d20a08d6c862dcd99b6f04400 RTX 2080Ti
157c46f3245781907678d796984d27f21b3828d434c770ed85f7ac03bf11fe9f RTX 3090
e2b643bf633651bd14296a6205d7ef7e7e0dbb8c837d492e56695e3999eae638 RTX A6000
68085f05ecc499655bf7923c8a7f65a2a05e72ec67dfc1fb342377a04163c78f RTX A6000
(runpod.io) 5bcb85507944291de0716edaa3e77c6bf56d2fa173c1d207089d3db089638e83

##  Appendix J CUDA code to produce clock fingerprint

⬇

1/* Copyright (c) 2024, Eleanor Clifford

2 * Copyright (c) 2022, NVIDIA CORPORATION. All rights reserved.

3 *

4 * Redistribution and use in source and binary forms, with or without

5 * modification, are permitted provided that the following conditions

6 * are met:

7 * * Redistributions of source code must retain the above copyright

8 * notice, this list of conditions and the following disclaimer.

9 * * Redistributions in binary form must reproduce the above

10 * copyright notice, this list of conditions and the following

11 * disclaimer in the documentation and/or other materials

12 provided with the distribution.

13 * * Neither the name of NVIDIA CORPORATION nor the names of its

14 * contributors may be used to endorse or promote products

15 * derived from this software without specific prior written

16 * permission.

17 *

18 * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ‘‘AS IS’’ AND

19 * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO,

20 * THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A

21 * PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE

22 * COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT,

23 * INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES

24 * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR

25 * SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION)

26 * HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,

27 * STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)

28 * ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED

29 * OF THE POSSIBILITY OF SUCH DAMAGE.

30 */

31

32// System includes

33#include <assert.h>

34#include <stdint.h>

35#include <stdio.h>

36#include <ieee754.h>

37#include <limits.h>

38

39// CUDA runtime

40#include <cuda_runtime.h>

41

42// helper functions and utilities to work with CUDA

43#include <helper_cuda.h>

44#include <helper_functions.h>

45

46#define INNER_LOOP 512

47#define NUM_BLOCKS 128

48#define NUM_THREADS 8

49#define NUM_TRIES 16

50

51// This kernel does something which doesn’t matter.

52// The timing results are stored in device memory.

53__global__ static void timedFunction(

54 const float *input, float *output,

55 clock_t *timer

56) {

57 extern __shared__ float shared[];

58

59 const int tid = threadIdx.x;

60 const int bid = blockIdx.x;

61

62 // Copy input.

63 shared[tid] = input[tid];

64 shared[tid + blockDim.x] = input[tid + blockDim.x];

65

66 if (tid == 0) timer[bid] = clock();

67

68 // Do some stuff

69 for (size_t i = 0; i < INNER_LOOP; i++) {

70 for (int d = blockDim.x; d > 0; d /= 2) {

71 __syncthreads();

72

73 if (tid < d) {

74 float f0 = shared[tid];

75 float f1 = shared[tid + d];

76

77 if (f1 < f0) {

78 shared[tid] = (f0 + f1);

79 }

80 }

81 }

82 }

83

84 // Write result.

85 if (tid == 0) output[bid] = shared[0];

86 __syncthreads();

87 if (tid == 0) timer[bid + gridDim.x] = clock();

88

89}

90

91// Start the main CUDA Sample here

92int main(int argc, char **argv) {

93 cudaSetDevice(2);

94

95 float *dinput = NULL;

96 float *doutput = NULL;

97 clock_t *dtimer = NULL;

98

99 clock_t timer[NUM_BLOCKS * 2];

100 float input[NUM_THREADS * 2];

101

102 for (int i = 0; i < NUM_THREADS * 2; i++) {

103 input[i] = (float)i;

104 }

105

106 long fastestClock = LONG_MAX;

107 for (int j = 0; j < NUM_TRIES; j++) {

108

109 checkCudaErrors(

110 cudaMalloc((void **)&dinput, sizeof(float) * NUM_THREADS * 2)

111 );

112 checkCudaErrors(

113 cudaMalloc((void **)&doutput, sizeof(float) * NUM_BLOCKS)

114 );

115 checkCudaErrors(

116 cudaMalloc((void **)&dtimer, sizeof(clock_t) * NUM_BLOCKS * 2)

117 );

118

119 checkCudaErrors(

120 cudaMemcpy(dinput, input, sizeof(float) * NUM_THREADS * 2,

121 cudaMemcpyHostToDevice)

122 );

123

124 timedFunction<<<

125 NUM_BLOCKS, NUM_THREADS, sizeof(float) * 2 * NUM_THREADS

126 >>>(dinput, doutput, dtimer);

127

128 checkCudaErrors(

129 cudaMemcpy(timer, dtimer, sizeof(clock_t) * NUM_BLOCKS * 2,

130 cudaMemcpyDeviceToHost)

131 );

132

133 checkCudaErrors(cudaFree(dinput));

134 checkCudaErrors(cudaFree(doutput));

135 checkCudaErrors(cudaFree(dtimer));

136

137 for (int i = 0; i < NUM_BLOCKS; i++) {

138 long t = (timer[i + NUM_BLOCKS] - timer[i]);

139 if (t < fastestClock) {

140 fastestClock = t;

141 }

142 }

143 }

144

145 printf("%x\n", (unsigned int)fastestClock);

146

147 return EXIT_SUCCESS;

148}

##  Appendix K Python code to produce finite precision fingerprint

⬇

1import torch

2import hashlib

3import numpy as np

4

5dev = "cuda"

6

7torch.manual_seed(0)

8inp = torch.randn(1, 50, 50, 100)

9m = torch.nn.Sequential(

10 torch.nn.Linear(100, 1000),

11 torch.nn.Linear(1000, 10000),

12 torch.nn.Linear(10000, 10),

13)

14

15m = m.to(dev)

16inp = inp.to(dev)

17

18out = []

19

20orig = None

21with torch.no_grad():

22 for i in range(1, 10):

23 _input = torch.cat(i * [inp]).clone()

24 output = m(_input)

25

26 if orig is None:

27 orig = output[0].clone().detach()

28

29 out.append(float((orig - output).sum()))

30

31print(hashlib.sha256(np.array(out)).hexdigest())

##  Appendix L Experiment Compute Resources

The soft locking experiments consume most of the compute resources. We
conducted all sparsity-aware locking and quantisation-aware locking on NVIDIA
V100 GPUs and 18-core Intel Xeon (Broadwell) processors. The fine-tuning took
around 1 GPU hours per trial on average, and in total, the fine-tuning time
was around 180 GPU hours. We spent additional time on preliminary and failed
experiments, which is around 40 GPU hours in total. The emulation cost
experiments were performed on three NVIDIA RTXA6000 GPUs with an AMD EPYC 7713
64-core processor. The emulation cost experiments took around 4 GPU hours. The
hard locking experiments were conducted on NVIDIA GTX1080Ti, RTX2080Ti,
RTX3090, RTXA6000 GPUs, and took 10 GPU hours in total.