Every wrapper in kernels.cu ends with CUDA_CHECK(cudaGetLastError()); return 0;.
That catches invalid launch configurations and nothing else. An illegal memory
access inside a kernel is asynchronous and only shows up at the next synchronising
call.
In practice an out of bounds permute_nd (see the expand, permute and rank>8
issues) reports itself as:
panicked at storage.rs:50: cuda: download failed
somewhere completely unrelated, with the op that actually caused it long gone.
Given how many kernels here take unvalidated indices, this makes the GPU path
really hard to debug.
Add an opt in FASTNN_CUDA_SYNC=1 that does cudaDeviceSynchronize() plus
cudaGetLastError() after each launch and names the kernel. Basically
detect_anomaly for CUDA.
Every wrapper in
kernels.cuends withCUDA_CHECK(cudaGetLastError()); return 0;.That catches invalid launch configurations and nothing else. An illegal memory
access inside a kernel is asynchronous and only shows up at the next synchronising
call.
In practice an out of bounds
permute_nd(see theexpand,permuteand rank>8issues) reports itself as:
somewhere completely unrelated, with the op that actually caused it long gone.
Given how many kernels here take unvalidated indices, this makes the GPU path
really hard to debug.
Add an opt in
FASTNN_CUDA_SYNC=1that doescudaDeviceSynchronize()pluscudaGetLastError()after each launch and names the kernel. Basicallydetect_anomalyfor CUDA.