Debugging Score-P with PyTorch DDP: A Field Guide to CUDA Error 802 and Other Surprises

Paramita Choudhury — Mon, 29 Jun 2026 05:17:28 +0000

When I set out to instrument my multi-GPU DNABERT-2 training runs with Score-P to analyse DDP communication overhead, I expected the hard part to be understanding the traces. Instead, the hard part turned out to be getting Score-P to coexist with PyTorch's torchrun-based DDP launch mechanism at all.

This post documents every error I hit and exactly how I fixed each one - in the hope that the next person trying to trace a PyTorch DDP workload with Score-P doesn't spend two days rediscovering the same root causes.

The setup: DNABERT-2 (117M-parameter genomic transformer), PyTorch 2.1.2, Score-P 8.1 with Python bindings, a SLURM cluster with A100-SXM4-40GB GPUs, 1/4/8-GPU configurations.

A companion post — Where does the time really go in multi-GPU training? — covers what the traces actually revealed once they worked. This post is purely the war stories of getting there.

Background: two things Score-P does that fight PyTorch DDP

The re-exec mechanism. When you run python -m scorep train.py, Score-P does not simply import itself and start tracing. It sets environment variables (including LD_PRELOAD) to load its C measurement library, then re-executes the entire Python process from scratch with those variables in place. Your script effectively starts twice: once as the launcher, once as the instrumented process. Anything that happens before the re-exec - including CUDA initialisation - happens in a process that then exits.

The two-layer CUDA model. CUDA has two separate APIs:

The driver API, used by nvidia-smi, nvmlInit(), etc. - what the kernel module exposes.
The runtime API, used by cudaGetDeviceCount(), torch.cuda.is_available(), torch.zeros(1).cuda() - what PyTorch and Score-P's CUDA adapter both use.

On a freshly allocated SLURM job, the driver API can respond immediately while the runtime API is still initialising - sometimes for tens of seconds. Score-P's C CUDA adapter, loaded via LD_PRELOAD at library-load time, probes the runtime API the moment the process starts. If the runtime isn't ready yet, the probe poisons the CUDA context for that process permanently.

With that context, here are the errors.

Error 1: FP16 ValueError - a hidden CUDA Error 802

Symptom. Training crashed immediately with:

ValueError: FP16 Mixed precision training with AMP or APEX ('--fp16') can only
be used on CUDA devices.

This looked like a config error - I was clearly on a GPU node, nvidia-smi showed four A100s, yet PyTorch claimed no CUDA device.

Real cause. Score-P's C CUDA adapter was loaded via LD_PRELOAD at startup, before PyTorch initialised CUDA. The adapter called cudaGetDeviceCount() while the runtime was still in the cudaErrorSystemNotReady (Error 802) state. That left the CUDA context permanently broken for the process; the later torch.cuda.is_available() returned False, and the FP16 ValueError was just a downstream symptom.

Fix direction. The CUDA runtime must be warmed up - forced to fully initialise - before python -m scorep sets LD_PRELOAD. Once Score-P's C library is loaded, it's too late.

Error 2: the nvidia-smi check was the wrong layer

First attempt. A pre-flight check that polled nvidia-smi until it responded:

for attempt in $(seq 1 10); do
    if nvidia-smi > /dev/null 2>&1 && [ "$ngpus_visible" -eq "$ngpus_expected" ]; then
        echo "CUDA ready after $attempt attempt(s)."
        break
    fi
    sleep 3
done

Why it wasn't enough. nvidia-smi uses the driver API. A successful call only proves the kernel module is responding - it says nothing about whether cudaGetDeviceCount() would succeed. On a fresh job, nvidia-smi passes on attempt 1 while the runtime API is still cudaErrorSystemNotReady. Wrong layer.

Error 3: `assert torch.cuda.is_available()` fails immediately

Second attempt. A Python check inside the warmup loop:

python -c "import torch; assert torch.cuda.is_available(); torch.zeros(1).cuda()"

