MacOS-only version with MPS support

.claude/settings.local.json

@@ -0,0 +1,9 @@
+{
+  "permissions": {
+    "allow": [
+      "Bash(grep -v \":0$\")",
+      "Bash(ls -la /Users/florian/Documents/Science-IA/libmolgrid/test/*.cu)",
+      "Bash(wc -l /Users/florian/Documents/Science-IA/libmolgrid/src/*.cu /Users/florian/Documents/Science-IA/libmolgrid/src/*.cpp)"
+    ]
+  }
+}

CMakeLists.txt

@@ -1,5 +1,5 @@
 cmake_minimum_required (VERSION 3.18)
-project(LibMolGrid LANGUAGES C CXX CUDA)
+project(LibMolGrid LANGUAGES C CXX OBJCXX)
 
 #version number
 set (VERSION_MAJOR 0)
@@ -18,71 +18,62 @@ option(BUILD_SHARED "Build shared library" ON)
 option(BUILD_STATIC "Build static library" ON)
 option(BUILD_COVERAGE "Build with code coverage" OFF)
 
-if(${BUILD_CONTAINED_PYTHON}) 
-  #if make a mostly statically linked python module, use static boost
-   set(BUILD_STATIC ON)
-   set(Boost_USE_STATIC_LIBS ON)
+if(${BUILD_CONTAINED_PYTHON})
+  set(BUILD_STATIC ON)
+  set(Boost_USE_STATIC_LIBS ON)
 endif()
 
 list(APPEND CMAKE_MODULE_PATH "${PROJECT_SOURCE_DIR}/cmake/Modules;${PROJECT_SOURCE_DIR}/cmake")
-# get git hash
 include(cmake/git_revision.cmake)
-
-# guard against in-source builds and bad build-type strings
 include(cmake/safeguards.cmake)
-
-# place binaries and libraries according to GNU standards
 include(GNUInstallDirs)
+
 set(CMAKE_ARCHIVE_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/${CMAKE_INSTALL_LIBDIR})
 set(CMAKE_LIBRARY_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/${CMAKE_INSTALL_LIBDIR})
 set(CMAKE_RUNTIME_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/${CMAKE_INSTALL_BINDIR})
 
 #dependencies
-find_package(CUDA REQUIRED)
-find_package(Boost REQUIRED COMPONENTS regex unit_test_framework program_options system filesystem iostreams serialization)
+find_package(Boost REQUIRED COMPONENTS regex unit_test_framework program_options filesystem iostreams serialization)
 find_package(OpenBabel3 REQUIRED)
 include_directories(SYSTEM ${OPENBABEL3_INCLUDE_DIR})
 find_package(ZLIB)
 
 # configure a header file to pass some of the CMake settings
-# to the source code
 configure_file (
   "${PROJECT_SOURCE_DIR}/cmake/config.h.in"
   "${PROJECT_BINARY_DIR}/include/libmolgrid/config.h"
   )
 
 set(CMAKE_CXX_STANDARD 17)
-set(CMAKE_CUDA_STANDARD 17)
+set(CMAKE_OBJCXX_STANDARD 17)
 
-if(CMAKE_CXX_COMPILER_ID MATCHES GNU)
-    set(CMAKE_CXX_FLAGS         "${CMAKE_CXX_FLAGS} -Wall -Wno-unknown-pragmas -Werror")
+if(CMAKE_CXX_COMPILER_ID MATCHES "GNU|Clang|AppleClang")
+    # OpenBabel 3.1.1 uses std::binary_function which was removed in C++17.
+    # Restore it via the compatibility shim in libc++.
+    set(CMAKE_CXX_FLAGS         "${CMAKE_CXX_FLAGS} -Wall -Wno-unknown-pragmas -D_LIBCPP_ENABLE_CXX17_REMOVED_UNARY_BINARY_FUNCTION")
     set(CMAKE_CXX_FLAGS_DEBUG   "-O0 -g3")
     if (BUILD_COVERAGE)
