36 docs tagged with "memory"

Allocators are objects that manage memory allocation for STL containers. They provide a standardized interface for customizing how containers acquire and release memory.

The classic tiled kernel loop looks like load → syncthreads() → compute → syncthreads() while stage i computes, stage i+1's load can already be in flight.

boost::intrusive_ptr is a reference-counting smart pointer that keeps the count inside the

boost: it owns a single heap object and deletes it when

boost: any number of sharedptr instances can

Boost.Pool is a fast memory allocator specialised for handing out many objects of the same fixed

A GPU's cache hierarchy looks superficially like a CPU's — an L1 per core-analog, a shared L2 behind it — but the access pattern it's optimized for is completely different. A CPU cache is tuned for one thread's temporal and spatial locality; a GPU's L1 and L2 exist to serve tens of thousands of threads issuing memory requests in 32-wide warps, and the granularity at which those requests are actually served is the fact that explains coalescing, wasted bandwidth, and most of what looks like "mysterious" memory performance on a GPU.

There is no server-side "memory" object that magically remembers a user. Memory is just the message list you re-send on every call. Each turn, you append the new human message, invoke the model, append its reply, and send the whole (possibly trimmed) list back next time.

A LangGraph agent that remembers a conversation across turns and processes, with a token budget so it doesn't grow unbounded — thread persistence and trimming applied together rather than explained again here.

Global memory's performance rules assume threads in a warp want different addresses and reward spreading them out into distinct sectors. Constant memory and the read-only data path invert that assumption: they're fast precisely when a warp's threads all want the same data, and the further a kernel's access pattern gets from that, the less either buys over an ordinary global read.

Overview

Shared Memory is scoped to a single block the cluster's combined on-chip shared memory, addressable across block boundaries, without routing through global memory at all.

UObjects are not managed with new/delete or reference counting — they're managed by a tracing

Thread Indexing establishes that the fastest-varying array index should track threadIdx.x, so that adjacent threads in a warp touch adjacent addresses. This page is the mechanism that makes that rule matter: how the hardware actually turns a warp's 32 addresses into memory transactions, and why the difference between "adjacent" and "scattered" can be an 8x difference in delivered bandwidth for the exact same amount of useful data.

Data arranged at addresses that are multiples of its size. Required for correctness on some architectures, critical for performance on all.

Why this matters

Every earlier page in this section has quietly leaned on one synchronization primitive or another — syncthreads(), a cluster.sync(), a pipeline's consumer_wait() — to make writes from one thread visible to reads from another. This page states the rule underneath all of them explicitly: CUDA's memory model is weakly ordered, and without an explicit fence or atomic, there is no guarantee about when, or even whether, one thread's writes become visible to another thread at all. Getting this wrong doesn't usually crash a kernel; it produces a result that's correct most of the time and silently wrong occasionally, which is far worse.

Overview

Every variable and every function in your firmware ends up in one of about six buckets, and which bucket it lands in is decided by two things: whether it is code or data, and whether its initial value is zero. That is nearly the whole rule. int counter; goes in .bss because its initial value is zero. int counter = 5; goes in .data because it is not. const int limit = 5; goes in .rodata because it never changes and can therefore stay in flash. Nobody chose those placements for your variable; the compiler applied that rule and emitted a section name.

A CUDA kernel does not have one undifferentiated pool of memory to work with — it has six, each with its own scope, lifetime, and performance profile, and picking the wrong one for a given piece of data is one of the most common ways a kernel ends up an order of magnitude slower than it should be. This page is the map: what each space is, who can see it, how long it lives, and the rough latency and bandwidth numbers that make the choice matter. The pages that follow work through each space in depth.

How C++ objects are arranged in memory: data members, padding, vtables, base class subobjects. Understanding layout is crucial for binary compatibility and optimization.

The SAXPY program in Your First Kernel paid three costs without the source code drawing attention to any of them ordinary malloc'd host memory is pageable, and pageable memory cannot be the source or destination of an asynchronous transfer. Pinned memory removes that restriction, and once transfers are asynchronous, they can be scheduled to overlap with compute instead of paying for it serially.

Constructs objects in pre-allocated memory without allocating. Separates construction from allocation for custom memory management.

Pointer arithmetic navigates contiguous memory with automatic scaling by type size. Essential for arrays but dangerous without bounds checking.

Overview

A pointer is a variable that stores a memory address, allowing indirect access to other variables.

The Register File and Occupancy covers how a kernel's register usage feeds directly into the occupancy calculation. This page is the other side of that same fact: what registers actually hold, what forces a value out of the register file and into memory instead, and how to tell when that's happened to your kernel.

Global memory is fast in aggregate but every access still pays a round trip through the memory system; when several threads in a block need the same piece of data, or a thread needs to hand a value to another thread in the same block, routing through global memory to do it wastes bandwidth on traffic that never needed to leave the chip. Shared memory exists for exactly that: a small, explicitly-managed, on-chip scratchpad that every thread in a block can read and write, fast enough to use as a staging area rather than just a cache.

Shared memory earns its speed by serving a whole warp in one cycle, but that promise depends on the warp's 32 addresses landing in 32 different pieces of hardware. When they don't, shared memory — normally close behind register speed — degrades to a fraction of it, and the kernel doesn't fail, it just quietly runs slower with no compiler warning to explain why.

Boost.SmartPtr is where modern C++ ownership semantics were invented. Long before std::shared_ptr

Storage duration defines when and where objects are created and destroyed. C++ has four storage durations: automatic, static, dynamic, and thread.

A microcontroller has one 4 GB address space and everything lives in it: flash, RAM, every peripheral register, the interrupt controller, the debug hardware. There is no MMU, so what you write in a pointer is the physical address the bus sees. That is the simplification that makes bare-metal firmware tractable — and it means the layout of that space is not a vendor's private business but part of the architecture you program against.

Before any of the CUDA syntax in the next section makes sense, one structural fact has to be settled: a discrete GPU is a separate computer. It has its own memory, its own processors, and no automatic view of what the CPU is doing — every byte the GPU touches had to arrive there deliberately, and every result has to leave the same way. This page names that structure once, independent of any specific API, so that the CUDA-specific mechanics in Your First Kernel land on a model you already have rather than a pile of new syntax to memorize.

A null-pointer write on a Cortex-M does not crash. Address 0x00000000 is real memory — it is the start of flash, or the boot alias — so *(uint32t )0 = 42 on a fresh chip silently does nothing at all, and the program carries on. A stack that overflows its intended region does not crash either; it grows down into .bss and corrupts variables that belong to something else, and the failure surfaces minutes later in code that is entirely innocent. Both are the same problem: on a bare Cortex-M, memory has no permissions, so a wrong access is indistinguishable from a right one.*

Every allocation covered so far draws a hard line between host and device memory: a pointer is valid on one side or the other, and moving data across the line is an explicit cudaMemcpy the programmer writes and pays for. Unified Memory erases that line for the source code — one pointer, usable from both host and device — while the underlying hardware and driver still have to physically move bytes between two separate memory systems whenever the data is touched from the "wrong" side. Understanding when that migration happens, and how to steer it, is the difference between Unified Memory being a convenience and Unified Memory being a performance trap.

Overview