Checkpoint: library or kernel?

By the end of this checkpoint, you will choose a framework op, a CUDA library call, or a custom kernel for a measured workload. Your decision must name the operation, the cost that controls it, and the evidence that would reverse your choice.

Start with a prediction

Day 39 measured cub::DeviceReduce::Sum at 0.254 ms. The course's hand-written reduction took 0.349 ms on the same input.

Day 40 measured its hand-written PageRank iteration at 1.7963 ms for 50 iterations. A Thrust and CUB version took 9.6324 ms.

Before reading on, write one sentence that explains why the library won the first comparison and lost the second. Do not use "generic code is slow" or "libraries are faster" as your reason.

Choose the owner before the implementation

A framework op is an operation already supplied by the program that owns the workload, such as torch.matmul inside a PyTorch model. It keeps the framework's shape checks, dispatch, streams, and automatic differentiation.

A CUDA library call names an operation such as GEMM, scan, reduction, or sparse matrix-vector multiplication. The library owns the kernel and its tuning policy.

A custom kernel gives you control over work that the other two choices cannot express. You also own its tests, architecture support, tuning, and future fixes.

Use this order:

1. Critical path: does this work take enough time to affect the result you care about?
2. Framework coverage: does the current framework already provide the exact operation and required gradient?
3. Library coverage: does a CUDA library compute the exact operation?
4. Data cost: does the call need a layout change, temporary buffer, or extra pass through global memory?
5. Neighbours: would the call split work that a custom kernel can fuse?
6. Ownership: does the measured gain pay for the tests and tuning you will maintain?

Stop if the work is off the critical path. Otherwise, prefer the first framework or library path that matches the operation without a costly layout change or split.

Write a custom kernel when no call expresses the work, or when a measured fusion gain pays for ownership. A stopwatch belongs near the end because timing a different operation cannot decide the choice.

Worked decision: matrix multiplication

Take day 44's FP32 matrix multiplication at N = 2048.

First, the multiplication is on the critical path. If this code runs inside PyTorch, use torch.matmul unless its contract prevents the required fusion or layout.

In a C++ program, cublasGemmEx computes the exact operation. Row-major storage needs an operand swap, not a transpose, so the layout change costs no device work.

The GEMM is standalone in this case. The hand-written kernel took 4.219 ms while cuBLAS took 2.858 ms, so the hand-written path reached 67.7 percent of the library's speed.

The decision is to call cuBLAS. A custom kernel would be slower and would add code that must be tuned again for each target architecture.

Check the rule

Now add a bias and ReLU after the GEMM. If the framework or library can fuse that epilogue, keep the higher-level call.

If it cannot, compare the fused custom path with the full unfused chain, including both extra reads and writes. Comparing only the GEMM kernels hides the cost that caused you to write the fused kernel.

// two calls: the middle value is written to global memory and read back
thrust::gather(policy, colIdx, colIdx + nnz, contrib, gathered);
cub::DeviceSegmentedReduce::Sum(temp, bytes, gathered, y, n, off, off + 1);

// one kernel: the middle value never leaves a register
sum += contrib[colIdx[e]];

Ten diagnostic questions

Answer all ten before opening the key. Each question tests a wrong rule that can lead to the wrong owner.

1. Day 39: generic algorithms

Why did CUB beat the course's reduction and scan?

• A. CUB uses architecture-specific tuning policies for operations it implements exactly.
• B. Any library kernel is faster than a custom kernel.
• C. CUB skipped the correctness checks used by the hand-written versions.

2. Day 39: CUB and Thrust

How do CUB and Thrust relate?

• A. They are two unrelated libraries, so you must choose one for a project.
• B. Thrust provides higher-level algorithms and can use CUB for CUDA work.
• C. CUB runs only inside kernels, while Thrust runs only on the host.

3. Day 81: tensor cores in cuBLAS

Does cublasGemmEx always use tensor cores?

• A. Yes, the GemmEx name requests tensor cores.
• B. No. The data type, compute type, math mode, and GPU decide which path is legal.
• C. Yes, but only for row-major matrices.

4. Day 82: random numbers

Which cuRAND setup gives each thread an independent part of one reproducible stream?

• A. Use the same seed and the thread index as the sequence number.
• B. Give every thread a different seed and leave the sequence at zero.
• C. Call the host generator once from each thread.

5. Day 84: what CUTLASS supplies

What do you do with a CUTLASS GEMM?

• A. Call one fixed runtime function that chooses every tile.
• B. Instantiate templates that name the operation, tile shapes, and architecture.
• C. Write PTX for each tensor-core instruction.

6. Day 85: CuPy and Numba

What is the main compiler difference between CuPy RawKernel and Numba CUDA?

• A. CuPy compiles CUDA source with NVRTC; Numba compiles a Python kernel through its CUDA compiler path.
• B. Both send the same CUDA source to nvcc during package installation.
• C. Neither compiles code on the first run.

