Memory access error using `md.pair.Table` with `md.nlist.Tree` on GPU

[mphoward]

Description

We have run into an issue where we get a memory access error when using a tabulated pair potential with a tree neighbor list on the GPU. With CUDA error checking enabled, we get an error using the reproducer below like:

Traceback (most recent call last):
  File "/home/mphoward/Documents/projects/peg_cg/setup/debug/test/reproducer.py", line 33, in <module>
    sim.run(0)
    ~~~~~~~^^^
  File "/home/mphoward/Documents/projects/peg_cg/setup/debug/env/lib/python3.14/site-packages/hoomd/simulation.py", line 561, in run
    self._cpp_sys.run(steps_int, write_at_start)
    ~~~~~~~~~~~~~~~~~^^^^^^^^^^^^^^^^^^^^^^^^^^^
RuntimeError: CUDA Error: an illegal memory access was encountered before /hoomd/md/NeighborListGPUTree.cc:442
unknown error
unknown error
unknown error
unknown error
unknown error

We don't get this error if we use md.pair.LJ instead. We also don't get this error if we use fewer points in the table, if we use md.nlist.Cell, or if we run on the CPU. This reproducer should not even have any pair interactions because there is only one particle.

Hence, I suspect something strange is going on between how the tabulated potential is being loaded into memory and how NeighborListGPUTree manages its memory, but I have no idea where to start looking to debug it!

Script

import hoomd
import numpy

# choose whether to use built-in WCA potential (True) or tabulated potential (False)
use_builtin = False

device = hoomd.device.GPU()
device.gpu_error_checking = True
sim = hoomd.Simulation(device=device)
snapshot = hoomd.Snapshot()
if snapshot.communicator.rank == 0:
    snapshot.configuration.box = [10, 10, 10, 0, 0, 0]
    snapshot.particles.types = ["A"]
    snapshot.particles.N = 1
sim.create_state_from_snapshot(snapshot)

nl = hoomd.md.nlist.Tree(buffer=0.4)
if use_builtin:
    wca = hoomd.md.pair.LJ(nlist=nl)
    wca.params[("A", "A")] = dict(epsilon=1.0, sigma=1.0)
else:
    r = numpy.linspace(1e-6, 2 ** (1 / 6), 1000, endpoint=False)
    u_lj = 4 * (numpy.power(r, -12) - numpy.power(r, -6) + 1 / 4)
    f_lj = 48 * (numpy.power(r, -13) - 0.5 * numpy.power(r, -7))
    wca = hoomd.md.pair.Table(nlist=nl)
    wca.params[("A", "A")] = dict(r_min=r[0], U=u_lj, F=f_lj)
wca.r_cut[("A", "A")] = 2 ** (1 / 6)

nve = hoomd.md.methods.ConstantVolume(filter=hoomd.filter.All())
integrator = hoomd.md.Integrator(dt=0.001, forces=[wca], methods=[nve])
sim.operations.integrator = integrator

sim.run(0)

Input files

No response

Output

Traceback (most recent call last):
  File "/home/mphoward/Documents/projects/peg_cg/setup/debug/test/reproducer.py", line 33, in <module>
    sim.run(0)
    ~~~~~~~^^^
  File "/home/mphoward/Documents/projects/peg_cg/setup/debug/env/lib/python3.14/site-packages/hoomd/simulation.py", line 561, in run
    self._cpp_sys.run(steps_int, write_at_start)
    ~~~~~~~~~~~~~~~~~^^^^^^^^^^^^^^^^^^^^^^^^^^^
RuntimeError: CUDA Error: an illegal memory access was encountered before /hoomd/md/NeighborListGPUTree.cc:442
unknown error
unknown error
unknown error
unknown error
unknown error

Expected output

Run without error.

Platform

Linux

Installation method

Conda-forge package

HOOMD-blue version

7.0.1

Python version

3.14.5

[joaander]

Does compute-sanitizer catch the error? Compile with the CUDA option -lineinfo so that compute-sanitizer can note the specific source code line that causes the illegal memory access.

[mphoward]

Sorry that it took me a while to get to doing this! Here is the initial output of compute-sanitizer:

========= Invalid __global__ read of size 4 bytes
=========     at __ldg(const unsigned int *)+0xb0 in sm_32_intrinsics.hpp:124
=========     by thread (0,0,0) in block (0,0,0)
=========     Access at 0x7ad34c5af638 is out of bounds
=========     and is 10,229,576,761 bytes after the nearest allocation at 0x7ad0eaa00000 of size 512 bytes
=========         Device Frame: hoomd::md::kernel::gpu_nlist_copy_primitives_kernel(unsigned int *, const unsigned int *, const unsigned int *, unsigned int)+0xb0 in NeighborListGPUTree.cu:304

That kernel is a simple transform:

__global__ void gpu_nlist_copy_primitives_kernel(unsigned int* d_traverse_order,
                                                 const unsigned int* d_indexes,
                                                 const unsigned int* d_primitives,
                                                 const unsigned int N)
    {                                
    // one thread per particle       
    const unsigned int idx = blockDim.x * blockIdx.x + threadIdx.x;
    if (idx >= N)                    
        return;                      
                                     
    const unsigned int primitive = d_primitives[idx];
    d_traverse_order[idx] = __ldg(d_indexes + primitive);
    }

The error looks like it is probably coming from primitive having some garbage value rather than the proper value (0, since there's only one particle in the test script).

Given that this runs fine without the Table, maybe something there is writing into the wrong memory (but compute-sanitizer isn't complaining about that...)? Another possibility could be that the memory isn't actually being initialized correctly, but it's not an issue when there is no Table for some reason.

[joaander]

If Table was writing out of bounds and corrupting data, compute-sanitizer should have caught it. I think that the use of Table leads to a specific arrangement of particles that causes Tree to fail.

I think the next step is to look at the kernel that computes d_primitives and see if we can identify why it generates nonsense values. Given that this is thread 0,0,0 in block 0,0,0, perhaps it is d_primitives[0] that is invalid - maybe the kernel that computes it is not even running?

[mphoward]

I thought something along these lines as well, but this reproducer runs without error (and passes compute-sanitizer) if I swap to a WCA (LJ) potential, and nothing changes about the particle configuration when doing that.

[mphoward]

I am going to be away for a week or so, but I agree we should check what the contents of d_primitives are next and if they are being set. Maybe we can also run a simple profiler to see a kernel timeline?

I may also experiment with skipping some of the code paths taken by table just to see if one of them is associated with the error.