100 DAYS OF GPU Course preview
Day 91Module 10
draft

Using more than one GPU

Kaggle hands you two T4s and nvidia-smi lists both. Your program uses one of them, and it will go on using one of them however big the array gets, because a CUDA program picks a device once and then stops asking. The repair looks like two calls: count the devices, set the device. Both are real and they are not enough. Which device is current when you create a stream, an event or an allocation decides which device owns it for life, and the card next to yours is a separate machine with a bus in between. By the end of this page you will have split one vector add across both cards, printed each card's own time rather than one number for the pair, and measured what it costs to move a buffer from one card to the other.

Hardware. This day needs two GPUs and the project's verification node has one Tesla T4, so the two-GPU numbers cannot come from it. The free two-GPU tier is Kaggle's T4 x2, and it is the only one (accelerator list at https://www.kaggle.com/docs/notebooks , checked 2026-08-29). The program counts the devices and runs a single-GPU fallback on one card, printing which path it took, so a Colab T4 still exercises the code.

Two devices, two of everything

The API surface people expect is the whole API surface: cudaGetDeviceCount tells you how many devices the process can see, and cudaSetDevice picks one. "A host thread can set the device it is currently operating on at any time by calling cudaSetDevice()", and until it does, the current device is device 0 (CUDA C++ Programming Guide, "Programming Systems with Multiple GPUs", https://docs.nvidia.com/cuda/cuda-programming-guide/03-advanced/multi-gpu-systems.html , checked 2026-09-01). Every program in this course so far has been using that default without knowing it.

What the two calls hide is that "current device" is ownership, not a hint. A cudaMalloc puts the allocation on whichever device is current, a stream created under device 1 belongs to device 1, and "a kernel launch will fail if it is issued to a stream that is not associated to the current device" (same page). Events are stricter: cudaEventRecord() "will fail if the input event and input stream are associated to different devices", and cudaEventElapsedTime() fails when its two events belong to different devices. That is why the program on this page cannot print one elapsed time for the pair. It prints one per device, takes the larger, and says why that is a lower bound.

So a second GPU is not a flag. It is a second set of allocations, a second stream, a second event pair, a copy in and a copy out, and host code deciding which half of the array each set is looking at. The kernel does not change: it is handed a base pointer and a count and has no idea it is looking at half of something.

Diagram: what belongs to which device. Two columns, one per device, with the host array drawn as a single bar across the top and dashed lines showing which slice feeds which column. Column 0, device 0: elements 0 to 8,388,913, its own stream, its own event pair, its own a, b and c allocations. Caption "8,388,914 elements, one shard." Column 1, device 1: elements 8,388,914 to 16,777,826, a second stream, a second event pair, a second set of allocations. Caption "8,388,913 elements, the shorter shard." Between the columns, a single arrow labelled with the host: every cross-device byte goes up to host memory and back down. Caption "no shared address space by default." Alt text: "One host array cut into two shards of 8,388,914 and 8,388,913 elements. Each device owns its own stream, events and three allocations, and the only path between the two columns runs through host memory."

The second GPU is not a bigger GPU

The intuition that fails here comes from adding SMs. Going from 40 SMs to 80 sounds like the same move as going from a T4 to a bigger card, and for the arithmetic it is. For the memory it is not. Two GPUs do not share an address space, and the road between them is the same PCIe bus day 53 already priced: on the project's T4, pinned host-to-device copies ran at 12.2 GB/s while a kernel reading device-resident memory ran at 244.8 GB/s (code/day53-pinned/evidence/run-2026-09-01.txt). A byte that has to cross from one card to the other through host memory pays that bus twice.

Which makes the first version people write wrong in a way the compiler will not catch:

cudaSetDevice(0);
vectorAdd<<<blocks, 256>>>(d_a, d_b, d_c, half);
cudaSetDevice(1);                       // the current device moved
vectorAdd<<<blocks, 256>>>(d_a + half, d_b + half, d_c + half, half);

d_a was allocated on device 0. The second launch runs on device 1 and dereferences it. There is no second allocation and no second stream, and both launches went to a default stream, so nothing would overlap even if the pointers were right. The two-line fix changes which device runs the kernel and nothing about where the data is.

Peer access is the mechanism that makes one device's pointer usable on another. cudaDeviceCanAccessPeer answers whether the topology allows it, cudaDeviceEnablePeerAccess turns it on, and "access granted by this call is unidirectional" (https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__PEER.html , checked 2026-09-01), so the reverse direction needs its own call. The guide states the payoff plainly: with peer access enabled, "peer-to-peer memory copies between these two devices no longer need to be staged through the host and are therefore faster". Without it the copy still works, which is the part that catches people out; it just takes the long road. And on the free tier there is no NVLink at the end of that road either: the T4 datasheet lists its system interface as "x16 PCIe Gen3" and names no NVLink connector (https://www.nvidia.com/content/dam/en-zz/Solutions/Data-Center/tesla-t4/t4-tensor-core-datasheet.pdf , checked 2026-09-01).

Splitting an array that does not divide by two

Full program: code/day91-multi-gpu/multi_gpu.cu.

The kernel is day 5's vector add, untouched apart from __restrict__. The element count is 16,777,827, which is odd on purpose. An even count lets n / 2 be written twice and still add up, and that is the bug: the last element belongs to nobody and nothing notices. The split gives the remainder away one element at a time and computes each shard's start from the previous shard's end, never from a formula:

    const size_t base = n / static_cast<size_t>(used);
    const size_t rem = n % static_cast<size_t>(used);
    size_t begin = 0;
    for (int d = 0; d < used; ++d) {
        shards[d].device = d;
        shards[d].begin = begin;
        shards[d].count = base + (static_cast<size_t>(d) < rem ? 1ull : 0ull);
        begin += shards[d].count;
    }

Each device gets its own everything, created while it is current. The stream, the event pair and the three allocations are made inside a cudaSetDevice for that device and are never used under another one:

        // Everything below is created while device d is current, and
        // belongs to device d for good. A stream, an event and an
        // allocation are not portable objects: "a kernel launch will fail
        // if it is issued to a stream that is not associated to the
        // current device".
        CUDA_CHECK(cudaSetDevice(s.device));
        CUDA_CHECK(cudaStreamCreate(&s.stream));
        CUDA_CHECK(cudaEventCreate(&s.start));
        CUDA_CHECK(cudaEventCreate(&s.stop));
        CUDA_CHECK(cudaMalloc(&s.d_a, bytes));
        CUDA_CHECK(cudaMalloc(&s.d_b, bytes));
        CUDA_CHECK(cudaMalloc(&s.d_c, bytes));

The correctness check never looks at the inputs. a[i] is i % 1024 and b[i] is twice that, so the answer is 3 * (i % 1024), a whole number under 3072 that a float holds exactly. The check recomputes that from the index alone, so a shard copied back at the wrong offset fails instead of agreeing with itself, and the output starts full of a sentinel so an element nobody wrote fails too. There is no tolerance because there is nothing to absorb: one float add of two exact small integers is exact, and a tolerance could only hide a misplaced shard.

Then the peer question, asked per ordered pair because the answer is not symmetric:

        std::printf("\npeer access, as reported by cudaDeviceCanAccessPeer\n");
        int canPeer = 0;
        for (int i = 0; i < used; ++i) {
            for (int j = 0; j < used; ++j) {
                if (i == j) {
                    continue;
                }
                int can = 0;
                CUDA_CHECK(cudaDeviceCanAccessPeer(&can, i, j));
                std::printf("  device %d -> device %d: %s\n", i, j,
                            can ? "yes" : "no");
                if (i == 0 && j == 1) {
                    canPeer = can;
                }
            }
        }

The program then moves 32 MiB from device 0 to device 1 four ways: down to pinned host memory and back up as two separately timed legs, then cudaMemcpyPeerAsync with peer access off and, if the query said yes, on. It does not overlap copies with compute, which is day 54's job, so this is the simple schedule rather than the good one.

Results

Partially re-run, not verified as a multi-GPU lesson. The single-GPU fallback ran again on the project's Tesla T4 with driver 580.173.02 and CUDA 13.0 (V13.0.88), matched exactly and exited 0. The required Kaggle T4 x2 transcript still does not exist, so the split-scaling and peer-copy claims remain predictions. Both fallback transcripts are listed in front matter; this table is CUDA 13.0.

Path Time (ms) GB/s vs whole array
whole array, with copies 16.789 12.0 1.00
split, with copies 16.793 12.0 1.00
whole array, kernel only 0.785 256.6 1.00
split, kernel only 0.784 256.7 1.00
Route, 32 MiB from device 0 to device 1 Time (ms) GB/s
leg 1: device 0 to pinned host not run not run
leg 2: pinned host to device 1 not run not run
staged total not run not run
cudaMemcpyPeer, access off not run not run
cudaMemcpyPeer, access on not run not run

What the fallback settled, and what remains:

  1. Held for the fallback only. All 16,777,827 elements passed the exact check and the process exited 0. The two-GPU boundary check is pending.
  2. Pending. One visible device gives the split one shard, so its kernel-only ratio was 1.00 by construction, not evidence for scaling.
  3. Held for the fallback clause. With copies, split over whole was 16.793 / 16.789, which rounds to 1.00. The predicted 1.0x to 1.8x two-GPU band is pending.
  4. Pending. No device-to-device route exists with one visible GPU.
  5. Pending. Peer access and both cudaMemcpyPeer rows were skipped.

The fallback proves the one-device branch and prices its machinery at noise: 0.784 ms split against 0.785 ms whole for the kernel. It does not prove that arithmetic scales or what crossing a bus costs.

Run it yourself

Kaggle, with the accelerator set to T4 x2, is the only free way to run the two-GPU path; the free P100 option is compute capability 6.0 and an sm_75 binary will not load on it (00-CORRECTIONS.md section 9). The quota is 30 GPU hours a week and sessions run up to 12 hours (https://www.kaggle.com/docs/notebooks , checked 2026-08-29), and this program needs about a minute of that. Write the file with %%writefile, build with !nvcc -std=c++17 -O3 -arch=sm_75 -o multi_gpu multi_gpu.cu, run it, and paste the transcript. On a Colab T4 or any single card you own, the same build runs the fallback path and produces every row except the peer table. See free GPU tiers for what each one gives you.

Exercise

The split hands each device an equal share, which is right only when the devices and their links are identical. Replace it with a weighted split: shard 0 takes a fraction you choose, shard 1 takes n minus what you already gave away rather than a second formula. Then sweep the weight and find where the two per-device times meet.

Time: 30 to 45 minutes. Submit: your createShards, the per-device rows at 0.50 and at your best weight, and one sentence saying whether the balance point moved and why.

Check: the program's own gate. Every element must equal 3 * (i % 1024) exactly; a miss prints the count of wrong elements, the first bad index, and got against want, then exits nonzero. A weight that loses or duplicates an element fails there, beside a shard boundary, before any timing prints.

Hint 1

You are looking for the point where neither device is waiting for the other. The program already prints the thing you are optimizing; you do not need a new measurement, only a second column of the one it has.

Hint 2

Two shards must satisfy exactly one invariant: begin[1] + count[1] == n and begin[1] == count[0]. If you compute both counts from the weight, rounding breaks that; if you compute the second as a subtraction, it cannot. Where does the sentinel show up when you get it wrong?

Solution

Take count[0] = static_cast<size_t>(w * n) and count[1] = n - count[0], with begin[1] = count[0], so the tail is a subtraction and the exact check cannot fail for arithmetic reasons. On two identical T4s the sweep should find 0.50, and the exercise then teaches nothing about balancing and everything about measuring: the per-device rows sit within noise of each other and the weight does not matter. Run it where the two cards differ, or where the second sits behind a slower link, and the same sweep pays for itself. A split is a guess about the hardware, and the per-device times are the only thing that knows whether it was right.

Pitfalls

The second GPU sits idle while the first works. A plain cudaMemcpy returns only when the copy is done, so device 0's whole sequence finishes before device 1's is issued and the split runs slower than one card. Use cudaMemcpyAsync on each device's own stream, which needs pinned host memory to be asynchronous at all: day 53 measured a pageable cudaMemcpyAsync call occupying 99.4 percent of its own copy's completion time, which is a blocking call wearing an async name. Both cards then pull from one host bandwidth anyway, which is what a with-copies ratio near 1.0 is telling you.

The launch fails as soon as you touch the second device. A stream created under device 0 and launched into while device 1 is current is refused with cudaErrorInvalidResourceHandle, and the same rule catches cudaEventRecord and cudaEventElapsedTime across a device boundary. Keep a device's stream, events and allocations in one struct, as this program does, and the mistake gets hard to type.

One element is missing and every test passes. With an even element count, n / 2 written twice adds up; on a real input length it does not. Compute the last shard by subtraction and pre-fill the output with a sentinel, so an element nobody wrote fails.

You expected the copy to fail without peer access and it succeeded. cudaMemcpyPeer works whether or not cudaDeviceEnablePeerAccess has been called; without it the runtime stages through host memory, which is correct, slow and silent. The query and the enable call are how you learn which one you got.

Your program suddenly sees one GPU, or the wrong one. CUDA_VISIBLE_DEVICES hides devices and renumbers the survivors from zero, so CUDA_VISIBLE_DEVICES=1 makes that card device 0 inside the process. A launcher, a scheduler or a stale notebook cell will all make cudaGetDeviceCount return 1 on a two-GPU box.

Go deeper

Next

Day 92 takes the case this program dodged: work where the devices have to talk to each other. An all-reduce over a gradient means every card needs every other card's numbers, and NCCL is the library that arranges those exchanges into a ring instead of a crowd. The peer table you print today is the input to that decision, because a collective on a link that stages through the host is a different animal from one on NVLink.