NVIDIA / NVIDIA/AMGX

AMGX 2.2.0 invalid write error with cuda-memcheck

Open
#180 4 comments 0 reactions 0 assignees View on GitHub

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

  1. Read the whole issue, then the project's contributing guide.
  2. Comment on the issue to say you are picking it up — it saves two people doing the same work.
  3. Fork the repository and make your change on a branch.
  4. 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

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.