How Many Operators Can Your GPU Run? Smoke-Testing a PyTorch GPU Environment with 55 Checks

Background

Previous article: ROCm on Windows Performance Troubleshooting: Why Is the Speedup Unimpressive When Running PyTorch on an RX 6650 XT?

When running self-compiled ROCm + PyTorch on Windows with an RX 6650 XT, I ran into a problem: LLM inference was indeed running on the GPU, but the speedup was only 1.7-2.0x, and GPU utilization in Task Manager was also quite low.

My first reaction was—could it be that many operations don't have a GPU path, and PyTorch quietly moved the compute back to the CPU?

To figure this out, I wrote a simple Python script. What it does can be stated in one sentence: for a bunch of common PyTorch operations, check whether the output is still on CUDA/HIP when the input is a CUDA/HIP tensor.

If an operation errors out immediately, or the output ends up on the CPU, that's worth digging into.

What the Script Does

The core logic is extremely simple:

x = torch.randn(128, 512, device="cuda")
out = some_operation(x)
print(out.device)  # 还在 cuda:0 上吗？

Combined with warmup + timing (with torch.cuda.synchronize() before and after), plus Python warning capture and error classification, and you have a complete smoke test.

But there's something that's easy to misunderstand here: the output tensor still being on the GPU does not mean the backend necessarily called a high-performance GPU kernel. PyTorch's device semantics themselves tend to keep CUDA tensor results on the same device; when a corresponding accelerator implementation is missing, the common behavior is to error out directly, rather than silently produce a CPU tensor. So the conclusion that "the output is still on cuda" is worth far less than "this proves the operator has a high-performance GPU kernel."

Which Operations Are Covered

The script tests 9 categories of common operations, roughly 55 checks:

The reason for choosing these operations is simple: they basically cover the most common compute patterns in mainstream deep learning models (including Transformers). They're sufficient, but they don't represent the performance quality of a complete workload.

How Timing Is Done

GPU operations are submitted asynchronously, so you must synchronize before and after timing:

for _ in range(3):       # warmup
    out = fn()

torch.cuda.synchronize()
t0 = time.perf_counter()
for _ in range(20):
    out = fn()
torch.cuda.synchronize()
elapsed = (time.perf_counter() - t0) / 20 * 1000

These timings can only be used as a rough reference. Some checks create random inputs inside a lambda, so the timing will mix in the cost of tensor allocation and random number generation. For serious performance analysis, you should pin the inputs separately, use a longer repeat, and combine it with a profiler.

What an Actual Run Looked Like

On an RX 6650 XT (gfx1032) + self-compiled ROCm/PyTorch environment, the result looked something like:

Summary: 54/55 returned CUDA/HIP tensors | 1 errors | 1 warnings

How to interpret this result:

• 54 checks successfully returned CUDA/HIP tensors. That's good news—it means these operations won't crash outright in the current environment.
• batch_norm errored. When MIOpen ran HIPRTC compilation it couldn't find <type_traits>; this is a missing C++ header issue, and installing the C++ toolchain from VS Build Tools will most likely fix it.
• sdpa can run, but PyTorch didn't enable memory efficient attention, so it took the math fallback. Functionally fine; performance is not optimal.

But this result cannot be written up as "98% operator coverage"—those 54 are only "the output is still on the GPU," not "proof of a high-performance GPU kernel." If you see someone taking similar numbers and saying "ROCm already has perfect support on Windows," that conclusion is overreaching.

Error Diagnosis

The script automatically classifies errors and prints suggestions right in the terminal. The categories it currently recognizes:

BatchNorm MIOpen/HIPRTC error — The most common issue. MIOpen uses HIPRTC to compile kernels at runtime, but on Windows HIPRTC often can't find C++ standard library headers. You'll see fatal error: 'type_traits' file not found in stderr. The fix is to install VS Build Tools and confirm that HIPRTC's include path is configured correctly. Note: this is not a CPU fallback; compilation fails outright.

SDPA: no flash attn — PyTorch wasn't compiled with memory efficient attention / FlashAttention enabled. SDPA takes the math fallback and still runs on GPU tensors, but for long-sequence, large-batch LLM workloads the performance gap is significant.

GPU OOM — Not enough VRAM; reduce the input size or close other GPU processes.

MIOpen runtime error — Other MIOpen errors. You can try clearing the kernel cache (rmdir /s /q %LOCALAPPDATA%\miopen), or turn on MIOpen_ENABLE_LOGGING=2 to see detailed logs.

What This Script Is Good For

• After setting up a new ROCm/PyTorch environment, quickly confirm that common operations don't error out immediately.
• After upgrading ROCm or PyTorch, run a regression check.
• As a first-layer screen before profiler-based diagnosis—first find out which operations have problems, then analyze in depth.

Not suitable for producing "operator coverage statistics" or a "final conclusion on performance bottlenecks."

If You Actually Want to Pin Down Performance Issues

Just looking at the output device isn't enough. To figure out why LLM inference is slow, I suggest also doing the following:

1. Use torch.profiler to record the time distribution of CPU and CUDA/HIP activity.
2. Time prefill and decode separately—the bottleneck of autoregressive decode is completely different from prefill.
3. Benchmark linear, RMSNorm, RoPE, SDPA, and the KV cache path separately using real LLM tensor shapes.
4. Turn on MIOpen/rocBLAS logging to confirm which library the key operations actually go through.
5. Try different batch sizes, prompt lengths, and max_new_tokens, and see how the performance curve changes.

Conclusions obtained this way are far more reliable than "checking whether the output is still on the GPU."