-	set(CMAKE_CXX_FLAGS_RELEASE "-O3 -g -fprofile-arcs -ftest-coverage")
+        set(CMAKE_CXX_FLAGS_RELEASE "-O3 -g -fprofile-arcs -ftest-coverage")
     else()
         set(CMAKE_CXX_FLAGS_RELEASE "-O3 -g")
-    endif()	
+    endif()
 endif()
 
-set(CMAKE_CUDA_ARCHITECTURES "all-major")
-
-set(CMAKE_CUDA_FLAGS_RELEASE "-O3 -g -lineinfo")
-set(CMAKE_CUDA_FLAGS_DEBUG "-O0 -g -G")
-
 # add the binary tree to the search path for include files
-# so that we will find libmolgrid_config.h
 include_directories("${PROJECT_BINARY_DIR}")
 include_directories(BEFORE ${CMAKE_SOURCE_DIR}/include)
 include_directories(${PROJECT_BINARY_DIR}/include)
-include_directories(SYSTEM ${CUDA_INCLUDE_DIRS})
 include_directories(SYSTEM ${Boost_INCLUDE_DIR})
 
+# Apple Silicon MPS requires these frameworks
+find_library(METAL_LIB   Metal   REQUIRED)
+find_library(FOUNDATION_LIB Foundation REQUIRED)
+
 add_subdirectory(src)
 add_subdirectory(python)
 add_subdirectory(docs)
 
 # enable testing
 include(CTest)
 enable_testing()
-# define tests
 add_subdirectory(test)

README.md