Why it failed. On a cold node, torch.cuda.is_available() can return False without raising - it just returns False silently. The assert then exits on the very first attempt, before the runtime had time to initialise.

Fix. Drop the assert. Call torch.zeros(1).cuda() directly inside try/except - let the exception be the "not ready yet" signal:

for attempt in $(seq 1 30); do
    if python -c "
import torch, sys
try:
    torch.zeros(1).cuda()
except Exception:
    sys.exit(1)
" 2>/dev/null; then
        echo "CUDA runtime ready after $attempt attempt(s)."
        break
    fi
    echo "CUDA runtime not ready (attempt $attempt/30), sleeping 10s..."
    sleep 10
done

This warmup runs without Score-P active - no LD_PRELOAD, no CUDA adapter - so it forces the runtime to initialise once. Every later process (including the Score-P-instrumented workers) then finds the runtime already warm.

Error 4: a node with a permanently broken CUDA runtime

Symptom. Even with the 30-attempt warmup (5 minutes), all attempts failed on one particular node, while nvidia-smi passed on attempt 1 every time.

Diagnosis. That node had a broken CUDA runtime install: the driver was fine, but cudaGetDeviceCount() never returned. A sysadmin problem, not an application one - and the scheduler kept landing my jobs there because it was first in the queue.

Fix. Exclude the bad node in the SLURM script:

#SBATCH --exclude=<broken_node>

After that, jobs landed on healthy nodes where both checks passed on attempt 1.

Error 5: the `scorep.user` import in DDP worker processes

Symptom. After fixing the node, 4- and 8-GPU runs still failed with Error 802 - this time in the worker processes spawned by torchrun, not the main process.

Cause. I had imported scorep.user at module level:

import scorep.user  # module-level import

When torchrun spawns one worker per GPU, each worker re-imports the module, and import scorep.user triggers Score-P's CUDA adapter init inside each freshly-spawned subprocess - before PyTorch sets up that worker's CUDA context. The parent-shell warmup does not carry into child processes.

Fix. Lazy import: defer import scorep.user until the first training_step(), by which point PyTorch has initialised CUDA for that worker:

# Module level - not imported yet
_scorep_user = None  # None = not yet tried; False = import failed

class ScorePTrainer(transformers.Trainer):
    def training_step(self, model, inputs):
        global _scorep_user
        if _scorep_user is None:
            try:
                import scorep.user
                _scorep_user = scorep.user
            except ImportError:
                _scorep_user = False
        if _scorep_user:
            _scorep_user.region_begin("dnabert_train_step")
        result = super().training_step(model, inputs)
        if _scorep_user:
            _scorep_user.region_end("dnabert_train_step")
        return result

The None guard means the import is attempted exactly once per process, on the first step - after CUDA is ready for that rank.

Error 6: Score-P memory limit exceeded

Symptom. With CUDA kernel tracing on (SCOREP_CUDA_ENABLE=kernel,memcpy,sync):

[Score-P] Warning: Too many memory requested. Score-P supports only up to,
but not including, 4 GiB of total memory per process. Reducing to its maximum value.

I had set SCOREP_TOTAL_MEMORY=4G. Score-P's hard per-process limit is strictly less than 4 GiB - exactly 4G hits the ceiling.

Fix. SCOREP_TOTAL_MEMORY=3500M - under the cap, with room for CUDA kernel traces across 8 ranks.

Error 7: `load_best_model_at_end` strategy conflict

Symptom. The short 50-step trace runs failed instantly:

ValueError: --load_best_model_at_end requires the save and eval strategy to match,
but found Evaluation strategy: NO / Save strategy: STEPS

I'd set --evaluation_strategy no to keep eval passes from distorting the trace timeline, but load_best_model_at_end=True requires matching save/eval strategies.

Fix. There is no "best model" for a 50-step diagnostic run:

--evaluation_strategy no \
--load_best_model_at_end False \
--save_steps 10000 \

Error 8: the trace contained the launcher, not the workers

Symptom. The runs completed, produced an OTF2 trace, and opened cleanly in Vampir - showing a single red bar: ...LocalElasticAgent._invoke_run. No kernels. No NCCL. No training steps. The process filter listed exactly one process.

Cause. The launcher was python -m scorep .../scorep_torchrun.py, where scorep_torchrun.py is just from torch.distributed.run import main; main() - i.e. plain torchrun. Its elastic agent spawns the GPU workers as separate child processes that start fresh python interpreters with no Score-P. Score-P therefore instrumented only the agent - the babysitter - which spends the whole run waiting. On disk the proof was unambiguous: the entire 8-GPU trace held a single process's events, and scorep-score showed only Python (USR) regions - no CUDA type:

$ ls traces/
0.def   0.evt                          # one process, not eight

$ scorep-score profile.cubex
flt  type  max_buf[B]   visits  time[s] time[%]  region
     ALL  16,266,573  625,628   38.47   100.0   ALL
     USR  16,266,302  625,627   38.04    98.9   USR     ← all Python, the agent waiting
  SCOREP        271        1    0.43     1.1   SCOREP

Fix. Stop letting torchrun spawn. Launch each rank yourself in a background loop, each as its own python -m scorep process with its own SCOREP_EXPERIMENT_DIRECTORY=scorep_rank_N, using a single-node rendezvous (MASTER_ADDR=localhost, per-rank RANK/LOCAL_RANK). This is exactly what torchrun does internally - fork N ranks, hand each its identity - except now every rank runs under Score-P. No srun, no spawn.

Error 9: `SCOREP_CUDA_ENABLE` captured zero kernels

Symptom. With per-rank launch working, the traces contained the training process - but zero CUDA kernels. scorep-score showed 98% USR (Python) regions and no CUDA type at all, despite SCOREP_CUDA_ENABLE=kernel,memcpy,sync.

Cause. The Score-P Python wrapper passes unknown flags to scorep-config, whose help is explicit: --cuda|--nocuda … On default cuda instrumentation is disabled. Setting SCOREP_CUDA_ENABLE only configures what the CUDA adapter records - but without --cuda the adapter is never loaded.

Fix. Add --cuda to the launch: python -m scorep --cuda --thread=pthread. A wrinkle: do not add the documented -- script separator - this wrapper version forwards it to scorep-config, which rejects it (Unknown option: '--'). After the fix, a CUDA type appears in scorep-score, with ~106 named GPU kernel regions per rank - and, crucially, the NCCL collectives:

$ scorep-score profile.cubex
flt  type  max_buf[B]     visits  time[s] time[%]  region
     ALL  73,352,737  3,001,818   41.63   100.0   ALL
     USR  73,351,928  2,821,228   36.83    88.5   USR
    CUDA   2,347,618     90,294    3.99     9.6   CUDA   ← GPU kernels now captured

$ scorep-score -r profile.cubex | grep CUDA | sort -k4 -rn | head
  CUDA    700 visits  2.38s  ncclKernel_AllReduce_RING_LL_Sum_float   ← gradient sync
  CUDA     62 visits  0.09s  ncclKernel_AllGather_RING_LL_Sum_int8_t
  CUDA  9,538 visits  0.09s  at::native::unrolled_elementwise_kernel<...>
  CUDA  8,376 visits  0.08s  at::native::elementwise_kernel<128, 4, ...>

Error 10: CUPTI buffer overflow at 8 ranks

Symptom. The 1- and 4-GPU traces were clean, but the 8-GPU run dropped records:

[CUPTI Activity] Dropped 85222 records. Current buffer size: 1048576 bytes
Proposed minimum SCOREP_CUDA_BUFFER=8889000

Cause. Eight ranks each profiling through CUPTI overran the default 1 MB per-process CUDA activity buffer between flushes, silently discarding kernel records.