7. Day 86: Triton support

What can a learner do when the available NVIDIA GPU has compute capability 7.5?

• A. Run the Triton GPU kernel because the CUDA toolkit supports that GPU.
• B. Use TRITON_INTERPRET=1 to check correctness on the CPU, but do not report a GPU time.
• C. Lower num_warps until the kernel becomes compatible.

8. Day 87: a PyTorch wrapper

What changes when you register a CUDA kernel as a PyTorch custom op?

• A. Registration makes the kernel faster.
• B. Registration supplies framework contracts such as dispatch, fake-tensor support, and autograd only when you implement them.
• C. Registration replaces the kernel with a built-in PyTorch op.

9. Day 88: reading a profile

Which profile row should you replace first?

• A. The row with the highest call count.
• B. The first row printed by the profiler.
• C. A critical-path row whose full operation can improve, including any fusion it enables.

10. Day 89: CUDA or Triton

What decides between CUDA and Triton for a new kernel?

• A. A fixed ranking that applies to every workload.
• B. Hardware support, existing library coverage, the operation's shape, and how much control you need.
• C. The language with fewer source lines.

Answer key and explanations

1. A. CUB has tuned policies for the operation and target architecture. B fails on the PageRank case, while C is false because both paths passed their checks.

2. B. Thrust is the higher-level layer and may dispatch CUDA work through CUB. A invents a project-wide choice, and C ignores CUB's device-wide and block-wide APIs.

3. B. The function name does not choose an instruction path by itself. A omits the compute settings, and row-major storage in C changes operand order rather than tensor-core support.

4. A. cuRAND documents parallel streams through the sequence number. B can create related streams with no spacing guarantee, while the host generator cannot run as a per-thread device call.

5. B. CUTLASS is a template library that produces a kernel from the choices you instantiate. A describes a conventional runtime library call, and C skips the abstractions CUTLASS provides.

6. A. CuPy RawKernel starts from CUDA source, while Numba starts from a Python function. B and C contradict the compile paths measured on day 85.

7. B. The interpreter checks results on the CPU and supplies no GPU timing. Toolkit support does not override Triton's compute capability floor, and num_warps cannot change that floor.

8. B. The wrapper connects a kernel to PyTorch, but you must register each contract. A confuses integration with speed, and C describes replacing the custom op rather than registering it.

9. C. Total effect on the workload matters more than call count or table order. A frequent row may take little time, and B gives display order meaning it does not have.

10. B. The decision depends on the target and operation. A ignores those constraints, while C treats source length as a performance and maintenance measure.

Route each miss

Do not turn the ten answers into a score. Use each miss to choose what to repeat.

Missed	Return to
1 or 2	Day 39, "Two beliefs that both send you the wrong way"
3	Day 81, "Four paths, and the bias with its own price tag"
4	Day 82, cuRAND pitfalls
5	Day 84, "Two programs: the algebra, then the GEMM"
6	Day 85, "Where each library compiles, and when"
7	Day 86, run requirements
8	Day 87, "The op is not the kernel"
9	Day 88, "Choose a row you can improve"
10	Day 89, the decision exercise

Mystery decision artifact

This checkpoint ships an audit transcript rather than an Nsight Systems timeline. Use it to practise ownership decisions, not timeline reading.

The transcript gives you five candidates from earlier runs. They do not form one measured pipeline; decide who should own each separate stage.

Choose framework op, CUDA library, custom kernel, or leave alone.

Stage	Evidence you have
A, reduction	hand-written 0.349 ms; CUB 0.254 ms; same operation
B, exclusive scan	hand-written 1.331 ms; CUB 0.571 ms; same operation
C, Sobel plus threshold	fused 1.323 ms; split 2.565 ms; no one call covers both stages
D, PageRank SpMV	hand-written chain 1.7963 ms; Thrust plus CUB 9.6324 ms; the library form materialises a gather
E, frame pipeline	the full device chain takes 0.715 ms inside a 16.667 ms frame budget

For each stage, fill four fields:

1. Who should own it?
2. Which contract or number decides the choice?
3. What would you change first?
4. What result would make you reverse the choice?

The word framework is valid only if you name the framework op and show that its contract matches. No stage in the supplied artifact proves that case, so choosing it here requires evidence from the surrounding application that the transcript does not contain.

Price a split before timing it

Day 48 gives a lower bound for a split chain. Divide the extra bytes by the measured memory bandwidth:

// The toll a library call pays when it cannot fuse with its neighbours:
// the extra bytes the split pushes through global memory, divided by the
// rate this card copies at. It is a prediction and not a measurement,
// and printing it beside the observed difference is the whole point,
// because where it is wrong it is wrong in one direction and for one
// reason.
static double predictedTollMs(size_t extraBytes, double copyGbs) {
    return static_cast<double>(extraBytes) / (copyGbs * 1.0e6);
}

