GPU Troubleshooting

The basics of GPU code are explained in Chapter 9. This page collects what to check when a problem occurs during building or running. Source paths are relative to the folder where you cloned MISO. The error messages were checked with nvcc from CUDA 12.9 and an RTX 4090, and may differ slightly between versions.

1. Stops during CMake configuration

Symptom	What to check
`No CMAKE_CUDA_COMPILER could be found`	Whether `nvcc` is on the PATH (`nvcc --version`). If it is not, specify it with `-DCMAKE_CUDA_COMPILER=/path/to/nvcc`
The settings are not applied after switching from the CPU version	Whether you are reusing the same build folder. Use separate folders for the CPU and the GPU, or delete the folder and recreate it
`CMAKE_CUDA_ARCHITECTURES` takes a different value even though you specified it	Whether, in your own `CMakeLists.txt`, you set this value before `enable_language(CUDA)` (Chapter 9, Section 4.4)
An error that no GPU is found with `native`	The machine you build on has no GPU (for example, a login node). Specify the value of the GPU you will use explicitly (such as `89`)

2. It does not compile

Error message (excerpt)	Cause and remedy
`__host__ or __device__ annotation on lambda requires --extended-lambda nvcc flag`	`MISO_LAMBDA` requires `nvcc`’s `--expt-extended-lambda` (the same as `--extended-lambda`). Check whether you link `MISO::MISO` and whether you turned MISO’s `MISO_USE_CUDA` ON
`The enclosing parent function ("work") for an extended __host__ __device__ lambda cannot have private or protected access within its class`	You wrote `MISO_LAMBDA` inside a `private` or `protected` member function. Move it to a `public` function or to a function outside the class
`... cannot be referenced -- it is a deleted function` (copying an `Array1D` or similar)	You use an owning array (such as `Array1D`) inside a lambda. Obtain the view with `view()` before the lambda and use the view
`calling a __host__ function("...") from a __host__ __device__ function("...") is not allowed` (warning #20011)	You call a function without `__device__` from a lambda or kernel. Do not leave it as it is even if it is only a warning. Add `__host__ __device__` to your own functions, and use `util::sqrt` and the like for math functions (`miso/include/miso/utility.hpp`)
`declaration is incompatible with previous "s"`	You declare `extern __shared__ T s[];` with a type in a templated kernel. As in MISO’s `reduce_shared_memory`, declare it once as a byte array and reinterpret it (the column in Chapter 9, MISO issue #190)
g++ reports `'__global__' does not name a type`	Whether you wrote the kernel outside `#ifdef __CUDACC__`, or you are not passing the `.cpp` to `nvcc`. Check CMake’s `set_source_files_properties(... LANGUAGE CUDA)`

Extended lambdas (lambdas marked __host__ __device__) have constraints on where they can be placed and how they capture, beyond those listed here. Check the section on extended lambdas in the NVIDIA CUDA C++ Programming Guide (https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html) for the version you use.

3. It stops at run time

3.1 Common errors

Message	Possible causes
`CUDA Error: an illegal memory access was encountered`	The lambda captures `this` (it uses a member variable directly). You pass a view of a host array to the CUDA version of `for_each`. An index is outside the array range
`no kernel image is available for execution on the device`	`CMAKE_CUDA_ARCHITECTURES` (`nvcc -arch`) points to a newer generation than the GPU you use. Rebuild to match the GPU’s compute capability (`nvidia-smi --query-gpu=compute_cap --format=csv`)
`CUDA driver version is insufficient for CUDA runtime version`	The CUDA Toolkit is newer than the driver. Use a CUDA Toolkit of a version the driver supports (the upper right of the `nvidia-smi` output shows the CUDA version the driver supports)
`No CUDA-capable devices were detected.`	MISO’s `cuda::Env` cannot find a GPU. Check with `nvidia-smi` that the GPU is visible, and that you are not hiding it with the environment variable `CUDA_VISIBLE_DEVICES`

Among illegal memory access errors, the hardest to find is capturing this. It works correctly in a CPU build, and it also happens to work in some GPU environments where HMM is enabled (MISO issue #170). Go back to the declaration and check whether each name used inside the lambda is a local variable or an argument (Chapter 9, Section 5).

3.2 Locating the error

Kernel launches are asynchronous, so the file and line printed by MISO_CUDA_CHECK are where the error was noticed. The kernel in which the error occurred is one launched earlier (Chapter 9, Section 7).

If you run with the environment variable CUDA_LAUNCH_BLOCKING=1, a kernel launch waits until the kernel finishes, and the error is reported by MISO_CUDA_CHECK(cudaGetLastError()) right after the launch. In the example of Chapter 9, Section 5.1, the output changed as follows.

# run as is
CUDA Error: an illegal memory access was encountered .../miso/include/miso/array1d.hpp 109
# run with CUDA_LAUNCH_BLOCKING=1
CUDA Error: an illegal memory access was encountered .../miso/include/miso/execution.hpp 91

array1d.hpp is inside a later copy_from, and execution.hpp is the check right after the launch of the CUDA version of for_each. To find further where for_each was called, temporarily add MISO_CUDA_CHECK(cudaDeviceSynchronize()); right after a for_each in your own code to narrow down the range.

If you run with compute-sanitizer (included in the CUDA Toolkit), it shows the kernel and thread that made the invalid read or write.

compute-sanitizer ./this_capture
mpiexec -np 2 compute-sanitizer ./build/mhd2d_vortex --config=config.yaml   # with MPI

This is the output (excerpt) for the example of Chapter 9, Section 5.1.

========= Invalid __global__ read of size 4 bytes
=========     at void miso::for_each_kernel<Scaler::apply_wrong(miso::Array1DView<float>) const::[lambda(int) (instance 1)]>(miso::Range1D, T1)+0xd0
=========     by thread (0,0,0) in block (0,0,0)
=========     Access to 0x7fff1107e0a4 is out of bounds

You can see that for_each_kernel, which runs the lambda inside apply_wrong, tried to read 4 bytes (the float member factor). An address starting with 0x7fff... is commonly used on Linux for the host stack (where local variables are placed), which is a clue that it tried to read a host object. Execution is much slower under compute-sanitizer, so try it with a smaller grid and fewer time steps.

4. When combined with MPI

Are you using CUDA-aware MPI? The CUDA version of MISO’s halo exchange passes the addresses of device buffers directly to MPI_Isend and MPI_Irecv (miso/include/miso/mhd_halo_exchange.hpp). With an MPI that does not support CUDA, the program hangs inside the communication or receives wrong values. Use a CUDA-aware MPI, such as the Open MPI included in the NVIDIA HPC SDK.
Mapping between ranks and GPUs. MISO selects the GPU by the remainder of the rank number within the node divided by the number of GPUs (cuda::Env in miso/include/miso/env.hpp). If there are more ranks than GPUs, multiple ranks share one GPU.
Order of waiting for completion. Before passing a buffer written by a GPU kernel to MPI, wait with cudaDeviceSynchronize. Before using the received values in a kernel, wait with MPI_Waitall. The two waits are different things.
When running inside a container. MISO’s demos load the environment variables in demo/shared/openmpi_config.env before running mpiexec.

5. When the numbers differ between the CPU and the GPU

First, make the input (config.yaml), the grid, the MPI decomposition, and the Real type (float/double) the same. Then examine the size and location of the difference.

How the difference looks	What to check
The last few digits of reduced values such as sums and means	A difference in the order of addition (the Host version goes in order from the front; the CUDA version uses a tournament within a block). It is not a mistake
The result of `reduce` differs greatly	Whether `init` is the identity element of `op` (0 for a sum, a sufficiently large value for a minimum). In the CUDA version, threads outside the range add `init` to the reduction (Chapter 9, Section 6.4)
Values at each grid point differ by about the rounding error	`nvcc` may merge a multiplication and an addition into one instruction (FMA), which changes the rounding from the CPU. The implementations of math functions also differ between the CPU and the GPU
It differs only near the boundary	The boundary condition, the update of the ghost cells, the halo exchange
It drifts along a particular axis	The order of indices, the range of `Range3D`, the array sizes
It differs greatly from the first step	Whether the initial condition has been copied to the device (the direction of `copy_from`), and the capturing of `this`

For how to compare and how to decide the tolerance, see Testing and Comparing Results. We do not assume that the CPU and GPU results match bit for bit.