Fix. Score-P told us the answer in the warning. SCOREP_CUDA_BUFFER=64M for generous headroom. No more dropped records.

Error 11: the NCCL watchdog vs. Score-P shutdown race

Symptom. The 8-GPU run finished training but then 6 of 8 ranks aborted during teardown:

terminate called after throwing an instance of 'c10::Error'
  what():  Should never been called   (dummyHasPrimaryContext)
  ... c10d::ProcessGroupNCCL::ncclCommWatchdog()

The aborts struck after training, killing the process before Score-P flushed its profile - so those ranks left no trace on disk.

Cause. A shutdown race: PyTorch's background NCCL watchdog thread runs its cleanup destructor (which touches the CUDA device) at interpreter exit, at the same time Score-P tears down its CUDA context. Whichever loses, crashes. It's non-deterministic - a later 4-GPU run lost the race where an 8-GPU run had won it.

Fix. Remove the race instead of fighting it: call torch.distributed.destroy_process_group() at the end of train(), so NCCL is torn down cleanly, before the interpreter (and Score-P) begin shutdown. With the process group gone, there's no watchdog destructor left to collide with Score-P's teardown, and all eight ranks flush reliably.

Final state: per-rank GPU traces collected

After all eleven fixes, every DDP rank ran under its own Score-P measurement, capturing each worker's GPU kernels and NCCL communication.

Config	Runtime	Samples/sec	Speedup
1 GPU	478.8 s	75.0	1×
4 GPU	103.8 s	345.9	4.61×
8 GPU	52.8 s	680.7	9.08×

On 1 GPU there is zero NCCL; from 4 GPUs the gradient AllReduce appears, and by 8 GPUs ncclKernel_AllReduce is the single largest GPU activity (~2.375 s, comparable to the entire backward pass) - yet it overlaps backward compute on a separate CUDA stream, which is why throughput still scales near-linearly. The Master Timeline makes the overlap visible: compute runs on the default stream CUDA[0:7] (CUDA_NULL_STREAM) while ncclKernel_AllReduce runs concurrently on a separate stream CUDA[0:20].

What that decomposition means is the subject of the companion post.

Lessons

Score-P's re-exec is not optional - design around it. Any CUDA init that must happen before Score-P's C adapter loads has to happen before the python -m scorep call in your shell script.
The driver and runtime APIs are different things. nvidia-smi passing is necessary but not sufficient. Test the runtime directly (torch.zeros(1).cuda()), and use try/except, not is_available().
Module-level imports of the Score-P user API break DDP workers. Each rank is a fresh subprocess; lazy-import inside the first method PyTorch guarantees runs after CUDA setup.
A broken node will burn your budget on CUDA timeouts. --exclude it as soon as you spot it.
Separate the diagnostic trace from the full run. CUDA-enabled --max_steps 50 with --evaluation_strategy no gives clean, size-controlled traces; the full run with CUDA off gives robust scaling stats. You need both.
To trace DDP workers, launch them yourself - don't let torchrun spawn. Replace it with a background loop where each rank is its own python -m scorep process. This is the single most important structural fix.
Setting an env var is not the same as loading the adapter. SCOREP_CUDA_ENABLE configures the CUDA adapter; --cuda loads it. Confirm a CUDA type appears in scorep-score.
Profile sums overstate communication - use the timeline for wall-clock truth. NCCL LL kernels busy-wait, and DDP overlaps AllReduce with backward compute on a separate stream, so summed kernel-time double-counts the overlap.
Tear down NCCL cleanly so it doesn't race your profiler at exit. torch.distributed.destroy_process_group() at the end of training removes the watchdog before interpreter shutdown.

The OTF2 traces behind this post were generated on a SLURM cluster (A100-SXM4-40GB) as part of a Score-P performance analysis of DDP training scaling for the DNABERT-2 genomic classifier.

DEV Community: Paramita Choudhury