{ "cells": [ { "cell_type": "markdown", "id": "0", "metadata": {}, "source": [ "Getting started with binary analysis\n", "====================================\n", "\n", "This notebook shows how to get started with binary analysis.\n", "\n", "Binary analysis examines compiled CUDA code.\n", "The CUDA toolkit provides the [CUDA binary utilities](https://docs.nvidia.com/cuda/cuda-binary-utilities/index.html) `cuobjdump` and `nvdisasm` for extracting and disassembling the contents of CUDA binary files.\n", "`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.\n", "\n", "As an example, we consider a kernel involving an atomic function call.\n", "We examine the CUDA assembly (SASS) instructions that this call is compiled to." ] }, { "cell_type": "markdown", "id": "1", "metadata": {}, "source": [ "About this page\n", "---------------\n", "\n", "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.\n", "\n", "To run the example yourself, {download}`download the notebook ` and open it in JupyterLab.\n", "Alternatively, copy-paste the code snippets successively into an interactive Python session.\n", "\n", "The requirements are Python 3.10 or newer and `ReProspect`.\n", "\n", "A CUDA Toolkit 12.8 or newer installation providing `nvcc`, `cuobjdump`, and `cu++filt` is also needed.\n", "Because the example is concerned with analysing the CUDA assembly code produced by compilation but does not run the kernel, no GPU is required." ] }, { "cell_type": "markdown", "id": "2", "metadata": {}, "source": [ "Source code\n", "-----------\n", "\n", "Let us consider a kernel that atomically adds each element of a source array to the corresponding element of a destination array.\n", "\n", "In this source code, `atomicAdd` is a legacy atomic function provided by CUDA as a C++ language extension, as described [here](https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/cpp-language-extensions.html#atomicadd).\n", "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.\n", "As a legacy atomic function, it has `memory_order_relaxed` C++ atomic memory semantics.\n", "Because the function name `atomicAdd` does not have a suffix, it is atomic at device scope.\n", "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." ] }, { "cell_type": "code", "execution_count": null, "id": "3", "metadata": {}, "outputs": [], "source": [ "CODE = \"\"\"\\\n", "#include \"cuda.h\"\n", "\n", "__global__ void my_kernel(int* __restrict__ const dst, const int* __restrict__ const src) {\n", " const auto index = blockIdx.x * blockDim.x + threadIdx.x;\n", " atomicAdd(&dst[index], src[index]);\n", "}\n", "\"\"\"" ] }, { "cell_type": "markdown", "id": "4", "metadata": {}, "source": [ "Compilation\n", "-----------\n", "\n", "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`).\n", "\n", "Later in this notebook, the atomic function call will be found to compile to different SASS instructions across these targets." ] }, { "cell_type": "code", "execution_count": null, "id": "5", "metadata": {}, "outputs": [], "source": [ "import pathlib\n", "import subprocess\n", "\n", "from reprospect.tools import architecture\n", "\n", "ARCHES = [\n", " architecture.NVIDIAArch.from_compute_capability(80),\n", " architecture.NVIDIAArch.from_compute_capability(90),\n", " architecture.NVIDIAArch.from_compute_capability(120),\n", "]\n", "\n", "print(subprocess.check_output(('nvcc', '--version')).decode().strip())\n", "\n", "workdir = pathlib.Path.cwd() / 'example_binary_analysis'\n", "workdir.mkdir(exist_ok=True)\n", "print(f'Working directory: {workdir}')\n", "\n", "for arch in ARCHES:\n", " source = workdir / f'atomic.{arch.as_sm}.cu'\n", " output = workdir / f'atomic.{arch.as_sm}'\n", "\n", " source.write_text(CODE)\n", "\n", " subprocess.check_call(('nvcc', f'--generate-code=arch={arch.as_compute},code=[{arch.as_sm}]', '-O3', '-c', source, '-o', output))" ] }, { "cell_type": "markdown", "id": "6", "metadata": {}, "source": [ "Extracting and disassembling CUDA binary code\n", "---------------------------------------------\n", "\n", "By default, the CUDA compiler driver `nvcc` embeds CUDA binary files into the compiled output file, as described [here](https://docs.nvidia.com/cuda/cuda-binary-utilities/index.html).\n", "A CUDA binary file is also referred to as a cubin.\n", "\n", "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.\n", "We will now extract this embedded cubin from the compiled output file and disassemble its CUDA binary code for each target.\n", "The CUDA toolkit provides the CUDA binary utility `cuobjdump` for this purpose.\n", "Here, we invoke it through the `ReProspect` class {py:class}`reprospect.tools.binaries.cuobjdump.CuObjDump`.\n", "The method {py:meth}`~reprospect.tools.binaries.cuobjdump.CuObjDump.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.\n", "The returned object has a member `functions`, a dictionary that maps each kernel's signature to its SASS code as a string.\n", "The kernel signature is decoded (demangled) from its low-level assembly name.\n", "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.\n", "\n", "It should be noted that the cubins embedded in a compiled output file follow a naming convention as in `...cubin`.\n", "The `ReProspect` class {py:class}`~reprospect.tools.binaries.cuobjdump.CuObjDump` has a method {py:meth}`~reprospect.tools.binaries.cuobjdump.CuObjDump.embedded_cubins` to inspect the names of the cubins embedded in a compiled output file." ] }, { "cell_type": "code", "execution_count": null, "id": "7", "metadata": {}, "outputs": [], "source": [ "from reprospect.tools.binaries import CuObjDump\n", "\n", "cuobjdump = {}\n", "\n", "for arch in ARCHES:\n", " compiled_output = workdir / f'atomic.{arch.as_sm}'\n", "\n", " cuobjdump[arch], _ = CuObjDump.extract(\n", " file=compiled_output,\n", " arch=arch,\n", " cwd=workdir,\n", " cubin=f'atomic.1.{arch.as_sm}.cubin',\n", " )\n", "\n", " print(cuobjdump[arch].functions.keys())" ] }, { "cell_type": "markdown", "id": "8", "metadata": {}, "source": [ "Like many classes in `ReProspect`, the class {py:class}`reprospect.tools.binaries.cuobjdump.CuObjDump` supports rich rendering.\n", "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." ] }, { "cell_type": "code", "execution_count": null, "id": "9", "metadata": {}, "outputs": [], "source": [ "for arch in ARCHES:\n", " print(cuobjdump[arch])" ] }, { "cell_type": "markdown", "id": "10", "metadata": {}, "source": [ "Decoding SASS instructions\n", "--------------------------\n", "\n", "Next, we use `ReProspect` to parse the SASS instructions.\n", "`ReProspect` provides the class {py:class}`reprospect.tools.binaries.sass.decoder.Decoder` for this purpose.\n", "The snippet below constructs for each target an instance of this class from the SASS code that was stored as a string previously.\n", "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 {py:class}`~reprospect.tools.binaries.sass.decoder.ControlCode`).\n", "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." ] }, { "cell_type": "code", "execution_count": null, "id": "11", "metadata": { "tags": [ "hide-output" ] }, "outputs": [], "source": [ "from reprospect.tools.binaries.sass import Decoder\n", "\n", "SIGNATURE = 'my_kernel(int *, const int *)'\n", "\n", "decoder: dict[architecture.NVIDIAArch, Decoder] = {}\n", "\n", "for arch in ARCHES:\n", " decoder[arch] = Decoder(code=cuobjdump[arch].functions[SIGNATURE].code)" ] }, { "cell_type": "markdown", "id": "12", "metadata": {}, "source": [ "The class {py:class}`~reprospect.tools.binaries.sass.decoder.Decoder` supports rich rendering.\n", "Printing shows for each SASS instruction its offset, its disassembled representation as a string, and its decoded control code." ] }, { "cell_type": "code", "execution_count": null, "id": "13", "metadata": {}, "outputs": [], "source": [ "for arch in ARCHES:\n", " print(decoder[arch])" ] }, { "cell_type": "markdown", "id": "14", "metadata": {}, "source": [ "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).\n", "\n", "Because its return value is unused, the compiler can lower `atomicAdd` to a reduction operation instruction.\n", "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*.\n", "\n", "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. \n", "It should be noted that when proceeding in this way, the opcode must be hard-coded per target architecture.\n", "The matchers introduced next will abstract such details away." ] }, { "cell_type": "code", "execution_count": null, "id": "15", "metadata": {}, "outputs": [], "source": [ "instruction_red = {}\n", "\n", "for arch in ARCHES:\n", " opcode = 'RED' if arch.compute_capability < 90 else 'REDG'\n", " instruction_red[arch] = [\n", " instr for instr in decoder[arch].instructions\n", " if f'{opcode}.' in instr.instruction\n", " ]\n", " assert len(instruction_red[arch]) == 1" ] }, { "cell_type": "markdown", "id": "16", "metadata": {}, "source": [ "Decomposing SASS instructions\n", "-----------------------------\n", "\n", "`ReProspect` provides an extensible matching framework for SASS instructions.\n", "At the lowest levels, matchers analyse SASS instructions and their components (instruction predicate, opcode, modifiers, and operands)." ] }, { "cell_type": "code", "execution_count": null, "id": "17", "metadata": {}, "outputs": [], "source": [ "from reprospect.testing.binaries.sass.instruction import AnyMatcher\n", "\n", "for arch in ARCHES:\n", " decomposed = AnyMatcher().match(inst=instruction_red[arch][0].instruction)\n", " print(f'Arch: {arch.as_sm}, Instruction: {decomposed}')" ] }, { "cell_type": "markdown", "id": "18", "metadata": {}, "source": [ "We can observe that for each considered target, the modifiers are `E`, `ADD`, `STRONG` and `GPU`.\n", "These modifiers indicate extended addressing (`E`), the reduction operation (`ADD`), the `STRONG` modifier, and the scope (`GPU`, i.e., `Device`).\n", "\n", "NVIDIA documents SASS instructions only partially.\n", "The interpretation of `STRONG` should be stated cautiously.\n", "It is compatible with NVIDIA's [memory consistency model](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html?highlight=atomic#operation-types), in which relaxed memory operations are classified as strong operations.\n", "\n", "NVIDIA does fully document the PTX instructions, which compilers generate before ultimately translating them to native target-architecture-specific SASS instructions.\n", "Consulting the PTX instruction set reference can thus also provide insight into the ultimately emitted SASS instructions; see, for instance, [this description](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html?highlight=atomic#parallel-synchronization-and-communication-instructions-red) of the PTX `red` instruction.\n", "\n", "\n", "\n", "Further, it should be noted that although the modifiers coincide across the three targets here, modifiers may depend on the target architecture in general." ] }, { "cell_type": "markdown", "id": "19", "metadata": {}, "source": [ "`ReProspect`'s matchers are implemented internally in terms of regex matching." ] }, { "cell_type": "code", "execution_count": null, "id": "20", "metadata": {}, "outputs": [], "source": [ "print(AnyMatcher().PATTERN)" ] }, { "cell_type": "markdown", "id": "21", "metadata": {}, "source": [ "Matching SASS instructions across target architectures\n", "------------------------------------------------------\n", "\n", "`ReProspect` provides SASS instruction matchers that abstract away target-architecture-specific details.\n", "\n", "Here, we use {py:class}`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.\n", "It builds its regex pattern as a function of the target architecture and of the requested reduction operation, scope, consistency, and value type." ] }, { "cell_type": "code", "execution_count": null, "id": "22", "metadata": {}, "outputs": [], "source": [ "import numpy\n", "\n", "from reprospect.testing.binaries.sass.instruction import ReductionMatcher, ThreadScope\n", "\n", "for arch in ARCHES:\n", " matcher_red = ReductionMatcher(\n", " arch=arch, operation='ADD',\n", " scope=ThreadScope.DEVICE,\n", " consistency='STRONG', dtype=numpy.int32,\n", " )\n", " print(f'Arch: {arch.as_sm}, Pattern: {matcher_red.pattern}')\n", " matched_red = [\n", " (inst, matched)\n", " for inst in decoder[arch].instructions\n", " if (matched := matcher_red.match(inst))\n", " ]\n", " assert len(matched_red) == 1\n", " print(f'Arch: {arch.as_sm}, Instruction: {matched_red}')" ] }, { "cell_type": "markdown", "id": "23", "metadata": {}, "source": [ "The `operation`, `scope`, and `dtype` arguments passed to {py:class}`~reprospect.testing.binaries.sass.instruction.atomic.ReductionMatcher` correspond directly to the source-level operation: 32-bit integer addition at device scope.\n", "The `consistency` argument matches the `STRONG` modifier in the emitted SASS instruction.\n", "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." ] }, { "cell_type": "markdown", "id": "24", "metadata": {}, "source": [ "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." ] }, { "cell_type": "markdown", "id": "25", "metadata": {}, "source": [ "Outlook\n", "-------\n", "\n", "This notebook illustrated how `ReProspect` supports a fully programmatic analysis of compiled CUDA code.\n", "Matching SASS instructions has practical value.\n", "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.\n", "Which implementation is selected may depend on intricate logic in the code base and the compiler tool chain, and may change as either evolves.\n", "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.\n", "With SASS instruction matchers, the intended mapping can become a test assertion that can run in CI/CD pipelines.\n", "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.\n", "This pattern is used in our {ref}`Kokkos atomics with desul case study `." ] } ], "metadata": { "kernelspec": { "display_name": "Python 3 (ipykernel)", "language": "python", "name": "python3" }, "language_info": { "codemirror_mode": { "name": "ipython", "version": 3 }, "file_extension": ".py", "mimetype": "text/x-python", "name": "python", "nbconvert_exporter": "python", "pygments_lexer": "ipython3", "version": "3.13.11" } }, "nbformat": 4, "nbformat_minor": 5 }