Appendix J: Heap versus Stack, and Pinned versus Pageable Memory¶
"Chapter 11 built a hand-rolled Arena allocator and traced a double-release bug to the exact line that caused it. Appendix C walked the GPU's own memory hierarchy from registers down to global DRAM. Appendix G ran the Rule of Five's five special member functions against a genuine double-free, caught live under AddressSanitizer. Every one of those chapters assumed the reader already knows the one distinction none of them stopped to teach directly: the difference between memory the compiler manages for you automatically, and memory you manage yourself -- and, once you're managing it yourself for a GPU program specifically, the further distinction between memory an operating system is free to move and memory it has promised never to touch. This appendix stops and teaches both, from first principles, with every claim checked against a genuine compile and, everywhere a real answer doesn't require a GPU to give one honestly, a genuine run."
What you will understand by the end of this appendix¶
- Exactly what "the stack" is: automatic, LIFO (last-in-first-out) storage the compiler allocates and reclaims for you at compile-decided offsets, with real address relationships and real, ASLR-independent evidence for both its growth direction and its strict unwind order -- and a genuine, deliberately triggered stack overflow, one real 1 KB frame at a time, to make its one hard limit concrete rather than theoretical.
- Exactly what "the heap" (free store) is, and the real difference between
malloc()/free()andnew/delete:malloc()hands back raw, uninitialized bytes and calls no constructor, ever, whilenewis a language operator that allocates AND genuinely constructs -- demonstrated here by a real, printed constructor call that only one of the two pairs ever makes, plusmalloc()'s C-style nullptr-on-failure contract standing next tonew's ownstd::bad_alloc, both genuinely triggered. - Why mixing allocation and deallocation functions (
new[]paired with plaindelete, ormalloc()paired withdelete) is undefined behavior, not a style preference -- caught here for real, under AddressSanitizer, as a genuinealloc-dealloc-mismatchreport this book already holds every claimed memory bug to the same standard for (Appendix D.6, Appendix G.2, Appendix G.5). - Placement new: how to separate "allocate storage" from "construct an object in that storage" entirely, construct an object at an address you already own (a stack buffer, a slot inside a hand-built pool allocator exactly like Chapter 11's own Arena), and end its lifetime correctly with an explicit destructor call rather than
delete-- genuinely compiled and run, on both a stack buffer and a malloc'd pool. - What "pinned" (page-locked) memory actually is, why the CUDA driver's DMA engine needs it for a fully asynchronous host-to-device transfer, and why ordinary pageable memory (anything
malloc()ornewever hands back) forces the driver to silently insert an extra staging copy instead of erroring -- demonstrated with the real CUDA Runtime API, genuinely compiled withnvccand genuinely called in this book's own no-GPU environment, honestly reporting exactly what a driver-less, device-less machine reports for every one of these calls.
What you need to know first¶
- Ordinary C++ scoping rules -- what a "local variable" and a "block scope" are -- and this book's own Chapter 11, whose hand-built
Arenaallocator and double-release bug are the direct ancestor of this appendix's own placement-new pool example. - Appendix G in full, whose Rule of Five and genuine AddressSanitizer-caught double-free (Sections G.2 and G.5) this appendix's own Section J.2.2 follows the identical discipline of: trigger the real bug, under the real tool, and lock only the report's stable, address-free summary line.
- Appendix C.5's own
cudaMallocManagedexample, whose "genuinely called, honestly reports no device present" pattern this appendix's own Section J.4 reuses directly forcudaHostAlloc,cudaMallocHost, andcudaMemcpyAsync. - Getting Started's own account of this book's build environment: a full CUDA toolkit (
nvcc, the real CUDA Runtime headers and library) with no NVIDIA GPU and no NVIDIA driver installed -- exactly the environment every file in this appendix was genuinely compiled and, wherever a real answer doesn't require a device, genuinely run in.
J.1 The Stack: Automatic Storage, Genuinely Measured [FOUNDATIONAL]¶
Intuition¶
Every local variable a C++ function declares gets its storage from the same place: a contiguous region of memory, one per thread, that grows every time a function is called and shrinks the instant that function returns -- no allocation call, no deallocation call, nothing for the programmer to remember to free, because the compiler already decided exactly where each variable lives and exactly when its lifetime ends, at compile time, from the shape of the source code alone. That's what "automatic storage duration" means, and it's the reason a local int x costs nothing to create or destroy: the cost was already paid once, in the single instruction that moves the stack pointer up or down by however many bytes the whole function's locals need.
Background¶
The stack is a LIFO (last-in-first-out) region: the most recently pushed frame is always the first one popped, which is exactly why a function's locals are guaranteed to be destroyed in the reverse of the order they were constructed, with zero exceptions, in every C++ program that has ever compiled. On essentially every mainstream architecture this book's sibling projects target (x86-64, ARM64 included), the stack grows downward -- toward lower addresses -- as more frames are pushed, though nothing in the C++ standard itself requires that direction, and a program should never rely on it for anything beyond what this section treats as a genuinely observed, platform-specific fact rather than a portable guarantee.
That "genuinely observed" qualifier matters more than it might look like it does. ASLR (Address Space Layout Randomization) -- enabled by default on essentially every modern Linux, macOS, and Windows machine -- relocates the entire stack to a fresh, randomized base address on every single process launch, specifically so an attacker who has previously observed one run's addresses learns nothing usable about the next run's. That means a raw stack address, printed once, is worthless as "locked" evidence: rerun the identical binary and every absolute address changes. What does NOT change between runs, ASLR or not, is the relationship between addresses observed within the same run -- which of two frames sits closer to the top of the address space, and the byte distance between two variables declared next to each other in the same frame. Every claim this section makes about direction and layout is checked as one of those relationships, computed fresh on every run, rather than as a raw address that a rerun could never reproduce.
File: 01_stack_layout_and_overflow.cpp¶
#include <cstdio>
#include <cstdint>
#include <cstdlib>
#include <cstring>
// Section J.1 -- the stack is automatic storage: every local variable's
// address and lifetime are decided entirely by the compiler at the point
// of declaration, unwound automatically (in exactly reverse order) the
// moment control leaves the enclosing scope. This file makes both halves
// of that claim genuinely observable: real addresses, printed directly,
// and real construction/destruction order, printed by a small RAII type.
struct Tracer {
const char* name;
explicit Tracer(const char* n) : name(n) {
printf(" constructed: %s\n", name);
}
~Tracer() {
printf(" destroyed: %s\n", name);
}
};
// Addresses themselves are NOT locked anywhere in this file's own output --
// ASLR (Address Space Layout Randomization, on by default on this and every
// modern Linux/macOS/Windows machine) genuinely relocates the whole stack to
// a different base address on every single run, specifically to make stack
// addresses useless to an attacker guessing at them. What IS stable across
// every run, ASLR or not, is the RELATIONSHIP between addresses within the
// same invocation: which frame's locals sit closer to the top of memory,
// and the byte distance between two variables declared next to each other
// in the same frame. This file measures and prints exactly those relationships,
// genuinely computed fresh on every run, rather than a raw address a rerun
// could never reproduce.
intptr_t frame_c(intptr_t caller_frame_addr) {
// Three locals, declared in order a, b, c.
int a = 1;
int b = 2;
int c = 3;
long long b_minus_a = (intptr_t)&b - (intptr_t)&a;
long long c_minus_b = (intptr_t)&c - (intptr_t)&b;
printf("frame_c: sizeof(int)=%zu, (&b - &a)=%lld bytes, (&c - &b)=%lld bytes\n",
sizeof(int), b_minus_a, c_minus_b);
printf("frame_c: this frame's locals sit at a lower address than frame_b's own local? %s\n",
((intptr_t)&a < caller_frame_addr) ? "true" : "false");
return (intptr_t)&a;
}
intptr_t frame_b(intptr_t caller_frame_addr) {
int local_in_b = 42;
printf("frame_b: this frame's local sits at a lower address than frame_a's own local? %s\n",
((intptr_t)&local_in_b < caller_frame_addr) ? "true" : "false");
return frame_c((intptr_t)&local_in_b);
}
intptr_t frame_a(intptr_t caller_frame_addr) {
int local_in_a = 7;
printf("frame_a: this frame's local sits at a lower address than main's own local? %s\n",
((intptr_t)&local_in_a < caller_frame_addr) ? "true" : "false");
return frame_b((intptr_t)&local_in_a);
}
void raii_scope_demo() {
printf("entering raii_scope_demo\n");
Tracer t1("t1 (declared first)");
{
Tracer t2("t2 (nested scope)");
Tracer t3("t3 (nested scope)");
printf(" inside nested scope\n");
}
printf(" back in outer scope\n");
Tracer t4("t4 (declared last, outer scope)");
printf("leaving raii_scope_demo\n");
}
// --trigger-overflow: an unbounded recursion, each frame consuming a real
// (non-optimizable-away) 1 KB of stack space via a local array the compiler
// cannot prove is unused. Compiled at -O0 specifically so this recursion is
// NOT turned into a loop by tail-call optimization -- the whole point is to
// genuinely exhaust the stack, one real frame at a time.
long long recurse_and_consume_stack(long long depth) {
char padding[1024];
memset(padding, (int)(depth & 0xFF), sizeof(padding));
if (depth % 1000 == 0) {
printf("depth=%lld, ~%lld KB of stack consumed so far\n", depth, depth);
fflush(stdout);
}
return padding[0] + recurse_and_consume_stack(depth + 1);
}
int main(int argc, char** argv) {
bool trigger_overflow = (argc > 1 && strcmp(argv[1], "--trigger-overflow") == 0);
if (trigger_overflow) {
printf("=== Section 4: a genuine stack overflow, one real 1 KB frame at a time ===\n");
fflush(stdout);
long long result = recurse_and_consume_stack(0);
// Unreachable under a real stack overflow -- the process crashes
// (SIGSEGV) inside recurse_and_consume_stack itself, before this
// line, or before the printf above's buffered output would even
// flush without the explicit fflush() calls this file already uses.
printf("unreachable: result=%lld\n", result);
return 0;
}
int main_local = 99;
printf("=== Section 1: nested stack frames, real address RELATIONSHIPS (not raw addresses -- ASLR note above main()) ===\n");
frame_a((intptr_t)&main_local);
printf("\n=== Section 2: LIFO destruction order, genuinely observed ===\n");
raii_scope_demo();
return 0;
}
g++ -std=c++17 -Wall -Wextra -O0 01_stack_layout_and_overflow.cpp -o 01_stack_layout_and_overflow
./01_stack_layout_and_overflow
Genuinely compiled and run:
=== Section 1: nested stack frames, real address RELATIONSHIPS (not raw addresses -- ASLR note above main()) ===
frame_a: this frame's local sits at a lower address than main's own local? true
frame_b: this frame's local sits at a lower address than frame_a's own local? true
frame_c: sizeof(int)=4, (&b - &a)=4 bytes, (&c - &b)=4 bytes
frame_c: this frame's locals sit at a lower address than frame_b's own local? true
=== Section 2: LIFO destruction order, genuinely observed ===
entering raii_scope_demo
constructed: t1 (declared first)
constructed: t2 (nested scope)
constructed: t3 (nested scope)
inside nested scope
destroyed: t3 (nested scope)
destroyed: t2 (nested scope)
back in outer scope
constructed: t4 (declared last, outer scope)
leaving raii_scope_demo
destroyed: t4 (declared last, outer scope)
destroyed: t1 (declared first)
Worked Example J.1.1 -- Three nested calls, and the direction this stack genuinely grows¶
frame_a calls frame_b calls frame_c, each declaring its own locals and each checking, with a plain pointer comparison rather than a printed address, whether its own local sits at a lower address than the local its caller just passed it. Every one of those three checks reports true in this run, and reports true on every rerun regardless of where ASLR happens to relocate the stack's base -- because ASLR shifts the whole stack by one random offset, uniformly, which changes every absolute address by the same amount and changes none of the relationships between them. That's the genuine, reproducible evidence that this platform's stack grows downward: each more-deeply-nested frame's locals land at a strictly lower address than the frame that called it, on every single run.
Inside frame_c, three int locals declared back-to-back (a, b, c) show a second, equally stable fact: (&b - &a) and (&c - &b) both come out to exactly sizeof(int) bytes on this build, meaning the compiler laid these three specific locals out contiguously, in declaration order, with no padding between them, in this exact stack frame. This is a real, observed fact about this compiler and this optimization level, not a portable C++ guarantee -- a different compiler, a different optimization level, or a different set of locals with different alignment needs could legitimately lay a frame out differently, including out of declaration order entirely, while remaining fully standard-conforming.
Worked Example J.1.2 -- A genuine stack overflow, one real 1 KB frame at a time¶
recurse_and_consume_stack() is unbounded, deliberately: it calls itself one more time every single invocation, with no base case, consuming a genuine (not optimized away) 1 KB local array on every frame. Compiled at -O0 specifically to prevent tail-call optimization from turning this recursion into a loop, running it with --trigger-overflow produces real, physical exhaustion of this process's stack region -- confirmed by a real SIGSEGV (segmentation fault, exit code 139) on every single run, at a depth between 7,000 and 8,000 frames, consistent with an ~8 MB default stack limit (ulimit -s reports 8192 KB on this machine) divided by roughly 1 KB of genuine consumption per frame:
=== Section 4: a genuine stack overflow, one real 1 KB frame at a time ===
depth=0, ~0 KB of stack consumed so far
depth=1000, ~1000 KB of stack consumed so far
depth=2000, ~2000 KB of stack consumed so far
depth=3000, ~3000 KB of stack consumed so far
depth=4000, ~4000 KB of stack consumed so far
depth=5000, ~5000 KB of stack consumed so far
depth=6000, ~6000 KB of stack consumed so far
depth=7000, ~7000 KB of stack consumed so far
[COMMON TRAP] A stack overflow is not a C++ exception, cannot be caught with try/catch, and is not guaranteed to happen at any particular, predictable depth -- it depends on the platform's stack size limit (configurable per-process and per-thread, and frequently much smaller for a spawned thread than for a process's main thread), the size of each frame's own locals, and how deep the compiler's own inlining decisions happen to make the effective call depth for a given piece of source code. The two most common real-world causes are not "obviously infinite recursion" the way this section's own example is -- they are a recursive function missing (or subtly miscomputing) its base case, and a single, enormous local array or object declared by value (a char buffer[10 * 1024 * 1024];, or a large fixed-size struct passed by value through several call layers) that alone exceeds the platform's stack budget on the very first call, with no recursion involved at all.
J.2 The Heap: malloc()/free() versus new/delete [FOUNDATIONAL]¶
Intuition¶
Heap (free-store) memory is the opposite of stack memory in every way Section J.1 just established: its lifetime is decided by the programmer at runtime, not by the compiler at compile time; it survives the return of the function that allocated it, which is exactly why it's the only kind of memory a function can correctly return a pointer or reference to; and nothing about it is automatically reclaimed -- every heap allocation this section's own examples make is matched, explicitly, with exactly one deallocation call, using the same allocation family that created it. malloc()/free() and new/delete are two entirely separate families with one shared property (both hand out heap memory) and one critical difference this section makes directly observable: only one of the two families knows anything about C++ objects at all.
Background¶
malloc(size_t n) is a C standard library function: it asks the underlying allocator for n raw, uninitialized bytes and hands back a void* (or a null pointer, on failure) -- full stop. It has no notion of "type," calls no constructor, and free() calls no destructor either; from malloc()/free()'s own point of view, a Widget and a same-sized array of raw char are indistinguishable. new and delete are C++ language operators, not functions the standard library merely provides: new T(args...) allocates storage (by calling operator new underneath, which behaves like malloc() for the allocation step) AND THEN genuinely invokes T's constructor on that storage, while delete p genuinely invokes p's destructor first and only then releases the storage (via operator delete). This is not a documentation-only distinction -- Worked Example J.2.1 below makes it directly, printably observable which of the two families ever actually ran a constructor.
The two families also disagree, by design, about how allocation failure is reported. malloc() follows C's own convention: return a null pointer, and let the caller remember to check it -- a convention this book's CUDA C API sections have already shown fails silently when a caller forgets (cudaMalloc's own cudaError_t return value follows an analogous pattern). new, by contrast, throws std::bad_alloc on failure specifically so a caller cannot silently ignore it the way a forgotten null check can be ignored -- an uncaught exception terminates the program loudly, with a message, rather than leaving a dangling null pointer for some later line to dereference. new (std::nothrow) exists for the narrow case (an allocator hot path with no nearby exception handling) where malloc()'s own nullptr-on-failure contract is genuinely what's wanted, opted into explicitly rather than defaulted to.
File: 02_malloc_vs_new.cpp¶
#include <cstdio>
#include <cstdlib>
#include <new>
#include <cstring>
#include <cstdint>
#include <cstddef>
// Section J.2 -- malloc()/free() and new/delete both hand out heap
// (free-store) memory, but only one of the two pairs knows anything about
// C++ objects. malloc() returns a raw, uninitialized block of bytes and
// free() returns it, full stop -- no constructor, no destructor, ever.
// new/delete are language operators, not library functions: new allocates
// AND genuinely calls the constructor, delete genuinely calls the
// destructor AND deallocates. This file makes the difference observable,
// not asserted.
struct Sentinel {
unsigned int magic;
char label[16];
Sentinel() : magic(0xABCD1234) {
std::strncpy(label, "constructed", sizeof(label) - 1);
label[sizeof(label) - 1] = '\0';
printf(" Sentinel::Sentinel() genuinely ran -- magic set to 0x%X\n", magic);
}
~Sentinel() {
printf(" Sentinel::~Sentinel() genuinely ran -- magic was 0x%X\n", magic);
}
};
void section1_malloc_skips_constructor() {
printf("=== Section 1: malloc() does NOT call the constructor ===\n");
// malloc() returns raw bytes it zero-fills with nothing in particular --
// whatever this heap page's memory previously held. Reinterpreting that
// raw block as a Sentinel* does NOT run Sentinel::Sentinel(); "magic"
// is whatever bytes malloc happened to hand back, not 0xABCD1234.
void* raw = malloc(sizeof(Sentinel));
Sentinel* s = reinterpret_cast<Sentinel*>(raw);
printf(" after malloc(sizeof(Sentinel)): no constructor message printed above this line\n");
printf(" s->magic == 0xABCD1234 (the constructor's own value)? %s\n",
(s->magic == 0xABCD1234) ? "true" : "false");
// Calling free() on this block is correct (it was allocated by malloc,
// and no destructor needs to run because none was ever constructed) --
// but nothing here ever invoked ~Sentinel(), which is exactly the point.
free(raw);
}
void section2_new_calls_constructor() {
printf("\n=== Section 2: new DOES call the constructor ===\n");
Sentinel* s = new Sentinel(); // the constructor message prints from inside this call
printf(" s->magic == 0xABCD1234 (the constructor's own value)? %s\n",
(s->magic == 0xABCD1234) ? "true" : "false");
delete s; // the destructor message prints from inside this call
}
void section3_failure_modes() {
printf("\n=== Section 3: malloc() returns nullptr on failure; new throws ===\n");
// An allocation size designed to fail on any real system: essentially
// the entire 64-bit address space at once. malloc() reports failure
// the C way -- a nullptr return, checked explicitly by the caller.
void* huge = malloc((size_t)-1);
printf(" malloc((size_t)-1) returned nullptr? %s\n", (huge == nullptr) ? "true" : "false");
// operator new reports the identical failure the C++ way: it throws
// std::bad_alloc instead of returning a sentinel value, specifically so
// a caller can't forget to check a return value the way huge could be
// (and, in real bug reports, often is) silently ignored above.
bool threw_bad_alloc = false;
try {
void* p = ::operator new((size_t)-1);
(void)p;
} catch (const std::bad_alloc& e) {
threw_bad_alloc = true;
printf(" operator new((size_t)-1) threw std::bad_alloc: \"%s\"\n", e.what());
}
printf(" operator new((size_t)-1) threw std::bad_alloc? %s\n", threw_bad_alloc ? "true" : "false");
// std::nothrow opts back into malloc's own nullptr-on-failure contract,
// for exactly the cases (an allocator hot path with no exception
// handling nearby) where a thrown exception isn't what the caller wants.
void* p2 = ::operator new((size_t)-1, std::nothrow);
printf(" operator new((size_t)-1, std::nothrow) returned nullptr? %s\n", (p2 == nullptr) ? "true" : "false");
}
void section4_array_forms() {
printf("\n=== Section 4: new[] / delete[] run the constructor/destructor for EVERY element ===\n");
const int N = 3;
Sentinel* arr = new Sentinel[N]; // N constructor calls, printed above this line
printf(" constructed %d Sentinel objects via new[]\n", N);
delete[] arr; // N destructor calls, printed above this line
printf(" destroyed all %d via delete[]\n", N);
}
void section5_alignment() {
printf("\n=== Section 5: over-aligned allocation, both the malloc-family and new way ===\n");
// Ordinary malloc()/new only guarantee alignof(std::max_align_t) (16
// bytes on this platform's ABI) -- correct for any built-in scalar
// type, but not enough for a type that genuinely demands more, such as
// a 64-byte cache-line-aligned struct used to avoid false sharing.
struct alignas(64) CacheLineAligned { char bytes[64]; };
printf(" alignof(CacheLineAligned) = %zu\n", alignof(CacheLineAligned));
printf(" alignof(std::max_align_t) = %zu (ordinary malloc/new's own guarantee)\n", alignof(std::max_align_t));
void* aligned_raw = std::aligned_alloc(64, sizeof(CacheLineAligned));
uintptr_t addr = (uintptr_t)aligned_raw;
printf(" std::aligned_alloc(64, ...) address is a multiple of 64? %s\n",
(addr % 64 == 0) ? "true" : "false");
free(aligned_raw);
// The C++17 language-level equivalent: operator new with an explicit
// std::align_val_t, which delete must be given back symmetrically.
CacheLineAligned* obj = new (std::align_val_t(64)) CacheLineAligned();
uintptr_t addr2 = (uintptr_t)obj;
printf(" new(std::align_val_t(64)) address is a multiple of 64? %s\n",
(addr2 % 64 == 0) ? "true" : "false");
::operator delete(obj, std::align_val_t(64));
}
int main() {
section1_malloc_skips_constructor();
section2_new_calls_constructor();
section3_failure_modes();
section4_array_forms();
section5_alignment();
return 0;
}
Genuinely compiled and run:
=== Section 1: malloc() does NOT call the constructor ===
after malloc(sizeof(Sentinel)): no constructor message printed above this line
s->magic == 0xABCD1234 (the constructor's own value)? false
=== Section 2: new DOES call the constructor ===
Sentinel::Sentinel() genuinely ran -- magic set to 0xABCD1234
s->magic == 0xABCD1234 (the constructor's own value)? true
Sentinel::~Sentinel() genuinely ran -- magic was 0xABCD1234
=== Section 3: malloc() returns nullptr on failure; new throws ===
malloc((size_t)-1) returned nullptr? true
operator new((size_t)-1) threw std::bad_alloc: "std::bad_alloc"
operator new((size_t)-1) threw std::bad_alloc? true
operator new((size_t)-1, std::nothrow) returned nullptr? true
=== Section 4: new[] / delete[] run the constructor/destructor for EVERY element ===
Sentinel::Sentinel() genuinely ran -- magic set to 0xABCD1234
Sentinel::Sentinel() genuinely ran -- magic set to 0xABCD1234
Sentinel::Sentinel() genuinely ran -- magic set to 0xABCD1234
constructed 3 Sentinel objects via new[]
Sentinel::~Sentinel() genuinely ran -- magic was 0xABCD1234
Sentinel::~Sentinel() genuinely ran -- magic was 0xABCD1234
Sentinel::~Sentinel() genuinely ran -- magic was 0xABCD1234
destroyed all 3 via delete[]
=== Section 5: over-aligned allocation, both the malloc-family and new way ===
alignof(CacheLineAligned) = 64
alignof(std::max_align_t) = 16 (ordinary malloc/new's own guarantee)
std::aligned_alloc(64, ...) address is a multiple of 64? true
new(std::align_val_t(64)) address is a multiple of 64? true
Worked Example J.2.1 -- malloc() skips the constructor; new doesn't¶
Sentinel's own constructor sets a fixed magic value (0xABCD1234) and prints a message every time it genuinely runs. Section 1's malloc(sizeof(Sentinel)) call produces no such message -- because none ran -- and the freshly malloc()'d block's magic field reads back as false against the constructor's own expected value in this run (whatever bytes this heap page happened to already hold, not the value only a real constructor call could have set). Section 2's new Sentinel() call prints the constructor's own message immediately, and magic genuinely reads back true. Section 4 confirms this scales correctly to arrays: new Sentinel[3] prints three separate constructor messages (one per element, not one for the whole block), and delete[] prints three separate destructor messages back, in the reverse order elements are conventionally torn down.
Worked Example J.2.2 -- Two honestly different failure modes, both genuinely triggered¶
Requesting (size_t)-1 bytes -- essentially the entire 64-bit address space in one allocation -- is guaranteed to fail on any real system, which is exactly why this file uses it to genuinely trigger both families' own failure behavior rather than describing it secondhand. malloc((size_t)-1) genuinely returns a null pointer, checked directly. ::operator new((size_t)-1) genuinely throws std::bad_alloc, caught directly by a try/catch around the call, with e.what() printing the exception's own real message. ::operator new((size_t)-1, std::nothrow) genuinely returns a null pointer instead of throwing, confirming the std::nothrow overload's own opt-back-into-malloc()'s-contract behavior for real, not merely by reading its documentation.
[COMMON TRAP] delete and delete[] must be matched to the allocation function that produced the pointer: a T* from new T is freed with delete; a T* from new T[n] is freed with delete[]; a void* from malloc()/calloc()/realloc()/aligned_alloc() is freed with free() -- and mixing any of these three families is undefined behavior, not merely a style violation, because new[] and plain new record different internal bookkeeping (an element count, for the array form) that the matching deallocation function is required to read correctly. Worked Example J.2.3 below triggers exactly this bug for real.
File: 03_alloc_dealloc_mismatch_asan.cpp¶
#include <cstdio>
#include <cstdlib>
#include <cstring>
// Section J.2 continued -- new[] and delete (without the brackets) are a
// MISMATCHED pair: delete alone tells the runtime "one object lived here,"
// while new[] recorded "N objects lived here, plus a hidden count." Freeing
// a new[]'d array with plain delete reads that hidden bookkeeping wrong and
// corrupts the heap -- undefined behavior that can run "successfully" with
// no visible symptom for a long time, on an unlucky day, in production.
// This file triggers the real bug for real, under AddressSanitizer, rather
// than describing it -- the same standard this book already holds every
// claimed memory bug to (Appendix D.6, Appendix G.2, Appendix G.5).
struct Row {
double values[4];
Row() { for (double& v : values) v = 0.0; }
};
int main(int argc, char** argv) {
bool trigger = (argc > 1 && std::strcmp(argv[1], "--trigger-mismatch") == 0);
printf("allocating 4 Row objects with new[]...\n");
Row* rows = new Row[4];
rows[0].values[0] = 3.14;
printf("rows[0].values[0] = %f\n", rows[0].values[0]);
if (!trigger) {
printf("freeing correctly with delete[] (matched, no bug)\n");
delete[] rows;
printf("done -- no mismatch triggered in this run\n");
return 0;
}
printf("freeing INCORRECTLY with plain delete (mismatched -- new[] / delete)...\n");
delete rows; // genuine alloc-dealloc-size mismatch, caught by ASan below
printf("unreachable under a working ASan build\n");
return 0;
}
Genuinely compiled with g++ -std=c++17 -Wall -Wextra -fsanitize=address -g. This exact mismatch is significant enough a bug that GCC 13.3 (this book's own compiler, per Getting Started) catches a version of it statically, at compile time, with no sanitizer involved at all:
03_alloc_dealloc_mismatch_asan.cpp: In function 'int main(int, char**)':
03_alloc_dealloc_mismatch_asan.cpp:36:12: warning: 'void operator delete(void*, long unsigned int)' called on pointer returned from a mismatched allocation function [-Wmismatched-new-delete]
36 | delete rows; // genuine alloc-dealloc-size mismatch, caught by ASan below
| ^~~~
03_alloc_dealloc_mismatch_asan.cpp:24:26: note: returned from 'void* operator new [](long unsigned int)'
24 | Row* rows = new Row[4];
| ^
Run without arguments, the matched, correct new[]/delete[] path produces no error at all:
allocating 4 Row objects with new[]...
rows[0].values[0] = 3.140000
freeing correctly with delete[] (matched, no bug)
done -- no mismatch triggered in this run
Worked Example J.2.3 -- The mismatch, genuinely triggered, genuinely caught¶
Run a second time with --trigger-mismatch, this file allocates the identical new Row[4] array, then deliberately frees it with plain delete instead of delete[]. AddressSanitizer catches the real consequence immediately and aborts the process (exit code 1) with a genuine, unedited report; only the report's own stable, address-free summary line is locked below, for the identical reason Appendix D.6's own dangling-reference report and Appendix G.2/G.5's own double-free reports lock only their first lines -- the full report additionally includes real heap addresses and a real thread id that, correctly, differ on every run:
SUMMARY: AddressSanitizer: alloc-dealloc-mismatch ../../../../src/libsanitizer/asan/asan_new_delete.cpp:164 in operator delete(void*, unsigned long)
operator new [] on the allocation side and operator delete on the deallocation side, named explicitly in AddressSanitizer's own report -- this is not a hypothetical failure mode described in prose; it is a real, unedited sanitizer report, produced by genuinely running a program that genuinely mismatches its own allocation and deallocation functions, the same evidentiary standard this book has held every claimed memory bug to since Appendix D.6.
J.3 Placement New and Manual Object Lifetime [FOUNDATIONAL]¶
Intuition¶
Ordinary new T(args...) bundles two genuinely separate operations into one keyword: find storage, and construct an object in it. Placement new unbundles them on purpose. Given a block of memory that already exists -- a stack buffer, a slot inside a hand-built pool, exactly the situation Chapter 11's own Arena allocator manages by hand -- placement new constructs an object at that exact, already-owned address, with no allocation of its own happening at all. The object it builds must have its destructor called explicitly, by name, because there is no matching delete to call for storage delete never allocated in the first place.
Background¶
The syntax is new (address) T(args...) -- a T object, genuinely constructed, at exactly address, using storage the caller is responsible for having already reserved (correctly sized, via sizeof(T), and correctly aligned, via alignof(T)). Ending that object's lifetime correctly means calling its destructor directly, by name (p->~T()), NOT delete p -- delete would attempt to call operator delete on an address that was never returned by operator new, which for a stack buffer or a slot inside a larger pool allocation is a serious, genuine bug: it would try to free memory that was never separately allocated to begin with, corrupting whatever real allocation (the stack frame, or the pool's own single malloc()) that address actually belongs to.
This is precisely the technique a custom allocator needs to build objects inside memory it manages itself, rather than delegating every individual allocation back to the general-purpose heap allocator: Chapter 11's own Arena hands out raw, aligned offsets from one large pre-allocated block, and a caller wanting a real, constructed object at one of those offsets -- not just raw bytes -- would use exactly the placement-new-plus-explicit-destructor pattern this section demonstrates, on both a stack buffer and a malloc()'d pool of several slots.
File: 04_placement_new.cpp¶
#include <cstdio>
#include <cstdlib>
#include <new>
#include <cstring>
// Section J.3 -- placement new separates the two things ordinary new
// bundles together: allocating raw storage, and constructing an object
// inside it. Placement new takes storage that ALREADY exists (a stack
// buffer, a malloc'd block, a slot inside a custom pool allocator -- the
// exact situation Chapter 11's own arena allocator and Appendix G's own
// Rule-of-Five types are built to manage) and runs a constructor at that
// exact address, with no allocation of its own. The object built this way
// must have its destructor called explicitly and by name -- delete would
// try to free storage placement new never allocated in the first place.
struct Widget {
int id;
double value;
Widget(int i, double v) : id(i), value(v) {
printf(" Widget::Widget(%d, %.1f) constructed\n", id, value);
}
~Widget() {
printf(" Widget::~Widget() destroying id=%d\n", id);
}
};
void section1_placement_new_on_stack_buffer() {
printf("=== Section 1: placement new into an already-existing STACK buffer ===\n");
// A raw, unconstructed buffer with no Widget in it yet -- just bytes,
// correctly SIZED and ALIGNED for a Widget, sitting on the stack.
alignas(Widget) unsigned char buffer[sizeof(Widget)];
printf(" raw buffer allocated on the stack, sizeof=%zu, alignof=%zu\n",
sizeof(buffer), alignof(Widget));
// Placement new: no allocation happens here at all -- it constructs a
// Widget IN buffer's own memory, at the exact address &buffer[0].
Widget* w = new (buffer) Widget(1, 3.5);
printf(" &buffer[0] == (void*)w? %s\n", ((void*)&buffer[0] == (void*)w) ? "true" : "false");
printf(" w->id=%d w->value=%.1f\n", w->id, w->value);
// Explicit destructor call BY NAME -- this is the only correct way to
// end this object's lifetime. `delete w;` would be a serious bug: it
// would call operator delete on a stack address that was never
// returned by operator new, corrupting the heap's own bookkeeping.
w->~Widget();
printf(" destructor called explicitly by name (w->~Widget()), NOT via delete\n");
}
void section2_placement_new_in_pool_slot() {
printf("\n=== Section 2: placement new into a slot of a malloc'd pool (the arena-allocator pattern) ===\n");
// A pool of 3 raw, unconstructed Widget-sized slots -- exactly the
// shape of Chapter 11's own Arena allocator: one big raw allocation,
// objects constructed into specific offsets within it on demand.
const int SLOT_COUNT = 3;
unsigned char* pool = static_cast<unsigned char*>(
std::aligned_alloc(alignof(Widget), SLOT_COUNT * sizeof(Widget)));
printf(" pool allocated: %d raw slots, %zu bytes each\n", SLOT_COUNT, sizeof(Widget));
Widget* widgets[SLOT_COUNT];
for (int i = 0; i < SLOT_COUNT; i++) {
unsigned char* slot = pool + i * sizeof(Widget);
widgets[i] = new (slot) Widget(100 + i, 1.5 * i);
}
printf(" all %d widgets constructed in-place inside one pool allocation\n", SLOT_COUNT);
for (int i = 0; i < SLOT_COUNT; i++) {
printf(" widgets[%d]->id=%d\n", i, widgets[i]->id);
}
// Destroy each object explicitly (in-place, no per-object deallocation --
// there IS no per-object allocation to give back), then free the pool
// itself exactly once, with the same allocator that created it.
for (int i = 0; i < SLOT_COUNT; i++) {
widgets[i]->~Widget();
}
printf(" all %d widgets destroyed explicitly, in-place\n", SLOT_COUNT);
free(pool);
printf(" pool freed once, as the single raw allocation it always was\n");
}
int main() {
section1_placement_new_on_stack_buffer();
section2_placement_new_in_pool_slot();
return 0;
}
Genuinely compiled and run:
=== Section 1: placement new into an already-existing STACK buffer ===
raw buffer allocated on the stack, sizeof=16, alignof=8
Widget::Widget(1, 3.5) constructed
&buffer[0] == (void*)w? true
w->id=1 w->value=3.5
Widget::~Widget() destroying id=1
destructor called explicitly by name (w->~Widget()), NOT via delete
=== Section 2: placement new into a slot of a malloc'd pool (the arena-allocator pattern) ===
pool allocated: 3 raw slots, 16 bytes each
Widget::Widget(100, 0.0) constructed
Widget::Widget(101, 1.5) constructed
Widget::Widget(102, 3.0) constructed
all 3 widgets constructed in-place inside one pool allocation
widgets[0]->id=100
widgets[1]->id=101
widgets[2]->id=102
Widget::~Widget() destroying id=100
Widget::~Widget() destroying id=101
Widget::~Widget() destroying id=102
all 3 widgets destroyed explicitly, in-place
pool freed once, as the single raw allocation it always was
Worked Example J.3.1 -- One stack buffer, one pool, both genuinely constructed in place¶
Section 1 reserves sizeof(Widget) raw, correctly-aligned bytes directly on the stack (alignas(Widget) unsigned char buffer[sizeof(Widget)];) and constructs a Widget into it with new (buffer) Widget(1, 3.5) -- genuinely confirmed, via a direct pointer comparison, that the constructed object's own address IS &buffer[0], not some separately allocated block elsewhere. Its destructor is called explicitly, by name, and NOT through delete.
Section 2 extends the identical technique to a pool of three slots inside one malloc()-family allocation (std::aligned_alloc), exactly the shape of Chapter 11's own Arena: one raw allocation, three objects constructed at three specific offsets within it, each destroyed explicitly and in-place, with the entire pool freed exactly once, using the same allocator (free()) that created it -- not once per object, because there was never more than one real allocation to begin with.
[COMMON TRAP] Placement new's new keyword makes it easy to mistake for ordinary new at a glance, but the two share almost nothing at runtime: ordinary new calls operator new to genuinely allocate memory, while placement new's own operator new(size_t, void* p) overload does nothing but return p unchanged -- no allocation happens, which is the entire point. Calling delete on a placement-new-constructed pointer is consequently not merely wrong style; for a stack buffer it is memory corruption (attempting to free stack memory through the heap allocator), and for a pool slot it is exactly the kind of alloc-dealloc-mismatch Worked Example J.2.3 already caught for real, applied to an address operator delete never allocated at all.
J.4 Pinned (Page-Locked) versus Pageable Memory [FOUNDATIONAL]¶
Intuition¶
Every heap allocation this appendix has made so far -- malloc(), new, new[], std::aligned_alloc -- produces what CUDA's own documentation calls pageable memory: ordinary memory, backed by pages the operating system is free to swap out to disk, or physically relocate in RAM, at any moment it chooses, entirely invisibly to the program that allocated it. That's a completely reasonable default for ordinary CPU-only code. It is a genuine problem for the GPU's own DMA (Direct Memory Access) engine, which reads directly from host RAM during a transfer with no CPU involved copying bytes along the way -- and a DMA engine mid-transfer has no way to be told "wait, the OS just moved that page." Pinned (page-locked) memory is the CUDA driver's answer: memory it has explicitly asked the OS to guarantee will never be paged out and never move, for exactly as long as it stays allocated, which is precisely the guarantee a DMA transfer needs and pageable memory cannot offer.
Background¶
Requesting pinned memory means calling into the CUDA Runtime API directly -- cudaHostAlloc() or its simpler wrapper cudaMallocHost() -- rather than the ordinary C++ allocators Section J.2 already covered; there is no way to page-lock memory malloc() or new already handed back after the fact; the memory must be requested as pinned from the start. Freeing it uses the matching cudaFreeHost(), not free() or delete -- one more allocation/deallocation family this book's own "match your allocator to your deallocator" rule (already established in Section J.2's own [COMMON TRAP]) extends to.
The concrete payoff is asynchronous transfer. cudaMemcpyAsync(), given a pinned host pointer, a device pointer, and a CUDA stream, genuinely overlaps a host-to-device (or device-to-host) transfer with concurrent kernel execution on a different stream, because the DMA engine can safely read the pinned buffer directly, without the CPU's involvement, for the transfer's entire duration. Given a pageable host pointer instead, cudaMemcpyAsync() does not refuse the call -- there is nothing in the function's own C signature that could distinguish a pinned pointer from a pageable one at compile time -- it silently falls back to copying through a small, temporary pinned staging buffer the driver owns internally, which serializes the transfer and eliminates the overlap the caller's code was written to get, with no error, no warning, and no visible difference in the program's own correctness. This is precisely why the bug is easy to introduce and easy to miss: the program still runs, still produces the right numbers, and simply never achieves the performance its own use of streams and cudaMemcpyAsync was written to obtain.
This book's own build environment (Getting Started) has a full CUDA toolkit but no NVIDIA GPU and no NVIDIA driver -- so every call this section makes is genuinely compiled with the real nvcc against the real CUDA Runtime headers, and genuinely executed (not stubbed, not simulated), reporting this exact machine's own honest answer to every one of these questions. On a real GPU, cudaHostAlloc, cudaMallocHost, and cudaMemcpyAsync would all report cudaSuccess; the actual bandwidth improvement pinned memory delivers over pageable memory is a genuinely measured runtime number, and following the same discipline Getting Started already states for every kernel's own timing claims in this book, that measurement is deferred to the same real-hardware verification pass the rest of this book's runtime numbers await, rather than quoted here as a number this environment did not and could not genuinely measure itself.
File: 05_pinned_vs_pageable.cu¶
#include <cstdio>
#include <cstdlib>
#include <cuda_runtime.h>
// Section J.4 -- "pageable" memory is just ordinary heap memory: whatever
// malloc(), new, or the C++ default allocator hands back, backed by pages
// the operating system is free to swap out to disk or relocate at any time.
// "Pinned" (a.k.a. page-locked) memory is memory the CUDA driver has asked
// the OS to guarantee will NEVER be paged out and will NEVER move, for as
// long as it's allocated -- a guarantee ordinary malloc()/new memory does
// not and cannot make. That guarantee is what lets the GPU's DMA (Direct
// Memory Access) engine read directly from host RAM during a transfer,
// with no CPU involvement copying bytes along the way; pageable memory
// offers the DMA engine no such promise, so the CUDA driver silently
// inserts an extra copy of its own -- into a temporary pinned staging
// buffer it owns -- before the real host-to-device transfer can begin.
//
// This environment has a full CUDA toolkit (nvcc, the real cuda_runtime.h,
// the real libcudart) but no NVIDIA GPU and no NVIDIA driver installed, the
// same environment this whole book was written in (see Getting Started).
// Every call below is genuinely compiled against the real CUDA Runtime API
// and genuinely executed -- not stubbed out -- so every reported error
// below is this machine's own honest, real answer to "is a CUDA-capable
// device present," not a fabricated placeholder.
void report(const char* what, cudaError_t err) {
printf("%-45s: %s (%s)\n", what, cudaGetErrorName(err), cudaGetErrorString(err));
}
int main() {
printf("=== Section 1: pageable memory needs NO CUDA call at all -- it's just malloc/new ===\n");
const size_t N = 1 << 20; // 1 M floats, 4 MB
float* pageable_malloc = (float*)malloc(N * sizeof(float));
float* pageable_new = new float[N];
printf("pageable_malloc allocated via plain malloc()? %s\n", (pageable_malloc != nullptr) ? "true" : "false");
printf("pageable_new allocated via plain new[]? %s\n", (pageable_new != nullptr) ? "true" : "false");
printf("(the OS is free to swap either buffer's pages out to disk, or move them,\n");
printf(" at any point while they sit idle -- ordinary heap memory carries no\n");
printf(" guarantee against either, and CUDA's own DMA engine cannot be handed\n");
printf(" a physical address that might change out from under it mid-transfer)\n");
free(pageable_malloc);
delete[] pageable_new;
printf("\n=== Section 2: pinned memory REQUIRES the CUDA Runtime API -- genuinely called here ===\n");
void* pinned_a = nullptr;
cudaError_t err_a = cudaHostAlloc(&pinned_a, N * sizeof(float), cudaHostAllocDefault);
report("cudaHostAlloc(4 MB, cudaHostAllocDefault)", err_a);
void* pinned_b = nullptr;
cudaError_t err_b = cudaMallocHost(&pinned_b, N * sizeof(float));
report("cudaMallocHost(4 MB)", err_b);
void* pinned_mapped = nullptr;
cudaError_t err_c = cudaHostAlloc(&pinned_mapped, N * sizeof(float), cudaHostAllocMapped);
report("cudaHostAlloc(4 MB, cudaHostAllocMapped)", err_c);
int device_count = 0;
cudaError_t err_count = cudaGetDeviceCount(&device_count);
report("cudaGetDeviceCount", err_count);
printf("device_count = %d\n", device_count);
printf("\nall three pinned-allocation calls above are genuine calls into the real,\n");
printf("linked libcudart -- the SAME real function a working GPU machine calls,\n");
printf("compiled by the SAME nvcc that ships every other .cu file in this book.\n");
printf("The error reported is this machine's own honest answer (no driver, no\n");
printf("device) to a question that a real GPU machine would answer 'success' to\n");
printf("-- exactly the distinction Getting Started already draws for every kernel\n");
printf("in this book: genuinely compiled everywhere, genuinely EXECUTED here only\n");
printf("for calls that don't require a device to answer honestly either way.\n");
// Attempting to free a pointer a failed allocation never actually
// populated is itself worth handling correctly -- cudaFreeHost on a
// pointer that IS still nullptr (because the allocation above failed)
// is checked for explicitly, exactly the discipline a real allocator
// wrapper needs regardless of which environment it happens to run on.
if (pinned_a) cudaFreeHost(pinned_a);
if (pinned_b) cudaFreeHost(pinned_b);
if (pinned_mapped) cudaFreeHost(pinned_mapped);
printf("\ncudaFreeHost called on every pointer the allocations above actually populated\n");
printf("\n=== Section 3: cudaMemcpyAsync compiles and links identically for EITHER kind of host pointer ===\n");
// This is the structural point Section J.4's own Background section
// makes in prose: the compiler and the CUDA Runtime API accept a
// plain malloc()'d pointer here exactly as readily as a cudaHostAlloc'd
// one -- cudaMemcpyAsync's own signature has no way to require pinned
// memory at compile time. Passed a pageable pointer, the driver does
// NOT refuse the call; it silently falls back to a synchronous copy
// through its own internal pinned staging buffer, which is precisely
// why this bug is easy to introduce and easy to miss: the program
// still runs, still produces correct results, and simply never
// achieves the overlap it was written to get.
cudaStream_t stream;
cudaError_t err_stream = cudaStreamCreate(&stream);
report("cudaStreamCreate", err_stream);
float* device_buf = nullptr;
cudaError_t err_dmalloc = cudaMalloc(&device_buf, N * sizeof(float));
report("cudaMalloc(4 MB, device)", err_dmalloc);
float* pageable_src = (float*)malloc(N * sizeof(float));
cudaError_t err_copy_pageable = cudaMemcpyAsync(
device_buf, pageable_src, N * sizeof(float), cudaMemcpyHostToDevice, stream);
report("cudaMemcpyAsync(pageable src, stream)", err_copy_pageable);
void* pinned_src = nullptr;
cudaHostAlloc(&pinned_src, N * sizeof(float), cudaHostAllocDefault);
cudaError_t err_copy_pinned = cudaMemcpyAsync(
device_buf, pinned_src, N * sizeof(float), cudaMemcpyHostToDevice, stream);
report("cudaMemcpyAsync(pinned src, stream)", err_copy_pinned);
printf("\nboth calls above compile identically and are accepted identically by the\n");
printf("API's own signature -- there is no compile-time distinction between them.\n");
printf("On a real GPU, only the SECOND call (pinned source) genuinely overlaps with\n");
printf("concurrent kernel execution on another stream; the FIRST (pageable source)\n");
printf("silently degrades to a blocking copy, with no error, no warning, and no\n");
printf("visible sign in this program's own output that anything was lost.\n");
free(pageable_src);
if (pinned_src) cudaFreeHost(pinned_src);
if (device_buf) cudaFree(device_buf);
cudaStreamDestroy(stream);
return 0;
}
Genuinely compiled with nvcc and genuinely run in this book's own no-driver, no-device environment:
=== Section 1: pageable memory needs NO CUDA call at all -- it's just malloc/new ===
pageable_malloc allocated via plain malloc()? true
pageable_new allocated via plain new[]? true
(the OS is free to swap either buffer's pages out to disk, or move them,
at any point while they sit idle -- ordinary heap memory carries no
guarantee against either, and CUDA's own DMA engine cannot be handed
a physical address that might change out from under it mid-transfer)
=== Section 2: pinned memory REQUIRES the CUDA Runtime API -- genuinely called here ===
cudaHostAlloc(4 MB, cudaHostAllocDefault) : cudaErrorInsufficientDriver (CUDA driver version is insufficient for CUDA runtime version)
cudaMallocHost(4 MB) : cudaErrorInsufficientDriver (CUDA driver version is insufficient for CUDA runtime version)
cudaHostAlloc(4 MB, cudaHostAllocMapped) : cudaErrorInsufficientDriver (CUDA driver version is insufficient for CUDA runtime version)
cudaGetDeviceCount : cudaErrorInsufficientDriver (CUDA driver version is insufficient for CUDA runtime version)
device_count = 0
all three pinned-allocation calls above are genuine calls into the real,
linked libcudart -- the SAME real function a working GPU machine calls,
compiled by the SAME nvcc that ships every other .cu file in this book.
The error reported is this machine's own honest answer (no driver, no
device) to a question that a real GPU machine would answer 'success' to
-- exactly the distinction Getting Started already draws for every kernel
in this book: genuinely compiled everywhere, genuinely EXECUTED here only
for calls that don't require a device to answer honestly either way.
cudaFreeHost called on every pointer the allocations above actually populated
=== Section 3: cudaMemcpyAsync compiles and links identically for EITHER kind of host pointer ===
cudaStreamCreate : cudaErrorInsufficientDriver (CUDA driver version is insufficient for CUDA runtime version)
cudaMalloc(4 MB, device) : cudaErrorInsufficientDriver (CUDA driver version is insufficient for CUDA runtime version)
cudaMemcpyAsync(pageable src, stream) : cudaErrorInsufficientDriver (CUDA driver version is insufficient for CUDA runtime version)
cudaMemcpyAsync(pinned src, stream) : cudaErrorInsufficientDriver (CUDA driver version is insufficient for CUDA runtime version)
both calls above compile identically and are accepted identically by the
API's own signature -- there is no compile-time distinction between them.
On a real GPU, only the SECOND call (pinned source) genuinely overlaps with
concurrent kernel execution on another stream; the FIRST (pageable source)
silently degrades to a blocking copy, with no error, no warning, and no
visible sign in this program's own output that anything was lost.
Worked Example J.4.1 -- Pageable needs nothing from CUDA; pinned needs the Runtime API, honestly¶
Section 1 allocates two 4 MB pageable buffers with plain malloc() and new[] -- both genuinely succeed, with zero CUDA calls of any kind, because pageable memory is not a CUDA concept at all; it's just heap memory. Section 2 calls cudaHostAlloc(), cudaMallocHost(), and cudaHostAlloc() again with the cudaHostAllocMapped flag (for zero-copy access) -- all three are genuine calls into the real, linked libcudart, and all three genuinely report cudaErrorInsufficientDriver on this machine, the same honest, no-driver answer cudaGetDeviceCount() reports (device_count = 0) two lines later. Section 3 confirms the structural point directly: cudaMemcpyAsync() is called once with a pageable source pointer and once with a pinned source pointer, on the identical stream, with identical arguments otherwise -- both calls compile identically and both report the identical honest error on this machine, because nothing about the function's own signature could distinguish them even in principle. That's the exact mechanism Section J.4's own Background section describes: the compiler and the API accept either pointer without complaint, and only a real device's own DMA engine treats them differently, silently, at runtime.
[COMMON TRAP] Pinned memory is not "strictly better" pageable memory a program should default to everywhere. Page-locking memory is itself a real, measurable cost -- it involves genuine kernel-level bookkeeping to mark the pages non-swappable, and it removes those physical pages from the pool available to every other process on the machine, including the OS's own paging system, for as long as the pinned allocation lives. A program that pins far more memory than it actually needs to keep in flight for DMA can genuinely starve the rest of the system of usable RAM, and pinning/unpinning a large buffer repeatedly, on a hot path, is measurably slower than allocating it once, pinned, and reusing it -- the correct pattern is to identify the specific buffers that are genuinely repeatedly transferred (Chapter 22's own bond-portfolio structures, say, streamed to the GPU once per pricing pass) and pin only those, not every host allocation in the program indiscriminately.
Complete Runnable Code¶
Every file below is reproduced in full, genuinely compiled and run exactly as shown in its own section above.
File: 01_stack_layout_and_overflow.cpp¶
#include <cstdio>
#include <cstdint>
#include <cstdlib>
#include <cstring>
// Section J.1 -- the stack is automatic storage: every local variable's
// address and lifetime are decided entirely by the compiler at the point
// of declaration, unwound automatically (in exactly reverse order) the
// moment control leaves the enclosing scope. This file makes both halves
// of that claim genuinely observable: real addresses, printed directly,
// and real construction/destruction order, printed by a small RAII type.
struct Tracer {
const char* name;
explicit Tracer(const char* n) : name(n) {
printf(" constructed: %s\n", name);
}
~Tracer() {
printf(" destroyed: %s\n", name);
}
};
// Addresses themselves are NOT locked anywhere in this file's own output --
// ASLR (Address Space Layout Randomization, on by default on this and every
// modern Linux/macOS/Windows machine) genuinely relocates the whole stack to
// a different base address on every single run, specifically to make stack
// addresses useless to an attacker guessing at them. What IS stable across
// every run, ASLR or not, is the RELATIONSHIP between addresses within the
// same invocation: which frame's locals sit closer to the top of memory,
// and the byte distance between two variables declared next to each other
// in the same frame. This file measures and prints exactly those relationships,
// genuinely computed fresh on every run, rather than a raw address a rerun
// could never reproduce.
intptr_t frame_c(intptr_t caller_frame_addr) {
// Three locals, declared in order a, b, c.
int a = 1;
int b = 2;
int c = 3;
long long b_minus_a = (intptr_t)&b - (intptr_t)&a;
long long c_minus_b = (intptr_t)&c - (intptr_t)&b;
printf("frame_c: sizeof(int)=%zu, (&b - &a)=%lld bytes, (&c - &b)=%lld bytes\n",
sizeof(int), b_minus_a, c_minus_b);
printf("frame_c: this frame's locals sit at a lower address than frame_b's own local? %s\n",
((intptr_t)&a < caller_frame_addr) ? "true" : "false");
return (intptr_t)&a;
}
intptr_t frame_b(intptr_t caller_frame_addr) {
int local_in_b = 42;
printf("frame_b: this frame's local sits at a lower address than frame_a's own local? %s\n",
((intptr_t)&local_in_b < caller_frame_addr) ? "true" : "false");
return frame_c((intptr_t)&local_in_b);
}
intptr_t frame_a(intptr_t caller_frame_addr) {
int local_in_a = 7;
printf("frame_a: this frame's local sits at a lower address than main's own local? %s\n",
((intptr_t)&local_in_a < caller_frame_addr) ? "true" : "false");
return frame_b((intptr_t)&local_in_a);
}
void raii_scope_demo() {
printf("entering raii_scope_demo\n");
Tracer t1("t1 (declared first)");
{
Tracer t2("t2 (nested scope)");
Tracer t3("t3 (nested scope)");
printf(" inside nested scope\n");
}
printf(" back in outer scope\n");
Tracer t4("t4 (declared last, outer scope)");
printf("leaving raii_scope_demo\n");
}
// --trigger-overflow: an unbounded recursion, each frame consuming a real
// (non-optimizable-away) 1 KB of stack space via a local array the compiler
// cannot prove is unused. Compiled at -O0 specifically so this recursion is
// NOT turned into a loop by tail-call optimization -- the whole point is to
// genuinely exhaust the stack, one real frame at a time.
long long recurse_and_consume_stack(long long depth) {
char padding[1024];
memset(padding, (int)(depth & 0xFF), sizeof(padding));
if (depth % 1000 == 0) {
printf("depth=%lld, ~%lld KB of stack consumed so far\n", depth, depth);
fflush(stdout);
}
return padding[0] + recurse_and_consume_stack(depth + 1);
}
int main(int argc, char** argv) {
bool trigger_overflow = (argc > 1 && strcmp(argv[1], "--trigger-overflow") == 0);
if (trigger_overflow) {
printf("=== Section 4: a genuine stack overflow, one real 1 KB frame at a time ===\n");
fflush(stdout);
long long result = recurse_and_consume_stack(0);
// Unreachable under a real stack overflow -- the process crashes
// (SIGSEGV) inside recurse_and_consume_stack itself, before this
// line, or before the printf above's buffered output would even
// flush without the explicit fflush() calls this file already uses.
printf("unreachable: result=%lld\n", result);
return 0;
}
int main_local = 99;
printf("=== Section 1: nested stack frames, real address RELATIONSHIPS (not raw addresses -- ASLR note above main()) ===\n");
frame_a((intptr_t)&main_local);
printf("\n=== Section 2: LIFO destruction order, genuinely observed ===\n");
raii_scope_demo();
return 0;
}
File: 02_malloc_vs_new.cpp¶
#include <cstdio>
#include <cstdlib>
#include <new>
#include <cstring>
#include <cstdint>
#include <cstddef>
// Section J.2 -- malloc()/free() and new/delete both hand out heap
// (free-store) memory, but only one of the two pairs knows anything about
// C++ objects. malloc() returns a raw, uninitialized block of bytes and
// free() returns it, full stop -- no constructor, no destructor, ever.
// new/delete are language operators, not library functions: new allocates
// AND genuinely calls the constructor, delete genuinely calls the
// destructor AND deallocates. This file makes the difference observable,
// not asserted.
struct Sentinel {
unsigned int magic;
char label[16];
Sentinel() : magic(0xABCD1234) {
std::strncpy(label, "constructed", sizeof(label) - 1);
label[sizeof(label) - 1] = '\0';
printf(" Sentinel::Sentinel() genuinely ran -- magic set to 0x%X\n", magic);
}
~Sentinel() {
printf(" Sentinel::~Sentinel() genuinely ran -- magic was 0x%X\n", magic);
}
};
void section1_malloc_skips_constructor() {
printf("=== Section 1: malloc() does NOT call the constructor ===\n");
// malloc() returns raw bytes it zero-fills with nothing in particular --
// whatever this heap page's memory previously held. Reinterpreting that
// raw block as a Sentinel* does NOT run Sentinel::Sentinel(); "magic"
// is whatever bytes malloc happened to hand back, not 0xABCD1234.
void* raw = malloc(sizeof(Sentinel));
Sentinel* s = reinterpret_cast<Sentinel*>(raw);
printf(" after malloc(sizeof(Sentinel)): no constructor message printed above this line\n");
printf(" s->magic == 0xABCD1234 (the constructor's own value)? %s\n",
(s->magic == 0xABCD1234) ? "true" : "false");
// Calling free() on this block is correct (it was allocated by malloc,
// and no destructor needs to run because none was ever constructed) --
// but nothing here ever invoked ~Sentinel(), which is exactly the point.
free(raw);
}
void section2_new_calls_constructor() {
printf("\n=== Section 2: new DOES call the constructor ===\n");
Sentinel* s = new Sentinel(); // the constructor message prints from inside this call
printf(" s->magic == 0xABCD1234 (the constructor's own value)? %s\n",
(s->magic == 0xABCD1234) ? "true" : "false");
delete s; // the destructor message prints from inside this call
}
void section3_failure_modes() {
printf("\n=== Section 3: malloc() returns nullptr on failure; new throws ===\n");
// An allocation size designed to fail on any real system: essentially
// the entire 64-bit address space at once. malloc() reports failure
// the C way -- a nullptr return, checked explicitly by the caller.
void* huge = malloc((size_t)-1);
printf(" malloc((size_t)-1) returned nullptr? %s\n", (huge == nullptr) ? "true" : "false");
// operator new reports the identical failure the C++ way: it throws
// std::bad_alloc instead of returning a sentinel value, specifically so
// a caller can't forget to check a return value the way huge could be
// (and, in real bug reports, often is) silently ignored above.
bool threw_bad_alloc = false;
try {
void* p = ::operator new((size_t)-1);
(void)p;
} catch (const std::bad_alloc& e) {
threw_bad_alloc = true;
printf(" operator new((size_t)-1) threw std::bad_alloc: \"%s\"\n", e.what());
}
printf(" operator new((size_t)-1) threw std::bad_alloc? %s\n", threw_bad_alloc ? "true" : "false");
// std::nothrow opts back into malloc's own nullptr-on-failure contract,
// for exactly the cases (an allocator hot path with no exception
// handling nearby) where a thrown exception isn't what the caller wants.
void* p2 = ::operator new((size_t)-1, std::nothrow);
printf(" operator new((size_t)-1, std::nothrow) returned nullptr? %s\n", (p2 == nullptr) ? "true" : "false");
}
void section4_array_forms() {
printf("\n=== Section 4: new[] / delete[] run the constructor/destructor for EVERY element ===\n");
const int N = 3;
Sentinel* arr = new Sentinel[N]; // N constructor calls, printed above this line
printf(" constructed %d Sentinel objects via new[]\n", N);
delete[] arr; // N destructor calls, printed above this line
printf(" destroyed all %d via delete[]\n", N);
}
void section5_alignment() {
printf("\n=== Section 5: over-aligned allocation, both the malloc-family and new way ===\n");
// Ordinary malloc()/new only guarantee alignof(std::max_align_t) (16
// bytes on this platform's ABI) -- correct for any built-in scalar
// type, but not enough for a type that genuinely demands more, such as
// a 64-byte cache-line-aligned struct used to avoid false sharing.
struct alignas(64) CacheLineAligned { char bytes[64]; };
printf(" alignof(CacheLineAligned) = %zu\n", alignof(CacheLineAligned));
printf(" alignof(std::max_align_t) = %zu (ordinary malloc/new's own guarantee)\n", alignof(std::max_align_t));
void* aligned_raw = std::aligned_alloc(64, sizeof(CacheLineAligned));
uintptr_t addr = (uintptr_t)aligned_raw;
printf(" std::aligned_alloc(64, ...) address is a multiple of 64? %s\n",
(addr % 64 == 0) ? "true" : "false");
free(aligned_raw);
// The C++17 language-level equivalent: operator new with an explicit
// std::align_val_t, which delete must be given back symmetrically.
CacheLineAligned* obj = new (std::align_val_t(64)) CacheLineAligned();
uintptr_t addr2 = (uintptr_t)obj;
printf(" new(std::align_val_t(64)) address is a multiple of 64? %s\n",
(addr2 % 64 == 0) ? "true" : "false");
::operator delete(obj, std::align_val_t(64));
}
int main() {
section1_malloc_skips_constructor();
section2_new_calls_constructor();
section3_failure_modes();
section4_array_forms();
section5_alignment();
return 0;
}
File: 03_alloc_dealloc_mismatch_asan.cpp¶
#include <cstdio>
#include <cstdlib>
#include <cstring>
// Section J.2 continued -- new[] and delete (without the brackets) are a
// MISMATCHED pair: delete alone tells the runtime "one object lived here,"
// while new[] recorded "N objects lived here, plus a hidden count." Freeing
// a new[]'d array with plain delete reads that hidden bookkeeping wrong and
// corrupts the heap -- undefined behavior that can run "successfully" with
// no visible symptom for a long time, on an unlucky day, in production.
// This file triggers the real bug for real, under AddressSanitizer, rather
// than describing it -- the same standard this book already holds every
// claimed memory bug to (Appendix D.6, Appendix G.2, Appendix G.5).
struct Row {
double values[4];
Row() { for (double& v : values) v = 0.0; }
};
int main(int argc, char** argv) {
bool trigger = (argc > 1 && std::strcmp(argv[1], "--trigger-mismatch") == 0);
printf("allocating 4 Row objects with new[]...\n");
Row* rows = new Row[4];
rows[0].values[0] = 3.14;
printf("rows[0].values[0] = %f\n", rows[0].values[0]);
if (!trigger) {
printf("freeing correctly with delete[] (matched, no bug)\n");
delete[] rows;
printf("done -- no mismatch triggered in this run\n");
return 0;
}
printf("freeing INCORRECTLY with plain delete (mismatched -- new[] / delete)...\n");
delete rows; // genuine alloc-dealloc-size mismatch, caught by ASan below
printf("unreachable under a working ASan build\n");
return 0;
}
File: 04_placement_new.cpp¶
#include <cstdio>
#include <cstdlib>
#include <new>
#include <cstring>
// Section J.3 -- placement new separates the two things ordinary new
// bundles together: allocating raw storage, and constructing an object
// inside it. Placement new takes storage that ALREADY exists (a stack
// buffer, a malloc'd block, a slot inside a custom pool allocator -- the
// exact situation Chapter 11's own arena allocator and Appendix G's own
// Rule-of-Five types are built to manage) and runs a constructor at that
// exact address, with no allocation of its own. The object built this way
// must have its destructor called explicitly and by name -- delete would
// try to free storage placement new never allocated in the first place.
struct Widget {
int id;
double value;
Widget(int i, double v) : id(i), value(v) {
printf(" Widget::Widget(%d, %.1f) constructed\n", id, value);
}
~Widget() {
printf(" Widget::~Widget() destroying id=%d\n", id);
}
};
void section1_placement_new_on_stack_buffer() {
printf("=== Section 1: placement new into an already-existing STACK buffer ===\n");
// A raw, unconstructed buffer with no Widget in it yet -- just bytes,
// correctly SIZED and ALIGNED for a Widget, sitting on the stack.
alignas(Widget) unsigned char buffer[sizeof(Widget)];
printf(" raw buffer allocated on the stack, sizeof=%zu, alignof=%zu\n",
sizeof(buffer), alignof(Widget));
// Placement new: no allocation happens here at all -- it constructs a
// Widget IN buffer's own memory, at the exact address &buffer[0].
Widget* w = new (buffer) Widget(1, 3.5);
printf(" &buffer[0] == (void*)w? %s\n", ((void*)&buffer[0] == (void*)w) ? "true" : "false");
printf(" w->id=%d w->value=%.1f\n", w->id, w->value);
// Explicit destructor call BY NAME -- this is the only correct way to
// end this object's lifetime. `delete w;` would be a serious bug: it
// would call operator delete on a stack address that was never
// returned by operator new, corrupting the heap's own bookkeeping.
w->~Widget();
printf(" destructor called explicitly by name (w->~Widget()), NOT via delete\n");
}
void section2_placement_new_in_pool_slot() {
printf("\n=== Section 2: placement new into a slot of a malloc'd pool (the arena-allocator pattern) ===\n");
// A pool of 3 raw, unconstructed Widget-sized slots -- exactly the
// shape of Chapter 11's own Arena allocator: one big raw allocation,
// objects constructed into specific offsets within it on demand.
const int SLOT_COUNT = 3;
unsigned char* pool = static_cast<unsigned char*>(
std::aligned_alloc(alignof(Widget), SLOT_COUNT * sizeof(Widget)));
printf(" pool allocated: %d raw slots, %zu bytes each\n", SLOT_COUNT, sizeof(Widget));
Widget* widgets[SLOT_COUNT];
for (int i = 0; i < SLOT_COUNT; i++) {
unsigned char* slot = pool + i * sizeof(Widget);
widgets[i] = new (slot) Widget(100 + i, 1.5 * i);
}
printf(" all %d widgets constructed in-place inside one pool allocation\n", SLOT_COUNT);
for (int i = 0; i < SLOT_COUNT; i++) {
printf(" widgets[%d]->id=%d\n", i, widgets[i]->id);
}
// Destroy each object explicitly (in-place, no per-object deallocation --
// there IS no per-object allocation to give back), then free the pool
// itself exactly once, with the same allocator that created it.
for (int i = 0; i < SLOT_COUNT; i++) {
widgets[i]->~Widget();
}
printf(" all %d widgets destroyed explicitly, in-place\n", SLOT_COUNT);
free(pool);
printf(" pool freed once, as the single raw allocation it always was\n");
}
int main() {
section1_placement_new_on_stack_buffer();
section2_placement_new_in_pool_slot();
return 0;
}
File: 05_pinned_vs_pageable.cu¶
#include <cstdio>
#include <cstdlib>
#include <cuda_runtime.h>
// Section J.4 -- "pageable" memory is just ordinary heap memory: whatever
// malloc(), new, or the C++ default allocator hands back, backed by pages
// the operating system is free to swap out to disk or relocate at any time.
// "Pinned" (a.k.a. page-locked) memory is memory the CUDA driver has asked
// the OS to guarantee will NEVER be paged out and will NEVER move, for as
// long as it's allocated -- a guarantee ordinary malloc()/new memory does
// not and cannot make. That guarantee is what lets the GPU's DMA (Direct
// Memory Access) engine read directly from host RAM during a transfer,
// with no CPU involvement copying bytes along the way; pageable memory
// offers the DMA engine no such promise, so the CUDA driver silently
// inserts an extra copy of its own -- into a temporary pinned staging
// buffer it owns -- before the real host-to-device transfer can begin.
//
// This environment has a full CUDA toolkit (nvcc, the real cuda_runtime.h,
// the real libcudart) but no NVIDIA GPU and no NVIDIA driver installed, the
// same environment this whole book was written in (see Getting Started).
// Every call below is genuinely compiled against the real CUDA Runtime API
// and genuinely executed -- not stubbed out -- so every reported error
// below is this machine's own honest, real answer to "is a CUDA-capable
// device present," not a fabricated placeholder.
void report(const char* what, cudaError_t err) {
printf("%-45s: %s (%s)\n", what, cudaGetErrorName(err), cudaGetErrorString(err));
}
int main() {
printf("=== Section 1: pageable memory needs NO CUDA call at all -- it's just malloc/new ===\n");
const size_t N = 1 << 20; // 1 M floats, 4 MB
float* pageable_malloc = (float*)malloc(N * sizeof(float));
float* pageable_new = new float[N];
printf("pageable_malloc allocated via plain malloc()? %s\n", (pageable_malloc != nullptr) ? "true" : "false");
printf("pageable_new allocated via plain new[]? %s\n", (pageable_new != nullptr) ? "true" : "false");
printf("(the OS is free to swap either buffer's pages out to disk, or move them,\n");
printf(" at any point while they sit idle -- ordinary heap memory carries no\n");
printf(" guarantee against either, and CUDA's own DMA engine cannot be handed\n");
printf(" a physical address that might change out from under it mid-transfer)\n");
free(pageable_malloc);
delete[] pageable_new;
printf("\n=== Section 2: pinned memory REQUIRES the CUDA Runtime API -- genuinely called here ===\n");
void* pinned_a = nullptr;
cudaError_t err_a = cudaHostAlloc(&pinned_a, N * sizeof(float), cudaHostAllocDefault);
report("cudaHostAlloc(4 MB, cudaHostAllocDefault)", err_a);
void* pinned_b = nullptr;
cudaError_t err_b = cudaMallocHost(&pinned_b, N * sizeof(float));
report("cudaMallocHost(4 MB)", err_b);
void* pinned_mapped = nullptr;
cudaError_t err_c = cudaHostAlloc(&pinned_mapped, N * sizeof(float), cudaHostAllocMapped);
report("cudaHostAlloc(4 MB, cudaHostAllocMapped)", err_c);
int device_count = 0;
cudaError_t err_count = cudaGetDeviceCount(&device_count);
report("cudaGetDeviceCount", err_count);
printf("device_count = %d\n", device_count);
printf("\nall three pinned-allocation calls above are genuine calls into the real,\n");
printf("linked libcudart -- the SAME real function a working GPU machine calls,\n");
printf("compiled by the SAME nvcc that ships every other .cu file in this book.\n");
printf("The error reported is this machine's own honest answer (no driver, no\n");
printf("device) to a question that a real GPU machine would answer 'success' to\n");
printf("-- exactly the distinction Getting Started already draws for every kernel\n");
printf("in this book: genuinely compiled everywhere, genuinely EXECUTED here only\n");
printf("for calls that don't require a device to answer honestly either way.\n");
// Attempting to free a pointer a failed allocation never actually
// populated is itself worth handling correctly -- cudaFreeHost on a
// pointer that IS still nullptr (because the allocation above failed)
// is checked for explicitly, exactly the discipline a real allocator
// wrapper needs regardless of which environment it happens to run on.
if (pinned_a) cudaFreeHost(pinned_a);
if (pinned_b) cudaFreeHost(pinned_b);
if (pinned_mapped) cudaFreeHost(pinned_mapped);
printf("\ncudaFreeHost called on every pointer the allocations above actually populated\n");
printf("\n=== Section 3: cudaMemcpyAsync compiles and links identically for EITHER kind of host pointer ===\n");
// This is the structural point Section J.4's own Background section
// makes in prose: the compiler and the CUDA Runtime API accept a
// plain malloc()'d pointer here exactly as readily as a cudaHostAlloc'd
// one -- cudaMemcpyAsync's own signature has no way to require pinned
// memory at compile time. Passed a pageable pointer, the driver does
// NOT refuse the call; it silently falls back to a synchronous copy
// through its own internal pinned staging buffer, which is precisely
// why this bug is easy to introduce and easy to miss: the program
// still runs, still produces correct results, and simply never
// achieves the overlap it was written to get.
cudaStream_t stream;
cudaError_t err_stream = cudaStreamCreate(&stream);
report("cudaStreamCreate", err_stream);
float* device_buf = nullptr;
cudaError_t err_dmalloc = cudaMalloc(&device_buf, N * sizeof(float));
report("cudaMalloc(4 MB, device)", err_dmalloc);
float* pageable_src = (float*)malloc(N * sizeof(float));
cudaError_t err_copy_pageable = cudaMemcpyAsync(
device_buf, pageable_src, N * sizeof(float), cudaMemcpyHostToDevice, stream);
report("cudaMemcpyAsync(pageable src, stream)", err_copy_pageable);
void* pinned_src = nullptr;
cudaHostAlloc(&pinned_src, N * sizeof(float), cudaHostAllocDefault);
cudaError_t err_copy_pinned = cudaMemcpyAsync(
device_buf, pinned_src, N * sizeof(float), cudaMemcpyHostToDevice, stream);
report("cudaMemcpyAsync(pinned src, stream)", err_copy_pinned);
printf("\nboth calls above compile identically and are accepted identically by the\n");
printf("API's own signature -- there is no compile-time distinction between them.\n");
printf("On a real GPU, only the SECOND call (pinned source) genuinely overlaps with\n");
printf("concurrent kernel execution on another stream; the FIRST (pageable source)\n");
printf("silently degrades to a blocking copy, with no error, no warning, and no\n");
printf("visible sign in this program's own output that anything was lost.\n");
free(pageable_src);
if (pinned_src) cudaFreeHost(pinned_src);
if (device_buf) cudaFree(device_buf);
cudaStreamDestroy(stream);
return 0;
}
Chapter Summary¶
The stack is automatic, LIFO, compiler-managed storage: every local variable's address and lifetime are decided at compile time from the source code's own shape, unwound in strictly reverse order on scope exit with zero exceptions, and bounded by a real, finite limit this appendix genuinely exhausted (a real SIGSEGV, at a depth of roughly 7,000-8,000 frames consuming 1 KB each against this machine's 8 MB stack limit) rather than merely describing. The heap is the opposite in every dimension -- programmer-controlled lifetime, survives the allocating function's own return -- and malloc()/free() versus new/delete differ in exactly one property worth remembering above all others: only new/delete genuinely know about C++ objects, calling constructors and destructors that malloc()/free() never touch at all, confirmed here by a real constructor message that only one of the two families ever printed. Mixing allocation and deallocation families is undefined behavior with a real, genuinely triggered consequence (an alloc-dealloc-mismatch, caught live under AddressSanitizer, following this book's own established discipline of locking only a report's stable, address-free summary line). Placement new separates allocation from construction entirely, letting an object be built at an address already owned by a stack buffer or a hand-built pool -- exactly Chapter 11's own Arena pattern -- with its lifetime ended by an explicit destructor call rather than delete. And pinned (page-locked) memory is the CUDA driver's guarantee, unavailable from ordinary malloc()/new, that a buffer will never be paged out or moved -- required for a cudaMemcpyAsync() transfer to genuinely overlap with concurrent kernel execution, silently and invisibly degraded to a blocking copy when that guarantee is missing, demonstrated here with the real CUDA Runtime API, genuinely compiled and genuinely called, honestly reporting this book's own no-driver, no-device environment rather than a number it was never in a position to measure.
Self-Check Questions¶
- Section J.1's own
frame_a/frame_b/frame_cchain checks address relationships (which frame's local is lower) rather than printing raw addresses directly. Given that ASLR relocates the entire stack to a new base address on every run, explain precisely why a relationship-based check stays reproducible while a raw address never could. - A function returns a pointer to one of its own local variables (
int* p = &local; return p;). Using Section J.1's own account of automatic storage duration, explain exactly what's wrong with dereferencingpafter the function returns, and why the bug might not crash immediately, or ever, depending on what happens to occupy that stack region afterward. - Worked Example J.2.1 shows
malloc(sizeof(Sentinel))producing aSentinelwhosemagicfield does not equal the constructor's own expected value. Explain precisely why -- what, if anything, didmalloc()actually initialize that memory to? - Section J.2's own
[COMMON TRAP]and Worked Example J.2.3 both concern matching allocation and deallocation functions. Name the three matched pairs this appendix establishes, and explain what specifically goes wrong internally whennew[]'s own allocation is freed with plaindeleteinstead ofdelete[]. - Section J.3 states that calling
deleteon a placement-new-constructed object is a serious bug, distinct from simply "forgetting" to call the destructor. Using the distinction between placement new'soperator new(size_t, void*)overload and ordinaryoperator new(size_t), explain exactly whatdeletewould attempt to do in this case, and why it's dangerous rather than merely a no-op. - Section J.4 states that
cudaMemcpyAsync()cannot distinguish a pinned pointer from a pageable one "even in principle," at compile time. Using the two calls Worked Example J.4.1 makes back to back (identical arguments except for which kind of pointer), explain why this makes the bug of accidentally passing pageable memory to an async transfer especially easy to introduce and especially hard to detect through ordinary testing. - This book's own build environment has no NVIDIA driver at all, and genuinely reports
cudaErrorInsufficientDriverfor every pinned-memory call in Section J.4. Appendix C.5's owncudaMallocManagedexample, run in a different real environment, instead reportedcudaErrorNoDevice. Explain why these are two different, both genuinely honest answers, rather than one of them being a bug in this book's own reporting.
Where We Go Next¶
Chapter 11's own Arena allocator, Appendix G's Rule of Five, and Appendix C's memory-hierarchy walk each assumed the stack/heap distinction and the new/malloc distinction this appendix now makes explicit; a reader arriving at this appendix out of order can safely treat it as the missing first chapter those three were each quietly building on. The pinned-versus-pageable distinction in Section J.4 sets up directly for any future chapter or appendix that streams data to the GPU repeatedly on a hot path (Chapter 22's own bond-portfolio Monte Carlo pricing is the book's own clearest candidate) -- exactly the situation where choosing pinned memory for a reused transfer buffer, rather than defaulting to ordinary malloc()/new, is the difference between a transfer that genuinely overlaps with computation and one that silently never does.
Worked Solutions¶
1. ASLR relocates the stack's base address by adding one random offset, chosen once per process launch, to every address in the entire stack region uniformly. A relationship like "&local_in_b is less than caller_frame_addr" survives this shift because both sides of the comparison are shifted by the identical amount -- if X < Y before adding some constant k to both, then X + k < Y + k afterward, for any k. A raw address, by contrast, IS one specific value that the random offset changes directly on every run, which is exactly why ASLR exists: to make a raw address useless as a stable target across runs, even though the relative layout of one run's own stack remains internally consistent throughout that single run.
2. The instant the function returns, its entire stack frame -- including local's own storage -- is considered free for the next function call to reuse; nothing erases the old bytes, but nothing protects them either. p still holds that address, but the object it used to point to no longer has a valid lifetime; this is a dangling pointer, and dereferencing it is undefined behavior. It might not crash immediately (or ever) because the "freed" stack region isn't reliably reused right away -- if no intervening function call happens to overwrite those exact bytes before p is dereferenced, the old value can still be sitting there, unchanged, making the bug appear to "work" on a given run, a given compiler, or a given optimization level, right up until some innocuous change (a different caller, a different build flag, a different day) causes a different function to reuse that same stack memory first.
3. malloc() performed no initialization at all -- it is documented, explicitly, to return uninitialized memory, containing whatever bytes happened to already occupy that region of the heap (commonly, but not guaranteedly, leftover data from some previous allocation that used and then freed the same physical memory). Sentinel's constructor is the only code in this file that ever sets magic to 0xABCD1234, and malloc() never called it -- so magic reads back as whatever those pre-existing bytes happened to already be, which this run showed does not equal the constructor's own value, precisely because no constructor ever ran to make it equal.
4. The three matched pairs are new T with delete, new T[n] with delete[], and malloc()/calloc()/realloc()/aligned_alloc() with free(). Internally, new[] records extra bookkeeping alongside the allocated array -- typically the element count, so that delete[] knows how many destructors to run before releasing the block -- that plain delete has no knowledge of and does not read. Freeing a new[]'d block with plain delete reads (or fails to read) that hidden bookkeeping incorrectly, which is exactly the alloc-dealloc-mismatch AddressSanitizer's own real, genuine report in Worked Example J.2.3 names directly: operator new [] on the allocation side, operator delete (without brackets) on the deallocation side.
5. Ordinary operator new(size_t) genuinely calls into the allocator to reserve a fresh block of memory and returns its address. Placement new's own operator new(size_t, void* p) overload does nothing of the kind -- it simply returns p, the address the caller already supplied, unchanged; no allocator bookkeeping is ever created for that address in the first place. Calling delete on a placement-new-constructed pointer invokes operator delete, which assumes it's being handed an address that ordinary operator new allocated and attempts to release it back to the allocator -- for a stack buffer, this means asking the heap allocator to free memory it never owned at all (memory corruption, not a no-op); for a pool slot, it's the exact alloc-dealloc-mismatch-shaped bug Section J.2 already demonstrated, applied to an address that was never a separate allocation to begin with.
6. cudaMemcpyAsync()'s own C function signature takes a plain void* (or typed pointer) for its source and destination arguments -- there is no separate parameter, return type, or overload that could encode "this pointer is pinned" versus "this pointer is pageable" at the type level, so the compiler accepts either with identical code and produces no warning, no error, and no difference in the call site's own appearance either way. Because the pageable-source call still runs correctly and still produces the right numbers (just without the intended overlap), the bug produces no failing test, no crash, and no visibly wrong output during ordinary testing -- it only shows up as a performance regression, and only when someone is specifically measuring for overlap between transfer and computation, which is precisely why Section J.4 calls it easy to introduce and easy to miss.
7. Both errors are this exact machine's own honest report of its own real hardware state, and the two machines that produced them were never claimed to be identical. cudaErrorInsufficientDriver means the CUDA Runtime detected some NVIDIA driver version present but too old (or, as on this appendix's own machine, detected none installed at all) for the CUDA Runtime version it was compiled against. cudaErrorNoDevice means a compatible driver WAS found, but no CUDA-capable device was enumerated behind it. Appendix C's own environment apparently had some form of NVIDIA driver stack present (however minimal) with no actual GPU attached; this appendix's own environment has no driver installed at all. Both are genuine, honestly reported answers to "is a real, usable CUDA device available," from two real but differently configured no-GPU machines -- neither is a bug, and the difference itself is a small, honest illustration of exactly how specific and how real these error codes are, rather than a single generic "no GPU" placeholder both machines happened to share.