AMGX 2.2.0 invalid write error with cuda-memcheck
Nobody has claimed this yet.
- Dominant language
- Cuda
- Stars
- 692
- Forks
- 197
- PR merge metrics
- No merged PRs in 30d
Description
Hello,
We recently move to AMGX 2.2.0 from 2.1.0 and we notice that some cases that usually run well where sometimes failing with this
error:
Thrust failure: after reduction step 2: cudaErrorIllegalAddress: an illegal memory access was encountered
File and line number are not available for this exception.
We check with cuda-memcheck and discover an invalid write throught a null pointer in amgx::distance2::compute_c_hat_kernel:
AMGX initialization, with GPU kernel...
AMGX version 2.2.0.132-opensource
Built on Apr 25 2022, 07:52:43
Compiled with CUDA Runtime 11.6, using CUDA driver 11.2
Cannot read file as JSON object, trying as AMGX config
AMGX: AmgX use gpu 0
AMGX: nrings = 2
========= Invalid global write of size 4
========= at 0x00001a30 in void amgx::distance2::compute_c_hat_kernel<int=8, int=256, int=128, int=32>(int, int const *, int const , int const , bool const , int const , int, bool const *, bool const *, int, bool const *, bool const *, bool const *)
========= by thread (227,0,0) in block (1,0,0)
========= Address 0x00000000 is out of bounds
Note that when there is no invalid write our solver did not converge while it did with 2.1.0
We look for uninitialized memory for both cpu and gpu parts but nothing appears.
Also the problem is not specific to a particular matrix.
running with cuda-gdb gives a similar error:
CUDA Exception: Warp Illegal Address
The exception was triggered at PC 0x1cd8fa50
Thread 1 "TestSolver.exe" received signal CUDA_EXCEPTION_14, Warp Illegal Address.
[Switching focus to CUDA kernel 0, grid 587, block (0,0,0), thread (192,0,0), device 0, sm 0, warp 6, lane 0]
0x000000001cd90788 in void amgx::distance2::compute_c_hat_kernel<8, 256, 128, 32>(int, int const*, int const*, int const*, bool const*, int const*, int*, int*, int*, int, int*, int*, int*)<<<(128,1,1),(256,1,1)>>> ()
A+
tg
Contributor guide
No contributing guide indexed for this repository
First steps
- Read the whole issue, then the project's contributing guide.
- Comment on the issue to say you are picking it up — it saves two people doing the same work.
- Fork the repository and make your change on a branch.
- Open a pull request that references the issue number.
Research direction
Start by reproducing the failure in amgx::distance2::compute_c_hat_kernel with cuda-memcheck or cuda-gdb, using the reported AMGX 2.2.0 environment. Compare the invalid write and solver convergence against AMGX 2.1.0; done means the null-pointer write and reported regression are explained and the affected cases run correctly.
Written by the indexing model from the issue text.
Assessment
- Domain
- hpc
- Issue type
- Bug
- Difficulty
- 4/5
- Estimated time
- 3-5 days
- Activity status
- Stale
- Clarity
- Needs clarification
- Newbie friendliness
- 25/100