This model predicted day 48's 1.6340 ms toll as 1.6448 ms. It under-predicted the Sobel split by 2.19 times and the PageRank split by 14.02 times because those stages did not already run at copy bandwidth.

Bytes give a floor, not a complete time model. A large miss tells you to inspect access pattern, useful lane work, and the exact operation each call performs.

Run the audit

The full program is library_or_kernel.cu. It launches no kernel and uses the measurements already stored in the repository.

The verdict function records the same decision order in code:

// The four gates, in the order that ends the audit soonest.
//
// Gate 0 is not one of the four. It is day 50's first checklist
// question, and an op that owns none of the run time is not worth
// replacing whatever the other answers are.
//
// Gate 2, layout, decides nothing here on purpose. A layout mismatch is
// a price, not a veto: cuBLAS wanting column-major costs a swap of the
// operands, not a rewrite. It becomes a veto only when the conversion
// costs more than the op, which is a COO matrix rebuilt into CSR on
// every call rather than once.
static Verdict verdict(const Candidate& c) {
    if (!c.onCriticalPath) {
        return kOffPath;  // gate 0, day 50 step 1
    }
    if (c.semantics == kNoCall) {
        return kKeepKernel;  // gate 1: nothing computes what you compute
    }
    if (!c.gapMeasured) {
        return kCallLibrary;  // gates 2 and 3 only price it
    }
    if (c.handWinRatio >= kWorthOwning) {
        return kKeepKernel;  // gate 4: the win pays for the maintenance
    }
    return kCallLibrary;
}

nvcc -std=c++17 -O3 -arch=sm_75 -o library_or_kernel library_or_kernel.cu
./library_or_kernel

The program should finish with these exact lines:

• 6 call the library, 2 keep the kernel, 1 off the path
• gate 4's threshold is 1.20x; the nearest measured gap sits 39 percent away
• every check passed

Exit code zero proves that the stored arithmetic and decision checks agree. It does not prove that an unmeasured library call is faster.

Independent gate

Submit the four fields for stages A through E before opening the key. Pass when all five owners match the evidence, every row cites a measured number or operation contract, and every reversal test could change the decision.

Then run the audit. Your transcript must end with the three lines above, and all sixteen recomputed rates must match their stored rates within 0.5 percent.

Decision key

A, CUDA library. Call cub::DeviceReduce::Sum; 0.254 ms beats 0.349 ms for the same work. Reverse the choice only if a surrounding fusion changes the operation or a same-work measurement changes the result.

B, CUDA library. Call cub::DeviceScan::ExclusiveSum; 0.571 ms beats 1.331 ms. The same reversal tests apply.

C, custom kernel. Keep Sobel and threshold fused because no one NPP call covers both stages and the split took 2.565 ms against 1.323 ms. A library or framework op that expresses the fused operation, then matches or beats 1.323 ms under the same checks, would reverse the choice.

D, custom kernel. Keep the PageRank chain because the measured Thrust and CUB composition writes a gathered array and took 9.6324 ms against 1.7963 ms. A measured cusparseSpMV row for the same CSR operation could reverse the choice, but that row does not exist yet.

E, leave alone. The complete device chain uses 0.715 ms of a 16.667 ms frame budget, so replacing one stage cannot change the sustained frame rate in the measured workload. Reconsider it if a profile shows that the device chain has moved onto the critical path.

Results

The audit ran on 2026-09-02 with CUDA 12.6, driver 580.173.02, and a Tesla T4 visible. It also ran with CUDA_VISIBLE_DEVICES=""; both runs printed the same rows and verdicts because the program launches no kernel.

The source was rebuilt with CUDA 13.0 V13.0.88 and reproduced all sixteen rate checks, three toll ratios, nine verdicts, and exit code zero. The three committed transcripts are linked from the lesson front matter and the Day 90 code README.

Check	Result
Recomputed rates	16 of 16 passed
Observed/predicted toll	0.99, 2.19, 14.02
Audit verdicts	6 library, 2 custom, 1 off path
Nearest measured gap to 1.20x threshold	39 percent

The audit transcript supports the decision exercise, but it cannot teach timeline reading without a five-stage .nsys-rep.

Sources

• cuBLAS cublasGemmEx reference, for data type, compute type, and algorithm controls (checked 2026-09-01)
• cuSPARSE generic API, for the cusparseSpMV operation and CSR inputs (checked 2026-09-01)
• NVIDIA Performance Primitives introduction, for separate filtering and thresholding operations (checked 2026-09-01)
• Triton compatibility, for its operating-system and GPU requirements (checked 2026-09-01)
• PyTorch custom C++ and CUDA operators, for registration, fake tensors, and autograd (checked 2026-09-01)

Next

Day 91 applies the same choice to work split across two GPUs. Libraries cover the collectives, but you still have to prove that the collective sits on the critical path and matches the data flow you need.