@@ -1,9 +1,12 @@
-libmolgrid
-==========
-[![Github CI](https://github.com/gnina/libmolgrid/actions/workflows/CI.yml/badge.svg)](https://github.com/gnina/libmolgrid/actions)
-[![codecov](https://codecov.io/gh/gnina/libmolgrid/branch/master/graph/badge.svg?token=gUFisf34rf)](https://codecov.io/gh/gnina/libmolgrid)
+libmolgrid — macOS / Apple Silicon fork
+========================================
 
-libmolgrid is under active development, but should be suitable for use by early adopters.
+> **This is a macOS-only fork of libmolgrid** with GPU acceleration via
+> [Metal Performance Shaders (MPS)](https://developer.apple.com/documentation/metalperformanceshaders)
+> for Apple Silicon (M1/M2/M3/M4). The CUDA backend has been replaced by Metal.
+>
+> For Linux / CUDA support, and information, refer to the **original project**:
+> [https://gnina.github.io/libmolgrid/](https://gnina.github.io/libmolgrid/)
 
 If you use libmolgrid in your research, please cite:
 
@@ -22,49 +25,57 @@ If you use libmolgrid in your research, please cite:
 }
 ```
 
-## Documentation
+## Installation (macOS / Apple Silicon)
 
-[https://gnina.github.io/libmolgrid/](https://gnina.github.io/libmolgrid/)
+### Prerequisites
 
-## Installation
+Install dependencies via [Homebrew](https://brew.sh):
 
-### PIP
-
-```pip install molgrid```
-
-### conda
-```conda install -c gnina molgrid```
-
-### Build from Source
-
-```apt install git build-essential libboost-all-dev python3-pip rapidjson-dev
+```bash
+brew install cmake boost open-babel rapidjson eigen python
 pip3 install numpy pytest pyquaternion
 ```
 
-[Install cmake 3.12 or higher.](https://cmake.org/install/)
-
-[Install CUDA.](https://developer.nvidia.com/cuda-downloads)
+Xcode Command Line Tools are required for the Metal shader compiler (`metal`, `metallib`):
 
-[Install OpenBabel 3.0.](https://github.com/openbabel/openbabel)
+```bash
+xcode-select --install
+```
 
-`apt install libeigen3-dev libboost-all-dev`
+### Build from Source
 
-```git clone https://github.com/gnina/libmolgrid.git
+```bash
+git clone https://github.com/gnina/libmolgrid.git
 cd libmolgrid
-mkdir build
-cd build
-cmake ..
-make -j8
+mkdir build && cd build
+cmake .. -DCMAKE_BUILD_TYPE=Release
+make -j$(sysctl -n hw.logicalcpu)
 sudo make install
 ```
 
+The Metal shaders (`src/metal/shaders.metal`) are compiled automatically into a
+`.metallib` during the build. No CUDA installation is required.
 
+On Apple Silicon, GPU and CPU share unified memory — no explicit host↔device
+copies occur at runtime.
 
+### Run Tests
 
-## Example
+Tests must be run from the `build/bin/` directory so that relative data paths resolve correctly:
+
+```bash
+cd build/bin
+./test_gridmaker_cpp
+./test_gridmaker_mps_mm
+./test_grid_mps_mm
+./test_mgrid_mps_mm
+./test_transform_mps_mm
 ```
+
+## Example
+
+```python
 import molgrid
-import pytest
 import numpy as np
 import torch
 import torch.nn as nn
@@ -110,36 +121,32 @@ def test_train_torch_cnn():
         if isinstance(m, nn.Conv3d) or isinstance(m, nn.Linear):
             init.xavier_uniform_(m.weight.data)
 
-    batch_size = 50
-    e = molgrid.ExampleProvider(data_root=datadir+"/structs",balanced=True,shuffle=True)
+    e = molgrid.ExampleProvider(data_root=datadir+"/structs", balanced=True, shuffle=True)
     e.populate(fname)
 
     gmaker = molgrid.GridMaker()
     dims = gmaker.grid_dimensions(e.num_types())
-    tensor_shape = (batch_size,)+dims
+    tensor_shape = (batch_size,) + dims
 
-    model = Net(dims).to('cuda')
+    model = Net(dims).to('mps')
     model.apply(weights_init)
 
     optimizer = optim.SGD(model.parameters(), lr=0.01, momentum=0.9)
 
-    input_tensor = torch.zeros(tensor_shape, dtype=torch.float32, device='cuda')
+    input_tensor = torch.zeros(tensor_shape, dtype=torch.float32, device='mps')
     float_labels = torch.zeros(batch_size, dtype=torch.float32)
 
     losses = []
     for iteration in range(100):
-        #load data
         batch = e.next_batch(batch_size)
-        gmaker.forward(batch, input_tensor, 0, random_rotation=False) #not rotating since convergence is faster this way
+        gmaker.forward(batch, input_tensor, 0, random_rotation=False)
         batch.extract_label(0, float_labels)
-        labels = float_labels.long().to('cuda')
+        labels = float_labels.long().to('mps')
 
         optimizer.zero_grad()
         output = model(input_tensor)
-        loss = F.cross_entropy(output,labels)
+        loss = F.cross_entropy(output, labels)
         loss.backward()
         optimizer.step()
         losses.append(float(loss))
-
 ```
-

include/libmolgrid/common.h

@@ -1,54 +1,71 @@
 /*
  * common.h
  *
- * Utility functions and definitions
+ * Utility functions and definitions for Apple Silicon MPS backend.
+ * Replaces the previous CUDA-based common.h.
  */
 
 #ifndef COMMON_H_
 #define COMMON_H_
 
 #include <memory>
 #include <cstring>
-#include <cuda_runtime_api.h>
-#include <cuda.h>
+#include <stdexcept>
+#include <iostream>
+#include <cmath>
 
-#ifdef __CUDACC__
-#define CUDA_CALLABLE_MEMBER __host__ __device__
-#define CUDA_DEVICE_MEMBER __device__
-#else
+// ---- Vector types (replaces CUDA built-ins) ----
+// These match the CUDA naming convention so the rest of the code is unchanged.
+
+#ifndef __METAL_VERSION__
+
+struct float2 { float x, y; };
+struct float3 { float x, y, z; };
+struct float4 { float x, y, z, w; };
+struct int2   { int x, y; };
+struct int3   { int x, y, z; };
+struct uint2  { unsigned int x, y; };
+struct uint3  { unsigned int x, y, z; };
+struct double3 { double x, y, z; };
+
+inline float2  make_float2(float x, float y)                        { return {x, y}; }
+inline float3  make_float3(float x, float y, float z)               { return {x, y, z}; }
+inline float4  make_float4(float x, float y, float z, float w)      { return {x, y, z, w}; }
+inline int2    make_int2(int x, int y)                               { return {x, y}; }
+inline int3    make_int3(int x, int y, int z)                        { return {x, y, z}; }
+inline uint2   make_uint2(unsigned x, unsigned y)                    { return {x, y}; }
+inline uint3   make_uint3(unsigned x, unsigned y, unsigned z)        { return {x, y, z}; }
+inline double3 make_double3(double x, double y, double z)            { return {x, y, z}; }
+
+#endif // __METAL_VERSION__
+
+// ---- Callable member annotation (was __host__ __device__) ----
+// On Apple Silicon, all code runs on CPU or Metal GPU; no annotation needed.
 #define CUDA_CALLABLE_MEMBER
 #define CUDA_DEVICE_MEMBER
-#endif
 
-//called in device code to perform a parallel operation
+// ---- Thread / block size constants (shared with Metal shaders via headers) ----
+#define LMG_CUDA_NUM_THREADS 512
+#define LMG_CUDA_BLOCKDIM    8
+#define LOG2_WARP_SIZE       5U
+#define WARP_SIZE            (1U << LOG2_WARP_SIZE)
+
+// ---- Block/thread count helpers ----
+#define LMG_GET_BLOCKS(N)  ((unsigned(N) + LMG_CUDA_NUM_THREADS - 1) / LMG_CUDA_NUM_THREADS)
+#define LMG_GET_THREADS(N) (unsigned(N) < LMG_CUDA_NUM_THREADS ? unsigned(N) : LMG_CUDA_NUM_THREADS)
+
+// ---- Kernel loop macro (used in CPU fallback paths) ----
+// In Metal this pattern is handled by thread_position_in_grid.
 #define LMG_CUDA_KERNEL_LOOP(i, n) \
-  for (int i = blockIdx.x * blockDim.x + threadIdx.x; \
-       i < (n); \
-       i += blockDim.x * gridDim.x)
+  for (int i = 0; i < static_cast<int>(n); i++)
 
-// CUDA: use 512 threads per block
-#define LMG_CUDA_NUM_THREADS 512
-#define LMG_CUDA_BLOCKDIM 8
-#define LOG2_WARP_SIZE 5U
-#define WARP_SIZE (1U << LOG2_WARP_SIZE)
-
-// CUDA: number of blocks for threads.
-#define LMG_GET_BLOCKS(N) ((unsigned(N) + LMG_CUDA_NUM_THREADS - 1) / LMG_CUDA_NUM_THREADS)
-// CUDA: combined with GET_BLOCKS, number of threads
-#define LMG_GET_THREADS(N) min(N,LMG_CUDA_NUM_THREADS)
-
-#ifndef __CUDA_ARCH__
-#define LMG_CUDA_CHECK(condition) \
-  /* Code block avoids redefinition of cudaError_t error */ \
-  do { \
-    cudaError_t error = condition; \
-    if(error != cudaSuccess) {                                          \
-        std::cerr << __FILE__ << ":" << __LINE__ << ": " << cudaGetErrorString(error); throw std::runtime_error(std::string("CUDA Error: ")+cudaGetErrorString(error)); } \
-  } while (0)
-#else
-// probably don't want to make API calls on the device.
-#define LMG_CUDA_CHECK(condition) condition
-#endif
+// ---- Error check (was LMG_CUDA_CHECK) ----
+// Metal errors are handled via NSError in metal_context.mm.
+// This macro is kept as a no-op placeholder so call sites still compile.
+#define LMG_CUDA_CHECK(condition) (condition)
 
+// ---- itoa helper used in error messages ----
+#include <string>
+inline std::string itoa(size_t v) { return std::to_string(v); }
 
 #endif /* COMMON_H_ */