[Python] cudaErrorIllegalAddress happens

[danielfrg]

Not 100% sure when it happens but I am hitting this when doing some benchmarks for cuda.compute here NVIDIA/cccl#7341 in the histogram/even Python implementation.

All previous run work fine but does seem to happen in the same one consistently.

Run:  [51/192] base [Device=0 T=I32 Elements=2^16 Bins=2048 Entropy=0.201]
/workspace/nvbench/detail/timestamps_kernel.cu:62: Cuda API call returned error: cudaErrorIllegalAddress: an illegal memory access was encountered
Command: 'cudaHostUnregister(&m_host_timestamps)'⏎

In general I think we need a way to ignore/skip these errors so that the run can finish regardless?

[danielfrg]

Ok, modifying my benchmark a bit i was able to get pass this but then i got:

Run:  [48/192] base [Device=0 T=F64 Elements=2^28 Bins=128 Entropy=0.201]
Pass: Cold: 2.398308ms GPU, 2.422008ms CPU, 7.21s total GPU, 7.44s total wall, 3008x
Pass: Batch: 2.318707ms GPU, 6.98s total GPU, 7.04s total wall, 3009x
Run:  [49/192] base [Device=0 T=I8 Elements=2^16 Bins=2048 Entropy=0.201]
Skip: Number of bins exceeds what SampleT can represent
Run:  [50/192] base [Device=0 T=I16 Elements=2^16 Bins=2048 Entropy=0.201]
Skip: CUDA error during warmup: cudaErrorIllegalAddress: an illegal memory access was encountered
Run:  [51/192] base [Device=0 T=I32 Elements=2^16 Bins=2048 Entropy=0.201]
Fail: Unexpected error: RuntimeError: /workspace/nvbench/cuda_stream.cuh:56: Cuda API call returned error: cudaErrorIllegalAddress: an illegal memory access was encountered
Command: 'cudaStreamCreate(&s)'

Which always fails when it changes to between those parameters: up to bins 128 is ok, change to 2048 fails. But if i remove the Bins 128, i can run the ones with Bins 2048 just fine.

The benchmark is: https://github.com/NVIDIA/cccl/blob/c0124542f2aa60f38c08f3ce6adbc4a2d7fa8eae/python/cuda_cccl/benchmarks/compute/histogram/even.py

[danielfrg]

Ok, that one was more a caching issue maybe nothing to be done here other than maybe allowing the failures to continue with the other benchmarks