Getting started with binary analysis
This notebook shows how to get started with binary analysis.
Binary analysis examines compiled CUDA code.
The CUDA toolkit provides the CUDA binary utilities cuobjdump and nvdisasm for extracting and disassembling the contents of CUDA binary files.
ReProspect enables a fully programmatic use of these tools: it launches them, collects their output into Python data structures, and supports analyses of these data that can go all the way to test assertions.
As an example, we consider a kernel involving an atomic function call. We examine the CUDA assembly (SASS) instructions that this call is compiled to.
About this page
This page is rendered from a Jupyter notebook, executed at documentation build time, located at docs/source/getting-started/example_binary_analysis.ipynb in the repository.
To run the example yourself, download the notebook and open it in JupyterLab.
Alternatively, copy-paste the code snippets successively into an interactive Python session.
The requirements are Python 3.10 or newer and ReProspect.
A CUDA Toolkit 12.8 or newer installation providing nvcc, cuobjdump, and cu++filt is also needed.
Because the example is concerned with analysing the CUDA assembly code produced by compilation but does not run the kernel, no GPU is required.
Source code
Let us consider a kernel that atomically adds each element of a source array to the corresponding element of a destination array.
In this source code, atomicAdd is a legacy atomic function provided by CUDA as a C++ language extension, as described here.
It reads the value from a given memory address, adds a value to it, and writes the result back to the same memory address, all in one atomic operation.
As a legacy atomic function, it has memory_order_relaxed C++ atomic memory semantics.
Because the function name atomicAdd does not have a suffix, it is atomic at device scope.
It should be noted that atomicAdd also returns the value that was previously stored at the address, but this kernel does not use this return value.
CODE = """\
#include "cuda.h"
__global__ void my_kernel(int* __restrict__ const dst, const int* __restrict__ const src) {
const auto index = blockIdx.x * blockDim.x + threadIdx.x;
atomicAdd(&dst[index], src[index]);
}
"""
Compilation
We compile the source code for three targets spanning the Ampere architecture (compute capability 8.0, target sm_80), Hopper (9.0, sm_90), and Blackwell (12.0, sm_120).
Later in this notebook, the atomic function call will be found to compile to different SASS instructions across these targets.
import pathlib
import subprocess
from reprospect.tools import architecture
ARCHES = [
architecture.NVIDIAArch.from_compute_capability(80),
architecture.NVIDIAArch.from_compute_capability(90),
architecture.NVIDIAArch.from_compute_capability(120),
]
print(subprocess.check_output(('nvcc', '--version')).decode().strip())
workdir = pathlib.Path.cwd() / 'example_binary_analysis'
workdir.mkdir(exist_ok=True)
print(f'Working directory: {workdir}')
for arch in ARCHES:
source = workdir / f'atomic.{arch.as_sm}.cu'
output = workdir / f'atomic.{arch.as_sm}'
source.write_text(CODE)
subprocess.check_call(('nvcc', f'--generate-code=arch={arch.as_compute},code=[{arch.as_sm}]', '-O3', '-c', source, '-o', output))
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2025 NVIDIA Corporation
Built on Fri_Nov__7_07:23:37_PM_PST_2025
Cuda compilation tools, release 13.1, V13.1.80
Build cuda_13.1.r13.1/compiler.36836380_0
Working directory: /tmp/tmp93kl3mdy/example_binary_analysis
Extracting and disassembling CUDA binary code
By default, the CUDA compiler driver nvcc embeds CUDA binary files into the compiled output file, as described here.
A CUDA binary file is also referred to as a cubin.
Because the example involves only a single kernel, compiled into a separate output file for each target, there is only a single embedded cubin for each target.
We will now extract this embedded cubin from the compiled output file and disassemble its CUDA binary code for each target.
The CUDA toolkit provides the CUDA binary utility cuobjdump for this purpose.
Here, we invoke it through the ReProspect class reprospect.tools.binaries.cuobjdump.CuObjDump.
The method extract() is a factory method: it extracts, from the compiled output file, the embedded cubin for the target architecture, and it disassembles the SASS instructions of each kernel in this cubin.
The returned object has a member functions, a dictionary that maps each kernel’s signature to its SASS code as a string.
The kernel signature is decoded (demangled) from its low-level assembly name.
Even though there is only a single kernel in this example and this kernel could therefore also be accessed as the first item in the dictionary, we will use this demangled signature below to access our kernel.
It should be noted that the cubins embedded in a compiled output file follow a naming convention as in <basename>.<index>.<sm_version>.cubin.
The ReProspect class CuObjDump has a method embedded_cubins() to inspect the names of the cubins embedded in a compiled output file.
from reprospect.tools.binaries import CuObjDump
cuobjdump = {}
for arch in ARCHES:
compiled_output = workdir / f'atomic.{arch.as_sm}'
cuobjdump[arch], _ = CuObjDump.extract(
file=compiled_output,
arch=arch,
cwd=workdir,
cubin=f'atomic.1.{arch.as_sm}.cubin',
)
print(cuobjdump[arch].functions.keys())
dict_keys(['my_kernel(int *, const int *)'])
dict_keys(['my_kernel(int *, const int *)'])
dict_keys(['my_kernel(int *, const int *)'])
Like many classes in ReProspect, the class reprospect.tools.binaries.cuobjdump.CuObjDump supports rich rendering.
Printing shows for each kernel its demangled signature, its SASS code, and information about its resource usage, such as the number of general purpose registers per thread that the kernel uses.
for arch in ARCHES:
print(cuobjdump[arch])
CuObjDump of /tmp/tmp93kl3mdy/example_binary_analysis/atomic.1.sm_80.cubin for architecture AMPERE80:
┌─────────────────┬────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────┐
│ Function │ my_kernel(int *, const int *) │
├─────────────────┼────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────┤
│ Symbol │ _Z9my_kernelPiPKi │
├─────────────────┼────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────┤
│ Code │ .headerflags @"EF_CUDA_SM80 EF_CUDA_VIRTUAL_SM(EF_CUDA_SM80)" │
│ │ /*0000*/ MOV R1, c[0x0][0x28] ; /* 0x00000a0000017a02 */ │
│ │ /* 0x000fe40000000f00 */ │
│ │ /*0010*/ S2R R2, SR_CTAID.X ; /* 0x0000000000027919 */ │
│ │ /* 0x000e220000002500 */ │
│ │ /*0020*/ HFMA2.MMA R7, -RZ, RZ, 0, 2.384185791015625e-07 ; /* 0x00000004ff077435 */ │
│ │ /* 0x000fe200000001ff */ │
│ │ /*0030*/ ULDC.64 UR4, c[0x0][0x118] ; /* 0x0000460000047ab9 */ │
│ │ /* 0x000fe40000000a00 */ │
│ │ /*0040*/ S2R R3, SR_TID.X ; /* 0x0000000000037919 */ │
│ │ /* 0x000e240000002100 */ │
│ │ /*0050*/ IMAD R2, R2, c[0x0][0x0], R3 ; /* 0x0000000002027a24 */ │
│ │ /* 0x001fca00078e0203 */ │
│ │ /*0060*/ IMAD.WIDE.U32 R4, R2, R7, c[0x0][0x168] ; /* 0x00005a0002047625 */ │
│ │ /* 0x000fcc00078e0007 */ │
│ │ /*0070*/ LDG.E.CONSTANT R5, [R4.64] ; /* 0x0000000404057981 */ │
│ │ /* 0x000ea2000c1e9900 */ │
│ │ /*0080*/ IMAD.WIDE.U32 R2, R2, R7, c[0x0][0x160] ; /* 0x0000580002027625 */ │
│ │ /* 0x000fca00078e0007 */ │
│ │ /*0090*/ RED.E.ADD.STRONG.GPU [R2.64], R5 ; /* 0x000000050200798e */ │
│ │ /* 0x004fe2000c10e184 */ │
│ │ /*00a0*/ EXIT ; /* 0x000000000000794d */ │
│ │ /* 0x000fea0003800000 */ │
│ │ /*00b0*/ BRA 0xb0; /* 0xfffffff000007947 */ │
│ │ /* 0x000fc0000383ffff */ │
│ │ /*00c0*/ NOP; /* 0x0000000000007918 */ │
│ │ /* 0x000fc00000000000 */ │
│ │ /*00d0*/ NOP; /* 0x0000000000007918 */ │
│ │ /* 0x000fc00000000000 */ │
│ │ /*00e0*/ NOP; /* 0x0000000000007918 */ │
│ │ /* 0x000fc00000000000 */ │
│ │ /*00f0*/ NOP; /* 0x0000000000007918 */ │
│ │ /* 0x000fc00000000000 */ │
│ │ /*0100*/ NOP; /* 0x0000000000007918 */ │
│ │ /* 0x000fc00000000000 */ │
│ │ /*0110*/ NOP; /* 0x0000000000007918 */ │
│ │ /* 0x000fc00000000000 */ │
│ │ /*0120*/ NOP; /* 0x0000000000007918 */ │
│ │ /* 0x000fc00000000000 */ │
│ │ /*0130*/ NOP; /* 0x0000000000007918 */ │
│ │ /* 0x000fc00000000000 */ │
│ │ /*0140*/ NOP; /* 0x0000000000007918 */ │
│ │ /* 0x000fc00000000000 */ │
│ │ /*0150*/ NOP; /* 0x0000000000007918 */ │
│ │ /* 0x000fc00000000000 */ │
│ │ /*0160*/ NOP; /* 0x0000000000007918 */ │
│ │ /* 0x000fc00000000000 */ │
│ │ /*0170*/ NOP; /* 0x0000000000007918 */ │
│ │ /* 0x000fc00000000000 */ │
├─────────────────┼────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────┤
│ Resource usage │ REG: 10, STACK: 0, SHARED: 0, LOCAL: 0, CONSTANT: {0: 368}, TEXTURE: 0, SURFACE: 0, SAMPLER: 0 │
└─────────────────┴────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────┘
CuObjDump of /tmp/tmp93kl3mdy/example_binary_analysis/atomic.1.sm_90.cubin for architecture HOPPER90:
┌─────────────────┬────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────┐
│ Function │ my_kernel(int *, const int *) │
├─────────────────┼────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────┤
│ Symbol │ _Z9my_kernelPiPKi │
├─────────────────┼────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────┤
│ Code │ .headerflags @"EF_CUDA_SM90 EF_CUDA_VIRTUAL_SM(EF_CUDA_SM90)" │
│ │ /*0000*/ LDC R1, c[0x0][0x28] ; /* 0x00000a00ff017b82 */ │
│ │ /* 0x000fe20000000800 */ │
│ │ /*0010*/ S2R R0, SR_TID.X ; /* 0x0000000000007919 */ │
│ │ /* 0x000e2e0000002100 */ │
│ │ /*0020*/ S2UR UR4, SR_CTAID.X ; /* 0x00000000000479c3 */ │
│ │ /* 0x000e300000002500 */ │
│ │ /*0030*/ LDC R7, c[0x0][RZ] ; /* 0x00000000ff077b82 */ │
│ │ /* 0x000e300000000800 */ │
│ │ /*0040*/ LDC.64 R4, c[0x0][0x218] ; /* 0x00008600ff047b82 */ │
│ │ /* 0x000e700000000a00 */ │
│ │ /*0050*/ LDC.64 R2, c[0x0][0x210] ; /* 0x00008400ff027b82 */ │
│ │ /* 0x000ea20000000a00 */ │
│ │ /*0060*/ IMAD R7, R7, UR4, R0 ; /* 0x0000000407077c24 */ │
│ │ /* 0x001fe2000f8e0200 */ │
│ │ /*0070*/ ULDC.64 UR4, c[0x0][0x208] ; /* 0x0000820000047ab9 */ │
│ │ /* 0x000fc60000000a00 */ │
│ │ /*0080*/ IMAD.WIDE.U32 R4, R7, 0x4, R4 ; /* 0x0000000407047825 */ │
│ │ /* 0x002fcc00078e0004 */ │
│ │ /*0090*/ LDG.E.CONSTANT R5, desc[UR4][R4.64] ; /* 0x0000000404057981 */ │
│ │ /* 0x000ee2000c1e9900 */ │
│ │ /*00a0*/ IMAD.WIDE.U32 R2, R7, 0x4, R2 ; /* 0x0000000407027825 */ │
│ │ /* 0x004fca00078e0002 */ │
│ │ /*00b0*/ REDG.E.ADD.STRONG.GPU desc[UR4][R2.64], R5 ; /* 0x000000050200798e */ │
│ │ /* 0x008fe2000c10e184 */ │
│ │ /*00c0*/ EXIT ; /* 0x000000000000794d */ │
│ │ /* 0x000fea0003800000 */ │
│ │ /*00d0*/ BRA 0xd0; /* 0xfffffffc00fc7947 */ │
│ │ /* 0x000fc0000383ffff */ │
│ │ /*00e0*/ NOP; /* 0x0000000000007918 */ │
│ │ /* 0x000fc00000000000 */ │
│ │ /*00f0*/ NOP; /* 0x0000000000007918 */ │
│ │ /* 0x000fc00000000000 */ │
│ │ /*0100*/ NOP; /* 0x0000000000007918 */ │
│ │ /* 0x000fc00000000000 */ │
│ │ /*0110*/ NOP; /* 0x0000000000007918 */ │
│ │ /* 0x000fc00000000000 */ │
│ │ /*0120*/ NOP; /* 0x0000000000007918 */ │
│ │ /* 0x000fc00000000000 */ │
│ │ /*0130*/ NOP; /* 0x0000000000007918 */ │
│ │ /* 0x000fc00000000000 */ │
│ │ /*0140*/ NOP; /* 0x0000000000007918 */ │
│ │ /* 0x000fc00000000000 */ │
│ │ /*0150*/ NOP; /* 0x0000000000007918 */ │
│ │ /* 0x000fc00000000000 */ │
│ │ /*0160*/ NOP; /* 0x0000000000007918 */ │
│ │ /* 0x000fc00000000000 */ │
│ │ /*0170*/ NOP; /* 0x0000000000007918 */ │
│ │ /* 0x000fc00000000000 */ │
├─────────────────┼────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────┤
│ Resource usage │ REG: 10, STACK: 0, SHARED: 0, LOCAL: 0, CONSTANT: {0: 544}, TEXTURE: 0, SURFACE: 0, SAMPLER: 0 │
└─────────────────┴────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────┘
CuObjDump of /tmp/tmp93kl3mdy/example_binary_analysis/atomic.1.sm_120.cubin for architecture BLACKWELL120:
┌─────────────────┬────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────┐
│ Function │ my_kernel(int *, const int *) │
├─────────────────┼────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────┤
│ Symbol │ _Z9my_kernelPiPKi │
├─────────────────┼────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────┤
│ Code │ .headerflags @"EF_CUDA_SM120 EF_CUDA_VIRTUAL_SM(EF_CUDA_SM120)" │
│ │ /*0000*/ LDC R1, c[0x0][0x37c] ; /* 0x0000df00ff017b82 */ │
│ │ /* 0x000fe20000000800 */ │
│ │ /*0010*/ S2R R0, SR_TID.X ; /* 0x0000000000007919 */ │
│ │ /* 0x000e2e0000002100 */ │
│ │ /*0020*/ S2UR UR6, SR_CTAID.X ; /* 0x00000000000679c3 */ │
│ │ /* 0x000e220000002500 */ │
│ │ /*0030*/ LDCU.64 UR4, c[0x0][0x358] ; /* 0x00006b00ff0477ac */ │
│ │ /* 0x000e6e0008000a00 */ │
│ │ /*0040*/ LDC R7, c[0x0][0x360] ; /* 0x0000d800ff077b82 */ │
│ │ /* 0x000e300000000800 */ │
│ │ /*0050*/ LDC.64 R4, c[0x0][0x388] ; /* 0x0000e200ff047b82 */ │
│ │ /* 0x000eb00000000a00 */ │
│ │ /*0060*/ LDC.64 R2, c[0x0][0x380] ; /* 0x0000e000ff027b82 */ │
│ │ /* 0x000ee20000000a00 */ │
│ │ /*0070*/ IMAD R7, R7, UR6, R0 ; /* 0x0000000607077c24 */ │
│ │ /* 0x001fc8000f8e0200 */ │
│ │ /*0080*/ IMAD.WIDE.U32 R4, R7, 0x4, R4 ; /* 0x0000000407047825 */ │
│ │ /* 0x004fcc00078e0004 */ │
│ │ /*0090*/ LDG.E.CONSTANT R5, desc[UR4][R4.64] ; /* 0x0000000404057981 */ │
│ │ /* 0x002ea2000c1e9900 */ │
│ │ /*00a0*/ IMAD.WIDE.U32 R2, R7, 0x4, R2 ; /* 0x0000000407027825 */ │
│ │ /* 0x008fca00078e0002 */ │
│ │ /*00b0*/ REDG.E.ADD.STRONG.GPU desc[UR4][R2.64], R5 ; /* 0x000000050200798e */ │
│ │ /* 0x004fe2000c12e104 */ │
│ │ /*00c0*/ EXIT ; /* 0x000000000000794d */ │
│ │ /* 0x000fea0003800000 */ │
│ │ /*00d0*/ BRA 0xd0; /* 0xfffffffc00fc7947 */ │
│ │ /* 0x000fc0000383ffff */ │
│ │ /*00e0*/ NOP; /* 0x0000000000007918 */ │
│ │ /* 0x000fc00000000000 */ │
│ │ /*00f0*/ NOP; /* 0x0000000000007918 */ │
│ │ /* 0x000fc00000000000 */ │
│ │ /*0100*/ NOP; /* 0x0000000000007918 */ │
│ │ /* 0x000fc00000000000 */ │
│ │ /*0110*/ NOP; /* 0x0000000000007918 */ │
│ │ /* 0x000fc00000000000 */ │
│ │ /*0120*/ NOP; /* 0x0000000000007918 */ │
│ │ /* 0x000fc00000000000 */ │
│ │ /*0130*/ NOP; /* 0x0000000000007918 */ │
│ │ /* 0x000fc00000000000 */ │
│ │ /*0140*/ NOP; /* 0x0000000000007918 */ │
│ │ /* 0x000fc00000000000 */ │
│ │ /*0150*/ NOP; /* 0x0000000000007918 */ │
│ │ /* 0x000fc00000000000 */ │
│ │ /*0160*/ NOP; /* 0x0000000000007918 */ │
│ │ /* 0x000fc00000000000 */ │
│ │ /*0170*/ NOP; /* 0x0000000000007918 */ │
│ │ /* 0x000fc00000000000 */ │
├─────────────────┼────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────┤
│ Resource usage │ REG: 10, STACK: 0, SHARED: 0, LOCAL: 0, CONSTANT: {0: 912}, TEXTURE: 0, SURFACE: 0, SAMPLER: 0 │
└─────────────────┴────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────┘
Decoding SASS instructions
Next, we use ReProspect to parse the SASS instructions.
ReProspect provides the class reprospect.tools.binaries.sass.decoder.Decoder for this purpose.
The snippet below constructs for each target an instance of this class from the SASS code that was stored as a string previously.
This instance has a member instructions, a list wherein each SASS instruction is represented by its offset, its hexadecimal encoding, its disassembled representation as a string, and its decoded control code (see also ControlCode).
The further decomposition of the disassembled representation as a string into its components (instruction predicate, opcode, modifiers, and operands) is performed by the instruction matchers introduced below.
from reprospect.tools.binaries.sass import Decoder
SIGNATURE = 'my_kernel(int *, const int *)'
decoder: dict[architecture.NVIDIAArch, Decoder] = {}
for arch in ARCHES:
decoder[arch] = Decoder(code=cuobjdump[arch].functions[SIGNATURE].code)
The class Decoder supports rich rendering.
Printing shows for each SASS instruction its offset, its disassembled representation as a string, and its decoded control code.
for arch in ARCHES:
print(decoder[arch])
┏━━━━━━━━┳━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━┳━━━━━━━┳━━━━━━━┳━━━━┳━━━━┳━━━━┳━━━━┳━━━━┳━━━━┓
┃ offset ┃ instruction ┃ stall ┃ yield ┃ b0 ┃ b1 ┃ b2 ┃ b3 ┃ b4 ┃ b5 ┃
┡━━━━━━━━╇━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━╇━━━━━━━╇━━━━━━━╇━━━━╇━━━━╇━━━━╇━━━━╇━━━━╇━━━━┩
│ 0000 │ MOV R1, c[0x0][0x28] │ 2 │ True │ │ │ │ │ │ │
│ 0010 │ S2R R2, SR_CTAID.X │ 1 │ True │ Wr │ │ │ │ │ │
│ 0020 │ HFMA2.MMA R7, -RZ, RZ, 0, 2.384185791015625e-07 │ 1 │ True │ │ │ │ │ │ │
│ 0030 │ ULDC.64 UR4, c[0x0][0x118] │ 2 │ True │ │ │ │ │ │ │
│ 0040 │ S2R R3, SR_TID.X │ 2 │ True │ Wr │ │ │ │ │ │
│ 0050 │ IMAD R2, R2, c[0x0][0x0], R3 │ 5 │ False │ Wa │ │ │ │ │ │
│ 0060 │ IMAD.WIDE.U32 R4, R2, R7, c[0x0][0x168] │ 6 │ False │ │ │ │ │ │ │
│ 0070 │ LDG.E.CONSTANT R5, [R4.64] │ 1 │ True │ │ │ Wr │ │ │ │
│ 0080 │ IMAD.WIDE.U32 R2, R2, R7, c[0x0][0x160] │ 5 │ False │ │ │ │ │ │ │
│ 0090 │ RED.E.ADD.STRONG.GPU [R2.64], R5 │ 1 │ True │ │ │ Wa │ │ │ │
│ 00a0 │ EXIT │ 5 │ True │ │ │ │ │ │ │
│ 00b0 │ BRA 0xb0 │ 0 │ False │ │ │ │ │ │ │
│ 00c0 │ NOP │ 0 │ False │ │ │ │ │ │ │
│ 00d0 │ NOP │ 0 │ False │ │ │ │ │ │ │
│ 00e0 │ NOP │ 0 │ False │ │ │ │ │ │ │
│ 00f0 │ NOP │ 0 │ False │ │ │ │ │ │ │
│ 0100 │ NOP │ 0 │ False │ │ │ │ │ │ │
│ 0110 │ NOP │ 0 │ False │ │ │ │ │ │ │
│ 0120 │ NOP │ 0 │ False │ │ │ │ │ │ │
│ 0130 │ NOP │ 0 │ False │ │ │ │ │ │ │
│ 0140 │ NOP │ 0 │ False │ │ │ │ │ │ │
│ 0150 │ NOP │ 0 │ False │ │ │ │ │ │ │
│ 0160 │ NOP │ 0 │ False │ │ │ │ │ │ │
│ 0170 │ NOP │ 0 │ False │ │ │ │ │ │ │
└────────┴─────────────────────────────────────────────────┴───────┴───────┴────┴────┴────┴────┴────┴────┘
┏━━━━━━━━┳━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━┳━━━━━━━┳━━━━━━━┳━━━━┳━━━━┳━━━━┳━━━━┳━━━━┳━━━━┓
┃ offset ┃ instruction ┃ stall ┃ yield ┃ b0 ┃ b1 ┃ b2 ┃ b3 ┃ b4 ┃ b5 ┃
┡━━━━━━━━╇━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━╇━━━━━━━╇━━━━━━━╇━━━━╇━━━━╇━━━━╇━━━━╇━━━━╇━━━━┩
│ 0000 │ LDC R1, c[0x0][0x28] │ 1 │ True │ │ │ │ │ │ │
│ 0010 │ S2R R0, SR_TID.X │ 7 │ True │ Wr │ │ │ │ │ │
│ 0020 │ S2UR UR4, SR_CTAID.X │ 8 │ True │ Wr │ │ │ │ │ │
│ 0030 │ LDC R7, c[0x0][RZ] │ 8 │ True │ Wr │ │ │ │ │ │
│ 0040 │ LDC.64 R4, c[0x0][0x218] │ 8 │ True │ │ Wr │ │ │ │ │
│ 0050 │ LDC.64 R2, c[0x0][0x210] │ 1 │ True │ │ │ Wr │ │ │ │
│ 0060 │ IMAD R7, R7, UR4, R0 │ 1 │ True │ Wa │ │ │ │ │ │
│ 0070 │ ULDC.64 UR4, c[0x0][0x208] │ 3 │ False │ │ │ │ │ │ │
│ 0080 │ IMAD.WIDE.U32 R4, R7, 0x4, R4 │ 6 │ False │ │ Wa │ │ │ │ │
│ 0090 │ LDG.E.CONSTANT R5, desc[UR4][R4.64] │ 1 │ True │ │ │ │ Wr │ │ │
│ 00a0 │ IMAD.WIDE.U32 R2, R7, 0x4, R2 │ 5 │ False │ │ │ Wa │ │ │ │
│ 00b0 │ REDG.E.ADD.STRONG.GPU desc[UR4][R2.64], R5 │ 1 │ True │ │ │ │ Wa │ │ │
│ 00c0 │ EXIT │ 5 │ True │ │ │ │ │ │ │
│ 00d0 │ BRA 0xd0 │ 0 │ False │ │ │ │ │ │ │
│ 00e0 │ NOP │ 0 │ False │ │ │ │ │ │ │
│ 00f0 │ NOP │ 0 │ False │ │ │ │ │ │ │
│ 0100 │ NOP │ 0 │ False │ │ │ │ │ │ │
│ 0110 │ NOP │ 0 │ False │ │ │ │ │ │ │
│ 0120 │ NOP │ 0 │ False │ │ │ │ │ │ │
│ 0130 │ NOP │ 0 │ False │ │ │ │ │ │ │
│ 0140 │ NOP │ 0 │ False │ │ │ │ │ │ │
│ 0150 │ NOP │ 0 │ False │ │ │ │ │ │ │
│ 0160 │ NOP │ 0 │ False │ │ │ │ │ │ │
│ 0170 │ NOP │ 0 │ False │ │ │ │ │ │ │
└────────┴────────────────────────────────────────────┴───────┴───────┴────┴────┴────┴────┴────┴────┘
┏━━━━━━━━┳━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━┳━━━━━━━┳━━━━━━━┳━━━━┳━━━━┳━━━━┳━━━━┳━━━━┳━━━━┓
┃ offset ┃ instruction ┃ stall ┃ yield ┃ b0 ┃ b1 ┃ b2 ┃ b3 ┃ b4 ┃ b5 ┃
┡━━━━━━━━╇━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━╇━━━━━━━╇━━━━━━━╇━━━━╇━━━━╇━━━━╇━━━━╇━━━━╇━━━━┩
│ 0000 │ LDC R1, c[0x0][0x37c] │ 1 │ True │ │ │ │ │ │ │
│ 0010 │ S2R R0, SR_TID.X │ 7 │ True │ Wr │ │ │ │ │ │
│ 0020 │ S2UR UR6, SR_CTAID.X │ 1 │ True │ Wr │ │ │ │ │ │
│ 0030 │ LDCU.64 UR4, c[0x0][0x358] │ 7 │ True │ │ Wr │ │ │ │ │
│ 0040 │ LDC R7, c[0x0][0x360] │ 8 │ True │ Wr │ │ │ │ │ │
│ 0050 │ LDC.64 R4, c[0x0][0x388] │ 8 │ True │ │ │ Wr │ │ │ │
│ 0060 │ LDC.64 R2, c[0x0][0x380] │ 1 │ True │ │ │ │ Wr │ │ │
│ 0070 │ IMAD R7, R7, UR6, R0 │ 4 │ False │ Wa │ │ │ │ │ │
│ 0080 │ IMAD.WIDE.U32 R4, R7, 0x4, R4 │ 6 │ False │ │ │ Wa │ │ │ │
│ 0090 │ LDG.E.CONSTANT R5, desc[UR4][R4.64] │ 1 │ True │ │ Wa │ Wr │ │ │ │
│ 00a0 │ IMAD.WIDE.U32 R2, R7, 0x4, R2 │ 5 │ False │ │ │ │ Wa │ │ │
│ 00b0 │ REDG.E.ADD.STRONG.GPU desc[UR4][R2.64], R5 │ 1 │ True │ │ │ Wa │ │ │ │
│ 00c0 │ EXIT │ 5 │ True │ │ │ │ │ │ │
│ 00d0 │ BRA 0xd0 │ 0 │ False │ │ │ │ │ │ │
│ 00e0 │ NOP │ 0 │ False │ │ │ │ │ │ │
│ 00f0 │ NOP │ 0 │ False │ │ │ │ │ │ │
│ 0100 │ NOP │ 0 │ False │ │ │ │ │ │ │
│ 0110 │ NOP │ 0 │ False │ │ │ │ │ │ │
│ 0120 │ NOP │ 0 │ False │ │ │ │ │ │ │
│ 0130 │ NOP │ 0 │ False │ │ │ │ │ │ │
│ 0140 │ NOP │ 0 │ False │ │ │ │ │ │ │
│ 0150 │ NOP │ 0 │ False │ │ │ │ │ │ │
│ 0160 │ NOP │ 0 │ False │ │ │ │ │ │ │
│ 0170 │ NOP │ 0 │ False │ │ │ │ │ │ │
└────────┴────────────────────────────────────────────┴───────┴───────┴────┴────┴────┴────┴────┴────┘
We can observe that the kernel defined above compiles to a sequence of SASS instructions, including load/store instructions (such as those with opcode LDCU, LDC, and LDG in the case of the sm_120 target), and integer instructions (such as those with opcode IMAD for sm_120, used here for address computation).
Because its return value is unused, the compiler can lower atomicAdd to a reduction operation instruction.
With the used CUDA Toolkit version, -O3, and the selected targets, it emits RED for sm_80 and REDG for sm_90 and sm_120; NVIDIA describes both opcodes as Reduction Operations on Generic Memory.
We can assert this behavior by checking that the decoded instruction list contains exactly one instruction whose disassembled representation as a string contains the expected opcode. It should be noted that when proceeding in this way, the opcode must be hard-coded per target architecture. The matchers introduced next will abstract such details away.
instruction_red = {}
for arch in ARCHES:
opcode = 'RED' if arch.compute_capability < 90 else 'REDG'
instruction_red[arch] = [
instr for instr in decoder[arch].instructions
if f'{opcode}.' in instr.instruction
]
assert len(instruction_red[arch]) == 1
Decomposing SASS instructions
ReProspect provides an extensible matching framework for SASS instructions.
At the lowest levels, matchers analyse SASS instructions and their components (instruction predicate, opcode, modifiers, and operands).
from reprospect.testing.binaries.sass.instruction import AnyMatcher
for arch in ARCHES:
decomposed = AnyMatcher().match(inst=instruction_red[arch][0].instruction)
print(f'Arch: {arch.as_sm}, Instruction: {decomposed}')
Arch: sm_80, Instruction: InstructionMatch(opcode='RED', modifiers=('E', 'ADD', 'STRONG', 'GPU'), operands=('[R2.64]', 'R5'), predicate=None, additional=None)
Arch: sm_90, Instruction: InstructionMatch(opcode='REDG', modifiers=('E', 'ADD', 'STRONG', 'GPU'), operands=('desc[UR4][R2.64]', 'R5'), predicate=None, additional=None)
Arch: sm_120, Instruction: InstructionMatch(opcode='REDG', modifiers=('E', 'ADD', 'STRONG', 'GPU'), operands=('desc[UR4][R2.64]', 'R5'), predicate=None, additional=None)
We can observe that for each considered target, the modifiers are E, ADD, STRONG and GPU.
These modifiers indicate extended addressing (E), the reduction operation (ADD), the STRONG modifier, and the scope (GPU, i.e., Device).
NVIDIA documents SASS instructions only partially.
The interpretation of STRONG should be stated cautiously.
It is compatible with NVIDIA’s memory consistency model, in which relaxed memory operations are classified as strong operations.
NVIDIA does fully document the PTX instructions, which compilers generate before ultimately translating them to native target-architecture-specific SASS instructions.
Consulting the PTX instruction set reference can thus also provide insight into the ultimately emitted SASS instructions; see, for instance, this description of the PTX red instruction.
Further, it should be noted that although the modifiers coincide across the three targets here, modifiers may depend on the target architecture in general.
ReProspect’s matchers are implemented internally in terms of regex matching.
print(AnyMatcher().PATTERN)
regex.Regex('(?:(?P<predicate>@!?U?P(?:T|[0-9]+)))?\\s*(?P<opcode>[A-Z0-9]+)(?:\\.(?P<modifiers>[A-Z0-9_]+))*\\s*(?:(?P<operands>[\\w!\\.\\[\\]\\+\\-\\|~]+)(?:,?\\s*(?P<operands>[\\w!\\.\\[\\]\\+\\-\\|~]+))*)?', flags=regex.V0)
Matching SASS instructions across target architectures
ReProspect provides SASS instruction matchers that abstract away target-architecture-specific details.
Here, we use reprospect.testing.binaries.sass.instruction.atomic.ReductionMatcher, whose internal implementation encodes how the components of a Reduction Operation on Generic Memory SASS instruction depend on the target architecture.
It builds its regex pattern as a function of the target architecture and of the requested reduction operation, scope, consistency, and value type.
import numpy
from reprospect.testing.binaries.sass.instruction import ReductionMatcher, ThreadScope
for arch in ARCHES:
matcher_red = ReductionMatcher(
arch=arch, operation='ADD',
scope=ThreadScope.DEVICE,
consistency='STRONG', dtype=numpy.int32,
)
print(f'Arch: {arch.as_sm}, Pattern: {matcher_red.pattern}')
matched_red = [
(inst, matched)
for inst in decoder[arch].instructions
if (matched := matcher_red.match(inst))
]
assert len(matched_red) == 1
print(f'Arch: {arch.as_sm}, Instruction: {matched_red}')
Arch: sm_80, Pattern: regex.Regex('(?P<opcode>RED)\\.(?P<modifiers>E)\\.(?P<modifiers>ADD)\\.(?P<modifiers>STRONG)\\.(?P<modifiers>GPU) (?P<address>(?P<operands>(?:\\[(?:R[0-9]+|UR[0-9]+)\\.64(?:\\+(?:-?0x[0-9A-Fa-f]+|UR[0-9]+))?\\]|desc\\[UR[0-9]+\\]\\[(?:R[0-9]+|UR[0-9]+)\\.64(?:\\+(?:-?0x[0-9A-Fa-f]+|UR[0-9]+))?\\]))), (?P<operands>R[0-9]+)', flags=regex.V0)
Arch: sm_80, Instruction: [(Instruction(offset=144, instruction='RED.E.ADD.STRONG.GPU [R2.64], R5', hex='0x000000050200798e', control=ControlCode(stall_count=1, yield_flag=True, read=7, write=7, wait=[False, False, True, False, False, False], reuse={'A': False, 'B': False, 'C': False, 'D': False})), InstructionMatch(opcode='RED', modifiers=('E', 'ADD', 'STRONG', 'GPU'), operands=('[R2.64]', 'R5'), predicate=None, additional={'address': ['[R2.64]']}))]
Arch: sm_90, Pattern: regex.Regex('(?P<opcode>REDG)\\.(?P<modifiers>E)\\.(?P<modifiers>ADD)\\.(?P<modifiers>STRONG)\\.(?P<modifiers>GPU) (?P<address>(?P<operands>desc\\[UR[0-9]+\\]\\[(?:R[0-9]+|UR[0-9]+)\\.64(?:\\+(?:-?0x[0-9A-Fa-f]+|UR[0-9]+|UR[0-9]+\\+-?0x[0-9A-Fa-f]+))?\\])), (?P<operands>R[0-9]+)', flags=regex.V0)
Arch: sm_90, Instruction: [(Instruction(offset=176, instruction='REDG.E.ADD.STRONG.GPU desc[UR4][R2.64], R5', hex='0x000000050200798e', control=ControlCode(stall_count=1, yield_flag=True, read=7, write=7, wait=[False, False, False, True, False, False], reuse={'A': False, 'B': False, 'C': False, 'D': False})), InstructionMatch(opcode='REDG', modifiers=('E', 'ADD', 'STRONG', 'GPU'), operands=('desc[UR4][R2.64]', 'R5'), predicate=None, additional={'address': ['desc[UR4][R2.64]']}))]
Arch: sm_120, Pattern: regex.Regex('(?P<opcode>REDG)\\.(?P<modifiers>E)\\.(?P<modifiers>ADD)\\.(?P<modifiers>STRONG)\\.(?P<modifiers>GPU) (?P<address>(?P<operands>desc\\[UR[0-9]+\\]\\[(?:R[0-9]+|UR[0-9]+)\\.64(?:\\+(?:-?0x[0-9A-Fa-f]+|UR[0-9]+|UR[0-9]+\\+-?0x[0-9A-Fa-f]+))?\\])), (?P<operands>R[0-9]+)', flags=regex.V0)
Arch: sm_120, Instruction: [(Instruction(offset=176, instruction='REDG.E.ADD.STRONG.GPU desc[UR4][R2.64], R5', hex='0x000000050200798e', control=ControlCode(stall_count=1, yield_flag=True, read=7, write=7, wait=[False, False, True, False, False, False], reuse={'A': False, 'B': False, 'C': False, 'D': False})), InstructionMatch(opcode='REDG', modifiers=('E', 'ADD', 'STRONG', 'GPU'), operands=('desc[UR4][R2.64]', 'R5'), predicate=None, additional={'address': ['desc[UR4][R2.64]']}))]
The operation, scope, and dtype arguments passed to ReductionMatcher correspond directly to the source-level operation: 32-bit integer addition at device scope.
The consistency argument matches the STRONG modifier in the emitted SASS instruction.
As already mentioned, it is compatible with the relaxed C++ memory ordering of legacy atomicAdd because in NVIDIA’s memory consistency model, operations with relaxed semantics are classified as strong operations.
Thus, the same architecture-aware matcher specification can be used for all targets: the architecture-specific opcode and operand syntax are encapsulated by the matcher.
Outlook
This notebook illustrated how ReProspect supports a fully programmatic analysis of compiled CUDA code.
Matching SASS instructions has practical value.
For example, staying in the context of atomic functions, a code base may provide several implementations of an atomic operation, such as a mapping to a hardware atomic instruction, a compare-and-swap loop, or a lock-based implementation, with substantially different performance.
Which implementation is selected may depend on intricate logic in the code base and the compiler tool chain, and may change as either evolves.
Verifying the selection by micro-benchmarking requires a physical device and suffers from runtime variability; the compiled CUDA code, by contrast, already contains the answer.
With SASS instruction matchers, the intended mapping can become a test assertion that can run in CI/CD pipelines.
With ReProspect’s target-architecture-aware SASS instruction matchers, such a test assertion can be written in terms of a single architecture-aware matcher specification.
This pattern is used in our Kokkos atomics with desul case study.