Skip to content
Earn a certificate
AssemblyBashCC#C++CUDADartElixirGoHaskellJavaJavaScriptJuliaKotlinLuaMATLABMojoPerlPHPPythonRRubyRustScalaSQLSwiftTritonTypeScriptWebAssemblyZig
Volume A · The TensorCore.si Library

Assembly

Talk to the processor directly.

Assembly language is a human-readable form of a processor's machine code. Each line is roughly one CPU instruction. You rarely write whole programs in it today, but reading it explains how computers really work, and AI engineers read GPU assembly (PTX and SASS) to understand why kernels are fast or slow.

Prologue

Why learn Assembly?

You do not need to write assembly every day, but being able to read it is a superpower for performance work.

  • Understand hardware: registers, memory, the stack and how function calls work.
  • Optimise: check what your compiler actually produced with tools like Compiler Explorer.
  • GPU insight: read PTX and SASS to see tensor core instructions in AI kernels.
Prologue

Set up and run your first program

This volume uses x86-64 Linux with NASM. On Apple silicon or Raspberry Pi you would use ARM64 assembly instead; the ideas are the same.

terminal
sudo apt install nasm
nasm -f elf64 hello.asm -o hello.o
ld hello.o -o hello
./hello
hello.asm
section .data
msg db "Hello, AI! Welcome to Assembly.", 10
len equ $ - msg

section .text
global _start
_start:
    mov rax, 1          ; syscall: write
    mov rdi, 1          ; file descriptor 1 = stdout
    mov rsi, msg        ; address of the text
    mov rdx, len        ; number of bytes
    syscall
    mov rax, 60         ; syscall: exit
    xor rdi, rdi        ; exit code 0
    syscall
TipComments start with ; in NASM. Try godbolt.org to see compiler output for any C code.
Chapter 1

Registers and mov

The CPU works on registers, tiny fast storage slots. x86-64 has 16 general-purpose 64-bit registers such as rax, rbx, rcx and rdx. mov copies data; arithmetic works in place.

registers.asm
    mov rax, 32         ; layers
    mov rbx, 4096       ; width
    imul rbx, rbx       ; rbx = width * width
    imul rax, rbx       ; rax = layers * width^2
    imul rax, 12        ; rax = 12 * layers * width^2 (approx. params)
    add rax, 1          ; rax += 1
    sub rax, 1          ; rax -= 1
Chapter 2

Memory and addressing

Square brackets read or write memory. Addresses can combine a base register, an index and a scale, which is perfect for walking arrays.

memory.asm
section .data
vram dq 80, 141, 180        ; three 64-bit numbers

section .text
    lea rsi, [vram]         ; rsi = address of the array
    mov rax, [rsi]          ; rax = 80
    mov rbx, [rsi + 8]      ; rbx = 141 (next 8-byte item)
    mov rcx, 2
    mov rdx, [rsi + rcx*8]  ; rdx = 180
    mov qword [rsi], 96     ; overwrite the first value
Chapter 3

Comparisons and jumps

There are no if statements or loops, only compare and jump. cmp sets flags, and conditional jumps like jl (jump if less) read them.

loop.asm
    ; sum = 1 + 2 + ... + 10
    xor rax, rax            ; sum = 0
    mov rcx, 1              ; i = 1
loop_top:
    add rax, rcx            ; sum += i
    inc rcx                 ; i++
    cmp rcx, 10
    jle loop_top            ; repeat while i <= 10
    ; rax now holds 55
Chapter 4

The stack and functions

call pushes the return address and jumps; ret pops it and jumps back. On Linux x86-64, arguments go in rdi, rsi, rdx, rcx, r8, r9 and the result returns in rax.

functions.asm
; long square(long x)  -> rdi in, rax out
square:
    mov rax, rdi
    imul rax, rdi
    ret

_start:
    mov rdi, 12
    call square             ; rax = 144
    push rax                ; save on the stack
    pop rbx                 ; restore into rbx
Chapter 5

Calling assembly from C

Real projects write only the hot function in assembly and call it from C. This is how fast math libraries are built.

add_f32.asm
; add_f32.asm   nasm -f elf64 add_f32.asm && gcc main.c add_f32.o
global add_f32
section .text
add_f32:                    ; float add_f32(float a, float b)
    addss xmm0, xmm1        ; floats arrive in xmm0, xmm1
    ret                     ; result returns in xmm0

; main.c:
; float add_f32(float, float);
; int main() { printf("%f\n", add_f32(1.5f, 2.25f)); }
Part II · AI with Assembly

Assembly and AI

AI performance comes down to instructions. On CPUs, SIMD instructions (AVX2, AVX-512, ARM NEON and AMX) process many numbers at once, and llama.cpp uses them for fast inference. On NVIDIA GPUs, CUDA compiles to PTX (a virtual assembly) and then SASS (the real machine code); tensor core work shows up as mma or wgmma instructions.

The essential libraries

Compiler ExplorerSee assembly for C, C++, Rust and CUDA online.
cuobjdump / nvdisasmDump SASS from compiled CUDA kernels.
Intel intrinsics guideEvery SIMD instruction explained.
perfFind which instructions are hot.

An AVX dot product

Eight float multiplies and adds per instruction with 256-bit AVX registers (requires a CPU with AVX and FMA).

dot8.asm
; float dot8(const float *a, const float *b)  ; exactly 8 floats
global dot8
section .text
dot8:
    vmovups ymm0, [rdi]         ; load 8 floats from a
    vmovups ymm1, [rsi]         ; load 8 floats from b
    vmulps  ymm0, ymm0, ymm1    ; 8 multiplies at once
    vextractf128 xmm1, ymm0, 1  ; upper 4 lanes
    vaddps  xmm0, xmm0, xmm1    ; add halves -> 4 values
    vhaddps xmm0, xmm0, xmm0    ; horizontal add -> 2
    vhaddps xmm0, xmm0, xmm0    ; -> 1
    vzeroupper
    ret
Part III · Build an AI agent

Assembly inside an AI agent

Nobody writes an LLM agent in assembly, but agents can call fast native tools. Here a Python agent uses the AVX dot8 function above (built as a shared library) as a tool for similarity scores.

Your goal→Model plans→Tool call→Your code runs the tool→Result back to model→Final answer
agent.py
# nasm -f elf64 dot8.asm && gcc -shared dot8.o -o libdot8.so
import ctypes, json
from openai import OpenAI

lib = ctypes.CDLL("./libdot8.so")
lib.dot8.restype = ctypes.c_float
F8 = ctypes.c_float * 8

def dot8(a, b):
    return float(lib.dot8(F8(*a), F8(*b)))

client = OpenAI(base_url="http://localhost:11434/v1", api_key="ollama")
tools = [{"type": "function", "function": {"name": "dot8", "description": "Dot product of two 8-number vectors",
    "parameters": {"type": "object", "properties": {"a": {"type": "array", "items": {"type": "number"}},
    "b": {"type": "array", "items": {"type": "number"}}}, "required": ["a", "b"]}}}]
messages = [{"role": "user", "content": "Dot product of [1,2,3,4,5,6,7,8] and [1,1,1,1,1,1,1,1]?"}]
for _ in range(5):
    msg = client.chat.completions.create(model="llama3.1", messages=messages, tools=tools).choices[0].message
    messages.append(msg)
    if not msg.tool_calls:
        print(msg.content); break
    for c in msg.tool_calls:
        messages.append({"role": "tool", "tool_call_id": c.id, "content": str(dot8(**json.loads(c.function.arguments)))})

The answer should be 36, computed by eight parallel multiplies in a single AVX instruction.

Run it locallyInstall Ollama, run ollama pull llama3.1, and keep it running. The code talks to its OpenAI-compatible API at http://localhost:11434/v1, so the same program also works with vLLM, llama.cpp server or a cloud API by changing the URL, model name and adding an API key.
Part IV · Your path

Assembly learning roadmap

Tick items off as you go. Your progress is saved in this browser.

Beginner

Intermediate

Advanced

Part IV · Your path

Projects to build

Hello syscallsPrint text and read input with Linux syscalls.
Array sumSum an array, then speed it up with AVX.
Kernel readerCompile a CUDA kernel and find its tensor core instructions in SASS.
Reference

Assembly cheat sheet

TaskSyntax
Copymov rax, rbx
Load memorymov rax, [rsi + 8]
Compare and jumpcmp rax, 10 / jl label
Call functioncall f / ret
Stackpush rax / pop rax
Exit (Linux)mov rax, 60 / syscall
Reference

Common mistakes

  • Assembly is specific to one CPU family: x86-64 code will not run on ARM.
  • Forgetting to save registers your caller expects to keep (rbx, rbp, r12–r15) causes strange bugs.
  • Intel and AT&T syntax put operands in opposite order.
Reference

Where to go next

On TensorCore.si: C volume, CUDA volume.

Keyboard: [ previous language · ] next language · / search the library