Dynamic loading + NVRTC #93

rubber-duck-debug · 2025-12-11T10:37:56Z

No description provided.

Co-authored-by: Guillaume Fraux <guillaume.fraux@epfl.ch>

…partial periodicity.

Resolved conflicts in mic_neighbourlist.cu by keeping the NVRTC-compatible flat array implementation (gpulite) instead of the double3 types (main). 🤖 Generated with [Claude Code](https://claude.com/claude-code) Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>

Fixes MinGW build failure where int32_t was not declared. 🤖 Generated with [Claude Code](https://claude.com/claude-code) Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>

macos-13 runners are being retired. 🤖 Generated with [Claude Code](https://claude.com/claude-code) Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>

macos-15 is ARM64 only and cannot build x86_64 wheels. 🤖 Generated with [Claude Code](https://claude.com/claude-code) Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>

macos-13 is retired, macos-15-intel is the replacement for Intel builds. 🤖 Generated with [Claude Code](https://claude.com/claude-code) Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>

This reverts commit 7ec8311.

This reverts commit bbfbb9b.

This reverts commit 54852da.

Luthaf

There are a lot of changes in vesin/src/mic_neighbourlist.cu that seems to undo b5559d1. Is there a specific reason for it?

Luthaf · 2025-12-16T14:26:54Z

python/vesin/tests/test_neighbors.py

+
+
+@pytest.mark.skipif(
+    not (HAS_CUPY and CUDA_AVAILABLE), reason="CuPy not available or CUDA not available"
+)
+@pytest.mark.parametrize("full_list", [False, True])
+def test_cupy_large_box_small_cutoff(full_list):
+    """Test CuPy with synthetic data - large box and small cutoff"""
+    # Use synthetic data with large box to avoid CUDA cutoff <= cell/2 limitation
+    points_np = np.array(
+        [
+            [0.0, 0.0, 0.0],
+            [0.0, 2.0, 0.0],
+            [0.0, 0.0, 2.0],
+            [-6.0, 0.0, 0.0],
+            [-6.0, -2.0, 0.0],
+            [-6.0, 0.0, -2.0],
+        ],
+        dtype=np.float64,
+    )
+
+    box_np = np.array(
+        [
+            [54.0, 0.0, 0.0],
+            [0.0, 54.0, 0.0],
+            [0.0, 0.0, 54.0],
+        ],
+        dtype=np.float64,
+    )
+
+    # NumPy computation
+    calculator_cpu = vesin.NeighborList(cutoff=2.1, full_list=full_list, sorted=True)
+    i_np, j_np, d_np, S_np = calculator_cpu.compute(
+        points=points_np,
+        box=box_np,
+        periodic=True,
+        quantities="ijdS",
+    )
+
+    # CuPy computation
+    calculator_gpu = vesin.NeighborList(cutoff=2.1, full_list=full_list, sorted=True)
+    points_cp = cp.asarray(points_np, dtype=cp.float64)
+    box_cp = cp.asarray(box_np, dtype=cp.float64)
+
+    i_cp, j_cp, d_cp, S_cp = calculator_gpu.compute(
+        points=points_cp,
+        box=box_cp,
+        periodic=True,
+        quantities="ijdS",
+    )
+
+    # Verify outputs are CuPy arrays
+    assert isinstance(i_cp, cp.ndarray)
+    assert isinstance(j_cp, cp.ndarray)
+    assert isinstance(d_cp, cp.ndarray)
+    assert isinstance(S_cp, cp.ndarray)
+
+    # Verify expected pairs based on full_list
+    pairs_np = list(zip(i_np.tolist(), j_np.tolist()))
+    if full_list:
+        expected_pairs = sorted(
+            [
+                (0, 1),
+                (0, 2),
+                (1, 0),
+                (2, 0),
+                (3, 4),
+                (3, 5),
+                (4, 3),
+                (5, 3),
+            ]
+        )
+    else:
+        expected_pairs = sorted(
+            [
+                (0, 1),
+                (0, 2),
+                (3, 4),
+                (3, 5),
+            ]
+        )
+    assert sorted(pairs_np) == expected_pairs
+
+
+@pytest.mark.skipif(
+    not (HAS_CUPY and CUDA_AVAILABLE), reason="CuPy not available or CUDA not available"
+)
+def test_cupy_no_neighbors():
+    """Test CuPy when there are no neighbors"""
+    points_np = np.array([[0.0, 0.0, 0.0], [0.5, 0.5, 0.5]], dtype=np.float64)
+    box_np = np.eye(3, dtype=np.float64)
+
+    calculator = vesin.NeighborList(cutoff=0.1, full_list=True)
+
+    # NumPy
+    i_np, j_np, S_np, D_np = calculator.compute(
+        points_np, box_np, True, quantities="ijSD"
+    )
+
+    # CuPy
+    calculator_gpu = vesin.NeighborList(cutoff=0.1, full_list=True)
+    points_cp = cp.asarray(points_np)
+    box_cp = cp.asarray(box_np)
+    i_cp, j_cp, S_cp, D_cp = calculator_gpu.compute(
+        points_cp, box_cp, True, quantities="ijSD"
+    )
+
+    # Both should have no neighbors
+    assert len(i_np) == 0
+    assert len(i_cp) == 0
+    assert isinstance(i_cp, cp.ndarray)


Suggested change

@pytest.mark.skipif(

not (HAS_CUPY and CUDA_AVAILABLE), reason="CuPy not available or CUDA not available"

)

@pytest.mark.parametrize("full_list", [False, True])

def test_cupy_large_box_small_cutoff(full_list):

"""Test CuPy with synthetic data - large box and small cutoff"""

# Use synthetic data with large box to avoid CUDA cutoff <= cell/2 limitation

points_np = np.array(

[

[0.0, 0.0, 0.0],

[0.0, 2.0, 0.0],

[0.0, 0.0, 2.0],

[-6.0, 0.0, 0.0],

[-6.0, -2.0, 0.0],

[-6.0, 0.0, -2.0],

],

dtype=np.float64,

)

box_np = np.array(

[

[54.0, 0.0, 0.0],

[0.0, 54.0, 0.0],

[0.0, 0.0, 54.0],

],

dtype=np.float64,

)

# NumPy computation

calculator_cpu = vesin.NeighborList(cutoff=2.1, full_list=full_list, sorted=True)

i_np, j_np, d_np, S_np = calculator_cpu.compute(

points=points_np,

box=box_np,

periodic=True,

quantities="ijdS",

)

# CuPy computation

calculator_gpu = vesin.NeighborList(cutoff=2.1, full_list=full_list, sorted=True)

points_cp = cp.asarray(points_np, dtype=cp.float64)

box_cp = cp.asarray(box_np, dtype=cp.float64)

i_cp, j_cp, d_cp, S_cp = calculator_gpu.compute(

points=points_cp,

box=box_cp,

periodic=True,

quantities="ijdS",

)

# Verify outputs are CuPy arrays

assert isinstance(i_cp, cp.ndarray)

assert isinstance(j_cp, cp.ndarray)

assert isinstance(d_cp, cp.ndarray)

assert isinstance(S_cp, cp.ndarray)

# Verify expected pairs based on full_list

pairs_np = list(zip(i_np.tolist(), j_np.tolist()))

if full_list:

expected_pairs = sorted(

[

(0, 1),

(0, 2),

(1, 0),

(2, 0),

(3, 4),

(3, 5),

(4, 3),

(5, 3),

]

)

else:

expected_pairs = sorted(

[

(0, 1),

(0, 2),

(3, 4),

(3, 5),

]

)

assert sorted(pairs_np) == expected_pairs

@pytest.mark.skipif(

not (HAS_CUPY and CUDA_AVAILABLE), reason="CuPy not available or CUDA not available"

)

def test_cupy_no_neighbors():

"""Test CuPy when there are no neighbors"""

points_np = np.array([[0.0, 0.0, 0.0], [0.5, 0.5, 0.5]], dtype=np.float64)

box_np = np.eye(3, dtype=np.float64)

calculator = vesin.NeighborList(cutoff=0.1, full_list=True)

# NumPy

i_np, j_np, S_np, D_np = calculator.compute(

points_np, box_np, True, quantities="ijSD"

)

# CuPy

calculator_gpu = vesin.NeighborList(cutoff=0.1, full_list=True)

points_cp = cp.asarray(points_np)

box_cp = cp.asarray(box_np)

i_cp, j_cp, S_cp, D_cp = calculator_gpu.compute(

points_cp, box_cp, True, quantities="ijSD"

)

# Both should have no neighbors

assert len(i_np) == 0

assert len(i_cp) == 0

assert isinstance(i_cp, cp.ndarray)

This is already in test_cuda.py

Luthaf · 2025-12-16T14:27:05Z

python/vesin/tests/test_neighbors.py

+try:
+    import cupy as cp
+
+    HAS_CUPY = True
+    # Check if CUDA is available
+    try:
+        cp.cuda.Device(0).compute_capability
+        CUDA_AVAILABLE = True
+    except cp.cuda.runtime.CUDARuntimeError:
+        CUDA_AVAILABLE = False
+except ImportError:
+    HAS_CUPY = False
+    CUDA_AVAILABLE = False
+    cp = None
+
+


Suggested change

try:

import cupy as cp

HAS_CUPY = True

# Check if CUDA is available

try:

cp.cuda.Device(0).compute_capability

CUDA_AVAILABLE = True

except cp.cuda.runtime.CUDARuntimeError:

CUDA_AVAILABLE = False

except ImportError:

HAS_CUPY = False

CUDA_AVAILABLE = False

cp = None

Luthaf · 2025-12-16T14:27:33Z

python/vesin/vesin/_c_api.py

please undo the changes on this file, they have been integrated differently

Luthaf · 2025-12-16T14:28:20Z

python/vesin/vesin/_neighbors.py

+        # Handle single bool, numpy bool, or 0-dim array
        if is_cupy:
-            periodic = cp.asarray(periodic, dtype=cp.bool_)
+            # For CuPy arrays, handle conversion differently
+            if isinstance(periodic, cp.ndarray):
+                if periodic.ndim == 0:
+                    val = bool(periodic.get())
+                    periodic = cp.array([val, val, val], dtype=cp.bool_)
+                else:
+                    periodic = cp.asarray(periodic, dtype=cp.bool_)
+            else:
+                # periodic is a Python bool or similar
+                periodic_arr = np.asarray(periodic)
+                if periodic_arr.ndim == 0:
+                    val = bool(periodic_arr)
+                    periodic = cp.array([val, val, val], dtype=cp.bool_)
+                else:
+                    periodic = cp.asarray(periodic_arr, dtype=cp.bool_)
        else:
-            periodic = np.asarray(periodic, dtype=np.bool_)
+            periodic_arr = np.asarray(periodic)
+            if periodic_arr.ndim == 0:
+                val = bool(periodic_arr)
+                periodic = np.array([val, val, val], dtype=np.bool_)
+            else:
+                periodic = np.asarray(periodic_arr, dtype=np.bool_)


I don't see why this change is required? Is this a leftover from previous version of the PR?

Luthaf · 2025-12-16T14:28:28Z

ruff.toml

 [lint]
 select = ["E", "F", "B", "I"]
-ignore = ["B018", "B904"]
+ignore = ["B018", "B904", "B905"]


Suggested change

ignore = ["B018", "B904", "B905"]

ignore = ["B018", "B904"]

Luthaf · 2025-12-16T14:33:58Z

vesin/src/vesin_cuda.cpp

+#include "cuda_cache.hpp"
+#include "dynamic_cuda.hpp"


Should these be hidden behind a gpulite header/directory?

Suggested change

#include "cuda_cache.hpp"

#include "dynamic_cuda.hpp"

#include <gpulite/cuda_cache.hpp>

#include <gpulite/dynamic_cuda.hpp>

or even

Suggested change

#include "cuda_cache.hpp"

#include "dynamic_cuda.hpp"

#include <gpulite.hpp>

Suggested change

#include "cuda_cache.hpp"

#include "dynamic_cuda.hpp"

#include <gpulite/gpulite.hpp>

Luthaf · 2025-12-16T14:35:42Z

vesin/src/vesin_cuda.cpp

-    if (neighbors.vectors && is_device_ptr(get_ptr_attributes(neighbors.vectors), "vectors")) {
-        CUDA_CHECK(cudaFree(neighbors.vectors));
+    if (neighbors.vectors && is_device_ptr(getPtrAttributes(neighbors.vectors), "vectors")) {
+        CUDART_SAFE_CALL(CUDART_INSTANCE.cudaFree(neighbors.vectors));


random API suggestion: should this be

Suggested change

CUDART_SAFE_CALL(CUDART_INSTANCE.cudaFree(neighbors.vectors));

CUDART_CALL(cudaFree(neighbors.vectors));

(or some other name like CUDART_DYNAMIC_CALL, GPULITE_CUDART_CALL, …)?

Basically hiding the global instance from users of gpu-lite

Luthaf · 2025-12-16T14:36:21Z

vesin/src/vesin_cuda.cpp

    neighbors.shifts = nullptr;
    neighbors.distances = nullptr;
    neighbors.vectors = nullptr;
+    extras->length_ptr = nullptr;


why is this required? I would think that the default constructor below would handle it

Luthaf · 2025-12-16T14:40:15Z

vesin/src/mic_neighbourlist.cu

+// Vector math helpers using flat array indexing
+__device__ inline double dot3(const double* a, const double* b) {
+    return a[0] * b[0] + a[1] * b[1] + a[2] * b[2];


Why undo these changes?

Luthaf · 2025-12-16T14:40:38Z

vesin/src/mic_neighbourlist.cu

+    double* vector,
+    int32_t* shift,
+    const double* box,
+    const double* inv_box,
+    const bool* periodic,


Same here, why use pointers instead of richer cuda types?

rubber-duck-debug and others added 30 commits June 6, 2025 16:39

initial commit

3e69f9b

initial build system functional

5edf607

added optional return values

8b9314d

small changes

c61401d

updated stub

325179b

updates with half implementation.

d9dd50f

stub update

6e0bd8c

stub update.

4c5b07a

Update vesin/src/vesin.cpp

44e0622

Co-authored-by: Guillaume Fraux <guillaume.fraux@epfl.ch>

updates to include better cuda ptr management

3f97a6f

Merge branch 'cuda' of https://github.com/Luthaf/vesin into cuda

d97454d

added better failure wrappers for cuda calls.

1b2f4e9

added static registry to avoid cuda tear-down issues.

981f0ff

Merge branch 'main' into cuda

51bbe0e

update with opaque struct

85277a9

remove cuda ptr manager

2a1c7c7

remove kernel sync.

805d8c1

revert public API

ef36d74

fix stub for CudaNeighborListExtras

3f7d48d

removed debug prints

616f9cc

multi-gpu support

2a77469

getting CPU-only to work.

44cdfd0

formatting.

a10b496

assert_close fix.

dd76071

Merge branch 'main' into cuda

cf7f514

ctype and format

c97945c

removed device count.

f863adc

added rcut vs. cell test incl. triclinic

162ccd0

added full vectorisation back.

53dd5e4

fixed typo.

3254b56

rubber-duck-debug and others added 27 commits October 1, 2025 15:34

Merge branch 'main' into cuda

c33b2c7

added broken target include dir

a6db0cc

fixes for device kind in c++ tests.

3767572

fixed doxygen types

f2518bd

initial commit with cupy tests.

e84a75d

formatting

c7102a8

Merge branch 'main' into cupy

dceb208

fixed issues with zip(..., strict=True)

e7e098c

odd, now complains about strict

c325b72

add B905 exception to ruff

5573639

Merge remote-tracking branch 'origin/main' into gpulite

af68340

added B905 exception

50bda52

removed explicit sync

4af9abb

fixed double free

c97d60f

Merge branch 'cupy' into gpulite

0852d8d

fixes to the CUDA implementation to support non-orthogonal cells and …

8936240

…partial periodicity.

Merge branch 'cuda_fixes' into gpulite

3d0c67c

Add missing #include <cstdint> to types.hpp

b48317b

Fixes MinGW build failure where int32_t was not declared. 🤖 Generated with [Claude Code](https://claude.com/claude-code) Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>

Update macOS runner from macos-13 to macos-15

54852da

macos-13 runners are being retired. 🤖 Generated with [Claude Code](https://claude.com/claude-code) Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>

Revert macOS runner to macos-13 for x86_64 builds

bbfbb9b

macos-15 is ARM64 only and cannot build x86_64 wheels. 🤖 Generated with [Claude Code](https://claude.com/claude-code) Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>

Use macos-15-intel for x86_64 macOS builds

7ec8311

macos-13 is retired, macos-15-intel is the replacement for Intel builds. 🤖 Generated with [Claude Code](https://claude.com/claude-code) Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>

Revert "Use macos-15-intel for x86_64 macOS builds"

cfa364e

This reverts commit 7ec8311.

Revert "Revert macOS runner to macos-13 for x86_64 builds"

ce141ff

This reverts commit bbfbb9b.

Revert "Update macOS runner from macos-13 to macos-15"

3848d71

This reverts commit 54852da.

formatting.

7868939

Merge branch 'main' into gpulite

f975dce

rubber-duck-debug force-pushed the gpulite branch from e796f6c to f975dce Compare December 11, 2025 22:36

rubber-duck-debug mentioned this pull request Dec 11, 2025

CUDA Cell List Implementation #95

Open

Luthaf reviewed Dec 16, 2025

View reviewed changes

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

Dynamic loading + NVRTC #93

Dynamic loading + NVRTC #93

Uh oh!

rubber-duck-debug commented Dec 11, 2025

Uh oh!

Luthaf left a comment

Uh oh!

Luthaf Dec 16, 2025

Uh oh!

Luthaf Dec 16, 2025

Uh oh!

Luthaf Dec 16, 2025

Uh oh!

Luthaf Dec 16, 2025

Uh oh!

Luthaf Dec 16, 2025

Uh oh!

Luthaf Dec 16, 2025

Uh oh!

Luthaf Dec 16, 2025

Uh oh!

Luthaf Dec 16, 2025

Uh oh!

Luthaf Dec 16, 2025

Uh oh!

Luthaf Dec 16, 2025

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

	#include "cuda_cache.hpp"
	#include "dynamic_cuda.hpp"
	#include <gpulite.hpp>

	#include "cuda_cache.hpp"
	#include "dynamic_cuda.hpp"
	#include <gpulite/gpulite.hpp>

	CUDART_SAFE_CALL(CUDART_INSTANCE.cudaFree(neighbors.vectors));
	CUDART_CALL(cudaFree(neighbors.vectors));

Dynamic loading + NVRTC #93

Are you sure you want to change the base?

Dynamic loading + NVRTC #93

Uh oh!

Conversation

rubber-duck-debug commented Dec 11, 2025

Uh oh!

Luthaf left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants