Merge pull request #1953 from glotzerlab/compute-sanitizer

Run `compute-sanitizer` on CI checks

.github/workflows/build_and_test.yaml

@@ -282,6 +282,21 @@ jobs:
       env:
         _HOOMD_DISALLOW_CUPY_: 1
 
+    ### compute-sanitizer checks: These checks are slow, so run them only when validation is
+    # requested, but no on the actual long validation tests. Every class's basic test should ensure
+    # that the GPU kernel is called for at least a few timesteps for these checks to be effective.
+    - name: Run memcheck
+      if: ${{ contains(inputs.config, 'cuda') && contains(github.event.pull_request.labels.*.name, 'validate') && inputs.validate == 'true' }}
+      run: compute-sanitizer --tool memcheck --error-exitcode 1 python3 -m pytest --pyargs hoomd -x -v -ra --durations=0 --durations-min=0.1
+
+    - name: Run racecheck
+      if: ${{ contains(inputs.config, 'cuda') && contains(github.event.pull_request.labels.*.name, 'validate') && inputs.validate == 'true' }}
+      run: compute-sanitizer --tool racecheck --error-exitcode 1 python3 -m pytest --pyargs hoomd -x -v -ra --durations=0 --durations-min=0.1
+
+    - name: Run synccheck
+      if: ${{ contains(inputs.config, 'cuda') && contains(github.event.pull_request.labels.*.name, 'validate') && inputs.validate == 'true' }}
+      run: compute-sanitizer --tool synccheck --error-exitcode 1 python3 -m pytest --pyargs hoomd -x -v -ra --durations=0 --durations-min=0.1
+
     ### Validation tests
     - name: Run pytest -m validate (serial)
       if: ${{ !contains(inputs.config, 'mpi') && contains(github.event.pull_request.labels.*.name, 'validate') && inputs.validate == 'true' }}

CHANGELOG.rst

@@ -21,6 +21,8 @@ Change Log
 * Read after write hazard in the GPU implementation of ``Dipole``, ``ALJ``, and all ``Patchy``
   potentials in ``hoomd.md.pair.ansiso``
   (`#1944 <https://github.com/glotzerlab/hoomd-blue/pull/1944>`__).
+* Read after write hazard in the GPU implementation of ``hoomd.md.mesh.conservation.Volume``
+  (`#1953 <https://github.com/glotzerlab/hoomd-blue/pull/1953>`__).
 
 *Added*
 

hoomd/md/VolumeConservationMeshForceComputeGPU.cu

@@ -164,6 +164,7 @@ __global__ void gpu_volume_reduce_partial_sum_kernel(Scalar* d_sum,
                 if (threadIdx.x < offs)
                     volume_sdata[threadIdx.x] += volume_sdata[threadIdx.x + offs];
                 offs >>= 1;
+                __syncthreads();
                 }
 
             // everybody sums up sum2K