diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index 3682ce0a..617f0c5e 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -2,9 +2,9 @@ name: CI on: push: - branches: [ v2.x ] + branches: [ master, v2.x ] pull_request: - branches: [ v2.x ] + branches: [ master, v2.x ] schedule: # Run every Sunday at midnight - cron: '0 0 * * 0' @@ -25,45 +25,72 @@ jobs: # Oldest supported versions - name: Linux (CUDA 10.2) cuda: "10.2.89" - gcc: "8.5.*" - nvcc: "10.2" + gcc: "7.*" + cuda12: 'false' # Latest supported versions - - name: Linux (CUDA 11.8) - cuda: "11.8.0" - gcc: "10.3.*" - nvcc: "11.8" - + - name: Linux (CUDA 12) + cuda: "12.*" + gcc: "11.*" + cuda12: 'true' steps: - name: Check out uses: actions/checkout@v2 - - - name: Install Mamba - uses: mamba-org/provision-with-micromamba@main - with: - environment-file: false - extra-specs: "" - cache-env: true - cache-downloads: true + - name: Manage disk space + run: | + sudo mkdir -p /opt/empty_dir || true + for d in \ + /opt/ghc \ + /opt/hostedtoolcache \ + /usr/lib/jvm \ + /usr/local/.ghcup \ + /usr/local/lib/android \ + /usr/local/share/powershell \ + /usr/share/dotnet \ + /usr/share/swift \ + ; do + sudo rsync --stats -a --delete /opt/empty_dir/ $d || true + done + sudo apt-get purge -y -f firefox \ + google-chrome-stable \ + microsoft-edge-stable gcc + sudo apt-get autoremove -y >& /dev/null + sudo apt-get autoclean -y >& /dev/null + sudo docker image prune --all --force + df -h - name: Install CUDA Toolkit + # CUDA 12 can be installed with mamba + if: ${{ matrix.cuda12 == 'false' }} uses: Jimver/cuda-toolkit@v0.2.10 with: cuda: ${{ matrix.cuda }} - linux-local-args: '["--toolkit", "--override"]' + linux-local-args: '["--toolkit", "--override"]' - - name: Prepare dependencies + - name: Prepare dependencies (CUDA <12) + if: ${{ matrix.cuda12 == 'false'}} run: | - sed -i -e "/cudatoolkit/c\ - cudatoolkit ${{ matrix.cuda }}" \ - -e "/gxx_linux-64/c\ - gxx_linux-64 ${{ matrix.gcc }}" \ - -e "/nvcc_linux-64/c\ - nvcc_linux-64 ${{ matrix.nvcc }}" \ + sed -i -e "/cuda-version/d" \ + -e "/cuda-libraries-dev/d" \ + -e "/cuda-nvcc/d" \ + -e "/gxx_linux-64/c\ - gxx_linux-64 ${{ matrix.gcc }}\n - gcc_linux-64" \ environment.yml + - name: Prepare dependencies (CUDA >=12) + if: ${{ matrix.cuda12 == 'true' }} + run: | + sed -i -e "/cuda-version/c\ - cuda-version ${{ matrix.cuda }}" \ + -e "/gxx_linux-64/c\ - gxx_linux-64 ${{ matrix.gcc }}" \ + environment.yml - name: Show dependency file run: cat environment.yml - - - name: Install dependencies - run: micromamba env create -n uammd -f environment.yml + + - name: Install Mamba + uses: mamba-org/setup-micromamba@v1 + with: + environment-file: environment.yml + cache-environment: false + cache-downloads: true env: CONDA_OVERRIDE_CUDA: ${{ matrix.nvcc }} diff --git a/.github/workflows/docs.yaml b/.github/workflows/docs.yaml new file mode 100644 index 00000000..aadd243d --- /dev/null +++ b/.github/workflows/docs.yaml @@ -0,0 +1,24 @@ +name: Build Documentation + +on: + push: + branches: [ master, v2.x ] + pull_request: + branches: [ master, v2.x ] + +jobs: + build-docs: + runs-on: ubuntu-latest + + steps: + - uses: actions/checkout@v2 + - name: Install docs dependencies + run: | + pip install -r docs/requirements.txt + shell: bash -el {0} + + - name: Build Sphinx Documentation + run: | + cd docs + make html + shell: bash -el {0} diff --git a/.readthedocs.yaml b/.readthedocs.yaml new file mode 100644 index 00000000..18bb84d1 --- /dev/null +++ b/.readthedocs.yaml @@ -0,0 +1,17 @@ + +version: 2 + +build: + os: "ubuntu-22.04" + tools: + python: "3.11" + +sphinx: + configuration: docs/conf.py + +python: + install: + - requirements: docs/requirements.txt + +formats: + - pdf diff --git a/README.md b/README.md index 1513b963..cd222948 100644 --- a/README.md +++ b/README.md @@ -9,13 +9,8 @@ **See the wiki for more info!** **You can find videos on the youtube channel** http://bit.ly/2r5WoOn -## DESCRIPTION - ----------------- -Raul P. Pelaez 2018-2022. (raul.perez(at)uam.es) - - A C++14+ header-only fast generic multiscale CUDA Molecular Dynamics framework made with moduarity, expandability and generality in mind. UAMMD is intended to be hackable and copy pastable. Although "Molecular Dynamics" is part of the name,the UAMMD framework allows for much more than that. To this moment multiple integrators are implemented allowing it to perform: @@ -42,15 +37,40 @@ Hop on to the examples folder for an introduction or check the [documentation](h See the documentation page at https://uammd.readthedocs.io for a full list of available modules. ---------------------- -## USAGE +## Usage ------------------- +You can use UAMMD as a library for integration into other codes or as a standalone engine. + +#### DEPENDENCIES + +--------------------- +Depends on: + + 1. CUDA 9.x+ : https://developer.nvidia.com/cuda-downloads + +Some modules make use of certain NVIDIA libraries included with CUDA: + + 1. cuBLAS + 2. cuFFT + +Some modules also make use of lapacke and cblas (which can be replaced by mkl). +Apart from this, any dependency is already included in the repository under the third_party folder. +See [Compiling UAMMD](https://uammd.readthedocs.io/en/latest/Compiling-UAMMD.html) in the documentation for more information. + +Every required dependency can be installed using conda with the environment file provided in the repository. We recommend using [mamba](https://mamba.readthedocs.io/en/latest/installation/mamba-installation.html) as a replacement for conda. + +```bash +mamba env create -f environment.yml +``` + +### Library mode + **UAMMD does not need to be compiled separatedly (it is header only)**. Some special flags might be needed to compile codes including with certain UAMMD headers, see [Compiling UAMMD](https://uammd.readthedocs.io/en/latest/Compiling-UAMMD.html). -Here you have a short example of how a typical UAMMD code looks like: - +Here you have a short example of how a typical UAMMD code looks like, encoding a simple Brownian dynamics simulation of non interacting particles.: ```c++ //Ideal brownian particles @@ -79,23 +99,40 @@ int main(int argc, char * argv[]){ ``` -Drop by the examples folder to get started with UAMMD or go to the [wiki](https://uammd.readthedocs.io/). +Drop by the examples folder to get started with UAMMD or go to the [documentation page for the examples](https://uammd.readthedocs.io/en/latest/Examples.html). + +### Stand alone engine + +The example `generic_md` includes almost every module available in UAMMD and can be configured from a parameter file. Go to `examples/generic_md` for instructions. -## DEPENDENCIES +## Running Tests --------------------- -Depends on: - 1. CUDA 9.x+ : https://developer.nvidia.com/cuda-downloads +The `tests` folder contains instructions on how to run the UAMMD correctness tests. You can also go to the [documentation page for tests](https://uammd.readthedocs.io/en/latest/Tests.html). -Some modules make use of certain NVIDIA libraries included with CUDA: - - 1. cuBLAS - 2. cuFFT - -Some modules also make use of lapacke and cblas (which can be replaced by mkl). -Apart from this, any dependency is already included in the repository under the third_party folder. -See [Compiling UAMMD](https://uammd.readthedocs.io/en/latest/Compiling-UAMMD.html) in the documentation for more information. + +## Repository structure + +------------------------- + +- The `docs` folder contains the documentation source files for [](https://uammd.readthedocs.io/) and scripts to build it. + +- The `examples` folder contains several examples of UAMMD usage. + +- The `src` folder contains the source code for UAMMD. Inside it, there are the following folders: +1. `global`: Definitions proper to the whole code. +2. `Integrator`: Source code for the [integrators](https://uammd.readthedocs.io/en/latest/Integrators.html). +3. `Interactor`: Source code for the [interactors](https://uammd.readthedocs.io/en/latest/Interactors.html). +4. `misc`: Miscellaneous source code. +5. `ParticleData`: Source code for [particle data](https://uammd.readthedocs.io/en/latest/ParticleData.html). +6. `System`: Source code for [system](https://uammd.readthedocs.io/en/latest/System.html). +7. `third_party`: Third party source code. +8. `utils`: Utilities for UAMMD. + +- The `test` folder contains the source code for the UAMMD tests. + +- The `extensions` is part of the subsystem for UAMMD extensions, currently a work in progress. ------------------------------------------ @@ -112,7 +149,8 @@ Raul P. Pelaez is the main developer of UAMMD. Other people that have contributed to UAMMD (thanks!): Marc Melendez Schofield +Pablo Ibañez Freire (https://github.com/PabloIbannez) +Pablo Palacios Alonso (http://github.com/PabloPalaciosAlonso) Sergio Panzuela Nerea Alcazar -Pablo Ibañez Freire (https://github.com/PabloIbannez) Salvatore Assenza diff --git a/compile_flags.txt b/compile_flags.txt deleted file mode 100644 index fa397140..00000000 --- a/compile_flags.txt +++ /dev/null @@ -1,11 +0,0 @@ --xcuda --DMAXLOGLEVEL=5 --std=c++14 --O3 --I/usr/local/cuda/include ---cuda-gpu-arch=sm_52 --Isrc --Isrc/third_party --I/usr/include/lapacke --I/usr/include/cblas ---cuda-path=/usr/local/cuda diff --git a/docs/.readthedocs.yaml b/docs/.readthedocs.yaml deleted file mode 100644 index 69f98b1f..00000000 --- a/docs/.readthedocs.yaml +++ /dev/null @@ -1,16 +0,0 @@ - -version: 2 - -formats: all - -build: - os: "ubuntu-20.04" - tools: - python: "3.10" - -sphinx: - configuration: conf.py - -python: - install: - - requirements: requirements.txt diff --git a/docs/Compiling-UAMMD.rst b/docs/Compiling-UAMMD.rst index 7dc501b2..697b2598 100644 --- a/docs/Compiling-UAMMD.rst +++ b/docs/Compiling-UAMMD.rst @@ -7,14 +7,16 @@ UAMMD is header-only, so a module is compiled only when its header is included. #. CUDA 8.0+ (https://developer.nvidia.com/cuda-downloads ) #. A C++ compiler with C++14 support (g++ 5+ will probably do) - #. Thrust (The version shipped with the cuda release will do) #. LAPACKE/CBLAS or Intel MKL (For some modules only) The newest compiler versions I have tested are g++-12.x and clang++-13.0 with cuda-12 in Fedora 37. You can even compile a source containing UAMMD code with clang++-7+ alone, without nvcc. +**These dependencies can be installed using conda with the provided environment.yaml file.** + Additionally, UAMMD makes use of the following external libraries, **which are already included in the repo under third_party**. You can either compile using these or place symlinks to your preferred versions (I have seen recent versions of cub not compiling on some platforms). - * Boost Preprocessor (http://www.boost.org/ ) (Extracted using bcp, just a few headers) + * Boost Preprocessor (http://www.boost.org/ ) (Extracted using bcp, just a few headers) + * Thrust (https://github.com/thrust/thrust ) * CUB 1.5.2+ (https://github.com/NVlabs/cub ) (Currently 1.8.0 in third_party/cub_bak*) * nod (https://github.com/fr00b0/nod ) (A lightweight C++11 signal/slot library) * SaruPRNG (http://dx.doi.org/10.1016/j.cpc.2012.12.003 ) (A parallel friendly RNG) diff --git a/docs/Integrator/BrownianHydrodynamics.rst b/docs/Integrator/BrownianHydrodynamics.rst index 41d2c18b..37f6d6e6 100644 --- a/docs/Integrator/BrownianHydrodynamics.rst +++ b/docs/Integrator/BrownianHydrodynamics.rst @@ -989,8 +989,26 @@ In addition to the previous ones, the integrator also requires: Both the solver and Integrator will fail if some particle lies beyond the domain limits in the z direction. +Computing average velocity in the plane directions +************************************************** +The class :code:`DPStokesSlab_ns::DPStokes` can also be used to compute the average velocity of a group of particles in the plane directions. + +.. cpp:function:: template std::vector DPStokes::computeAverageVelocity(PosIterator pos, ForceIterator forces, int numberParticles, int direction = 0, cudaStream_t st = 0) + + + Computes the average velocity :math:`\langle v(z) \rangle_{x/y}` a group of particles in the plane directions. + + :param PosIterator pos: Iterator to the positions of the particles. + :param ForceIterator forces: Iterator to the forces acting on the particles. + :param int numberParticles: Number of particles. + :param int direction: Direction of the average velocity. 0 for x, 1 for y. + :param cudaStream_t st: CUDA stream where the computation will be performed. + :returns: A vector with the average velocity (size n.z). + + + .. rubric:: References: .. [1] An Introduction to Dynamics of Colloids. Dhont, J.K.G. 1996. https://www.elsevier.com/books/an-introduction-to-dynamics-of-colloids/dhont/978-0-444-82009-9 diff --git a/docs/Tests.rst b/docs/Tests.rst index 803443a6..3d771d7e 100644 --- a/docs/Tests.rst +++ b/docs/Tests.rst @@ -3,14 +3,15 @@ Tests The folder "test" in the main UAMMD repository contains test scripts for the different solvers. Mostly these consist on a bash script accompanied by a readme and an UAMMD code reproducing a certain known solution. The test script will typically generate messages stating the result of the tests in addition to some figures that can be inspected to evaluate the correctness of the solver. -In order to run them simply go into the folder of a certain test and run :code:`bash test.bash`. Note that these tests can typically take a lot of time, as averaging is required. Check the test script for a certain solver for more information and customization. You may have to tweak the Makefiles. Finally, test scripts usually include the possibility of using many GPUs. +In order to run them go into the folder of a certain test and run :code:`bash test.bash`. Note that these tests can typically take a lot of time, as averaging is required. Check the test script for a certain solver for more information and customization. You may have to tweak the Makefiles. Finally, test scripts usually include the possibility of using many GPUs. -Recent Unit Tests are written using `GTest `. In these instances a CMakeLists.txt file is present. To compile and run the tests use the standard steps: +Recent Unit Tests are written using `GTest `_. In these instances a CMakeLists.txt file is present. To compile and run the tests use the standard steps: .. code:: bash mkdir build && cd build cmake .. #Solve any dependency issues raised by CMake - make make test + +Note that some tests are stochastic in nature and can thus fail intermittently. Typically rerunning these tests solves the issue. diff --git a/environment.yml b/environment.yml index 38c87124..70fc0573 100644 --- a/environment.yml +++ b/environment.yml @@ -1,12 +1,13 @@ +name: uammd channels: - conda-forge dependencies: - cmake >=3.22 - - cudatoolkit 11.8.* - - gxx_linux-64 10.3.* + - cuda-version 12.* + - gxx_linux-64 11.* + - cuda-libraries-dev + - cuda-nvcc - make - - nvcc_linux-64 11.8 - - sysroot_linux-64 2.17 - mkl-devel - pybind11 - - python 3.10.* + - python 3.11.* diff --git a/examples/README.md b/examples/README.md index 1d178f8e..db448bde 100644 --- a/examples/README.md +++ b/examples/README.md @@ -3,8 +3,6 @@ In these folders you will find a lot of examples showcasing the different capabi UAMMD is presented with two distinct identities, it can work as a standalone molecular simulation GPU code (such as gromacs or hoomd) or as a library exposing GPU algorithms for you to use to accelerate your own code. Keep in mind that UAMMD is designed to be hackable and copy pastable, so feel free to take what you want and give nothing back! -The Makefile in this folder simply goes into each folder and runs the Makefiles there. - This folder is organized in the following way: ### basic_concepts @@ -16,8 +14,7 @@ Stuff not covered by the basic tutorial with more complex and obscure functional ### generic_md This code has almost every module that UAMMD offers in a single file that you can compile and then tweak via a data.main parameter file. -If you are not looking to learn how UAMMD works and how to modify it or extend it get in here. -You might be able to encode your simulation by tweaking the data.main. +If you are not looking to learn how UAMMD works and how to modify it or extend it get in here. You might be able to encode your simulation by tweaking the data.main. ### integration_schemes The basic tutorial covers only a couple of integration modules, in this folder you will find copy pastable functions to create any UAMMD Integrator. From Brownian Dynamics to Smoothed Particle Hydrodynamics. @@ -31,3 +28,27 @@ This family of examples shows off want you can do outside the UAMMD simulation e ### misc Examples that do not fit in the other categories. +## Compiling the examples + +For convenience and learning purposes the examples can be compiled either with Make or CMake. This folder contains both a Makefile and a CMakeLists.txt. + +I suggest using the CMake system, as it tends to be more robust regarding install locations of the different dependencies. When running cmake a Makefile will be generated. + +To compile every example, run: + +```bash + mkdir -p build + cd build + cmake .. + make -j3 +``` + +Binaries for the examples will be available under the created `bin` folder. + +Refer to [Compiling UAMMD](https://github.com/RaulPPelaez/UAMMD/wiki/Compiling-UAMMD) for additional information if the building process fails. + +You can pass to CMake any UAMMD related definitions, for instance: + +```shell +cmake -DDOUBLE_PRECISION .. +``` diff --git a/examples/generic_md/README.md b/examples/generic_md/README.md index 2fa4d966..636ebc61 100644 --- a/examples/generic_md/README.md +++ b/examples/generic_md/README.md @@ -29,10 +29,8 @@ With this components as-is one can simulate systems from a microscale LJ-liquid ### USAGE -Compile with ```$ make```, you might have to customize the Makefile first to adapt it to your system. -If you are having trouble with it, start by going into [Compiling UAMMD](https://github.com/RaulPPelaez/UAMMD/wiki/Compiling-UAMMD) -Then execute ```$ ./generic```. If you have a valid CUDA environment this will generate an example data.main file and then halt. +After compiling, execute ```$ ./generic```. If you have a valid CUDA environment this will generate an example data.main file and then halt. Modify this file as desired and then run ```$ ./generic``` again. The default data.main will perform a LJ-liquid simulation with a VerletNVT integrator. @@ -50,11 +48,11 @@ If you need to create a BD integrator, you can copy paste the function createInt Need to read parameters from a file? just copy paste the function readParameters. And so on. -If you would like a more bottom up approach to UAMMD, you can surf the examples/basic folder, which will give you through UAMMD with an increasingly complex set of example codes. -Additionally, you can drop by the wiki: -https://github.com/RaulPPelaez/UAMMD/wiki +If you would like a more bottom up approach to UAMMD, you can surf the `examples/basic_concepts` folder, which will give you through UAMMD with an increasingly complex set of example codes. +Additionally, you can drop by the docs: +https://uammd.readthedocs.io/ Which has a lot of information. From basic functionality to descriptions and references for the algorithms implemented by each module. This code is an example on how to work inside UAMMD's simulation ecosystem, however UAMMD is not restricted to this kind of usage. It is designed to be as hackable as possible and exposes a lot of the core functionality as stand alone with the intention of making it work as an external GPU accelerator in other codes. -See the folder examples/uammd_as_a_library for examples on how to do this. +See the folder `examples/uammd_as_a_library` for examples on how to do this. diff --git a/src/misc/IBM.cuh b/src/misc/IBM.cuh index df5f3352..6581a834 100644 --- a/src/misc/IBM.cuh +++ b/src/misc/IBM.cuh @@ -54,7 +54,7 @@ REFERENCES: #include "IBM.cu" #include "IBM_utils.cuh" #include "utils/Grid.cuh" - +#include "third_party/type_names.h" namespace uammd{ namespace IBM_ns{ struct LinearIndex3D{ diff --git a/test/BDHI/DPStokes/dpstokes_test.cu b/test/BDHI/DPStokes/dpstokes_test.cu index 8d3eecf6..c749790b 100644 --- a/test/BDHI/DPStokes/dpstokes_test.cu +++ b/test/BDHI/DPStokes/dpstokes_test.cu @@ -102,9 +102,9 @@ TEST(DPStokes, ReproducesOpenBoundarySelfMobilityWithLargeDomainSlitChannel){ force[0] = {1,1,1}; auto mf = dpstokes->Mdot(pos.data().get(), force.data().get(), 1); real3 disp = mf[0]; - ASSERT_THAT(disp.x, ::testing::DoubleNear(1.0, 1e-2)); - ASSERT_THAT(disp.y, ::testing::DoubleNear(1.0, 1e-2)); - ASSERT_THAT(disp.z, ::testing::DoubleNear(1.0, 1e-2)); + ASSERT_THAT(disp.x, ::testing::DoubleNear(1.0, 0.1)); + ASSERT_THAT(disp.y, ::testing::DoubleNear(1.0, 0.1)); + ASSERT_THAT(disp.z, ::testing::DoubleNear(1.0, 0.1)); } TEST(DPStokes, ReproducesOpenBoundarySelfMobilityWithLargeDomainBottomWall){ @@ -120,9 +120,9 @@ TEST(DPStokes, ReproducesOpenBoundarySelfMobilityWithLargeDomainBottomWall){ force[0] = {1,1,1}; auto mf = dpstokes->Mdot(pos.data().get(), force.data().get(), 1); real3 disp = mf[0]; - ASSERT_THAT(disp.x, ::testing::DoubleNear(1.0, 1e-2)); - ASSERT_THAT(disp.y, ::testing::DoubleNear(1.0, 1e-2)); - ASSERT_THAT(disp.z, ::testing::DoubleNear(1.0, 1e-2)); + ASSERT_THAT(disp.x, ::testing::DoubleNear(1.0, 0.1)); + ASSERT_THAT(disp.y, ::testing::DoubleNear(1.0, 0.1)); + ASSERT_THAT(disp.z, ::testing::DoubleNear(1.0, 0.1)); } TEST(DPStokes, ReproducesOpenBoundarySelfMobilityWithLargeDomainOpen){ @@ -138,9 +138,9 @@ TEST(DPStokes, ReproducesOpenBoundarySelfMobilityWithLargeDomainOpen){ force[0] = {1,1,1}; auto mf = dpstokes->Mdot(pos.data().get(), force.data().get(), 1); real3 disp = mf[0]; - ASSERT_THAT(disp.x, ::testing::DoubleNear(1.0, 1e-2)); - ASSERT_THAT(disp.y, ::testing::DoubleNear(1.0, 1e-2)); - ASSERT_THAT(disp.z, ::testing::DoubleNear(1.0, 1e-2)); + ASSERT_THAT(disp.x, ::testing::DoubleNear(1.0, 0.1)); + ASSERT_THAT(disp.y, ::testing::DoubleNear(1.0, 0.1)); + ASSERT_THAT(disp.z, ::testing::DoubleNear(1.0, 0.1)); } //TESTS for the Integrator start here diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index d1579cef..89741b4e 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -62,12 +62,15 @@ FetchContent_Declare( GIT_TAG release-1.12.1 ) FetchContent_MakeAvailable(googletest) - +enable_testing() +include(CTest) add_subdirectory(BDHI/DPStokes) add_subdirectory(BDHI/FCM) add_subdirectory(BDHI/quasi2D) add_subdirectory(utils) add_subdirectory(misc/chebyshev) +add_subdirectory(misc/ibm) +add_subdirectory(misc/lanczos) IF (CMAKE_BUILD_TYPE MATCHES "Debug") if(CMAKE_CUDA_COMPILER_ID STREQUAL "NVIDIA") diff --git a/test/misc/ibm/CMakeLists.txt b/test/misc/ibm/CMakeLists.txt new file mode 100644 index 00000000..6bb4f7aa --- /dev/null +++ b/test/misc/ibm/CMakeLists.txt @@ -0,0 +1,10 @@ +add_executable(ibm test_ibm.cu) +target_include_directories(ibm PUBLIC ${CMAKE_CURRENT_SOURCE_DIR}) +set_target_properties(ibm PROPERTIES RUNTIME_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/bin") +target_link_libraries( + ibm + GTest::gtest_main + GTest::gmock_main +) +include(GoogleTest) +gtest_discover_tests(ibm) diff --git a/test/misc/ibm/Makefile b/test/misc/ibm/Makefile deleted file mode 100644 index e517885d..00000000 --- a/test/misc/ibm/Makefile +++ /dev/null @@ -1,25 +0,0 @@ - - -#Default log level is 5, which prints up to MESSAGE, 0 will only print critical errors and 14 will print everything up to the most low level debug information -LOG_LEVEL=5 - -NVCC=nvcc -CXX=g++ - -CUDA_ROOT=$(shell dirname `which nvcc`)/.. -UAMMD_ROOT=../../../ -#Uncomment to compile in double precision mode -DOUBLE_PRECISION=-DDOUBLE_PRECISION -INCLUDEFLAGS=-I$(CUDA_ROOT)/include -I$(UAMMD_ROOT)/src -I$(UAMMD_ROOT)/src/third_party -NVCCFLAGS=-ccbin=$(CXX) -std=c++14 -O3 $(INCLUDEFLAGS) -DMAXLOGLEVEL=$(LOG_LEVEL) $(DOUBLE_PRECISION) --extended-lambda - -all: $(patsubst %.cu, %, $(wildcard *.cu)) - -%: %.cu Makefile - $(NVCC) $(NVCCFLAGS) $< -o $(@:.out=) - - -clean: $(patsubst %.cu, %.clean, $(wildcard *.cu)) - -%.clean: - rm -f $(@:.clean=) diff --git a/test/misc/ibm/ibm.cu b/test/misc/ibm/ibm.cu deleted file mode 100644 index 7ad0cd72..00000000 --- a/test/misc/ibm/ibm.cu +++ /dev/null @@ -1,150 +0,0 @@ - -#include "utils.cuh" -#include - - -real phi(real r, real rmax, real sigma){ - r= abs(r); - real res = 0; - if(r<=rmax) - res=(1.0/(sigma*sqrt(2*M_PI)))*exp(-0.5*r*r/(sigma*sigma)); - return res; -} - -void runSpreadingTest(real3 qi, real3 L, int3 n, int supp){ - real3 h = L/make_real3(n); - real sigma = h.x; - complex Fi = complex{1,0}; - bool error = false; - int3 ci = chebyshev::doublyperiodic::Grid(Box(L), n).getCell(qi); - System::log("Spreading a particle at %g %g %g (cell %d %d %d)", - qi.x, qi.y, qi.z, - ci.x, ci.y, ci.z); - std::vector fr = spreadParticles({qi}, {Fi}, n, sigma, supp, L); - complex maxErr{}; - int3 cellOfMaxErr={-1,-1,-1}; - Box box(L); - box.setPeriodicity(1, 1, 0); - for(int i = 0; i 1e-13 or (norm(fr[id]) == 0 and norm(fijk) != 0) or - (norm(fijk) == 0 and norm(fr[id]) != 0)){ - error = true; - System::log("Difference in cell %d %d %d: Found %g %g, expected %g %g (error %g)", - i,j,k, fr[id].real(), fr[id].imag(), fijk.real(), fijk.imag(), - thrust::norm(err)); - } - } - } - } - if(cellOfMaxErr.x >= 0 and norm(maxErr) > 1e-14){ - int id = cellOfMaxErr.x + (cellOfMaxErr.y+cellOfMaxErr.z*n.y)*n.x; - real rmax = supp*h.x*0.5; - real3 r; - r.x = -L.x*0.5 + cellOfMaxErr.x*h.x-qi.x; - r.y = -L.y*0.5 + cellOfMaxErr.y*h.y-qi.y; - r.z = 0.5*L.z*cospi(cellOfMaxErr.z/(n.z-1.0))-qi.z; - r= box.apply_pbc(r); - complex fijk = {0,0}; - fijk = Fi*phi(r.x, rmax, sigma)*phi(r.y, rmax, sigma)*phi(r.z, rmax, sigma); - System::log("Maximum error %g %g found at cell %d %d %d (expected %.15g %.15g, got %.15g %.15g)", - maxErr.real(), maxErr.imag(), - cellOfMaxErr.x, cellOfMaxErr.y, cellOfMaxErr.z, - fijk.real(), fijk.imag(), - fr[id].real(), fr[id].imag() - ); - } - if(not error){ - System::log("[SUCCESS]"); - } -} - -//Spreads particles forces and checks that the correct force density has been created. -void checkSpreading(){ - std::random_device r; - std::default_random_engine e1(1234); - std::uniform_real_distribution uniform(-0.5, 0.5); - real3 L = {1,1,1}; - int3 n = {16,16,16}; - //I check an even and an odd support - int supp = 5; - int ntest = 5; - for(int test = 0; test("Interpolating particle at %g %g %g\n", qi.x, qi.y, qi.z); - complex Fi = complex{1,0}; - real sigma = 0.1; - std::vector fr = spreadParticles({qi}, {Fi}, n, sigma, supp, L); - auto res = interpolateField({qi}, fr, n, sigma, supp, L); - real l = L.x; - real solution = (9278850.0* pow(M_PI,1.5)*pow(erf(0.5*l/sigma),3))/(2301620723.0*pow(sigma,3)); - real error = abs(res[0].real()-solution)/solution; - if(error>1e-11){ - System::log("Too much error in quadrature"); - System::log("Error in quadrature %g, got %.15g, expected %.15g", - error, res[0].real(), solution); - } - else{ - System::log("[SUCCESS]"); - } - -} -//Checks that F = JSF, i.e interpolating after spreading returns the original quantity. -void checkJSF(){ - System::log("Interpolation after spreading (quadrature) test"); - int3 n = {128,128,128}; - real l = 1; - real3 L = {l,l,l}; - int supp = n.x; - real3 qi = make_real3(0,0,0)*L; - runInterpolationTest(qi, L, n, supp); -} - - -int main(int argc, char* argv[]){ - { - System::log("Checking spreading mechanism"); - checkSpreading(); - System::log("Checking interpolation mechanism"); - checkJSF(); - } - return 0; -} diff --git a/test/misc/ibm/test_ibm.cu b/test/misc/ibm/test_ibm.cu new file mode 100644 index 00000000..758de58e --- /dev/null +++ b/test/misc/ibm/test_ibm.cu @@ -0,0 +1,167 @@ + +#include "utils.cuh" +#include "gmock/gmock.h" +#include +#include +#include + +real phi(real r, real rmax, real sigma) { + r = abs(r); + real res = 0; + if (r <= rmax) + res = + (1.0 / (sigma * sqrt(2 * M_PI))) * exp(-0.5 * r * r / (sigma * sigma)); + return res; +} + +real phiz(real r, real rmax, real sigma, real H, real pz) { + if (abs(r) <= rmax) { + real bot_rimg = fabs(-H - pz + r); + return phi(r, rmax, sigma) - + ((bot_rimg >= rmax) ? real(0.0) : phi(bot_rimg, rmax, sigma)); + } else + return 0; +} + +real3 distanceToCell(int3 cell, real3 pos, int3 n, real3 L) { + real3 h = L / make_real3(n); + Box box(L); + box.setPeriodicity(1, 1, 0); + real3 r; + r.x = -L.x * 0.5 + (cell.x) * h.x - pos.x; + r.y = -L.y * 0.5 + (cell.y) * h.y - pos.y; + r.z = L.z * cospi(cell.z / (n.z - 1.0)) - pos.z; + r = box.apply_pbc(r); + return r; +} + +bool runSpreadingTest(real3 qi, real3 L, int3 n, int supp) { + real sigma = L.x / n.x; + real rmax = supp * L.x * 0.5 / n.x; + complex Fi = complex{1, 0}; + bool error = false; + int numberParticles = 1; + thrust::device_vector pos(numberParticles); + thrust::device_vector forces(numberParticles); + pos[0] = qi; + forces[0] = Fi; + auto fr = spreadParticles(pos, forces, n, sigma, supp, L); +#ifdef DOUBLE_PRECISION + constexpr real errorThreshold = 1e-12; +#else + constexpr real errorThreshold = 1e-5; +#endif + for (int i = 0; i < n.x; i++) { + for (int j = 0; j < n.y; j++) { + for (int k = 0; k < n.z; k++) { + real3 r = distanceToCell({i, j, k}, qi, n, L); + complex fijk = {0, 0}; + fijk = Fi * phi(r.x, rmax, sigma) * phi(r.y, rmax, sigma) * + phiz(r.z, rmax, sigma, L.z, qi.z); + int id = i + (j + k * n.y) * n.x; + real errnorm = std::max(fijk.real(), fijk.imag()); + if (errnorm == 0) { + errnorm = std::max(fr[id].real(), fr[id].imag()); + if (errnorm == 0) + errnorm = 1; + } + complex err = (fr[id] - fijk) / errnorm; + err.real(abs(err.real())); + err.imag(abs(err.imag())); + if (norm(err) > errorThreshold or + (norm(fr[id]) == 0 and norm(fijk) != 0) or + (norm(fijk) == 0 and norm(fr[id]) != 0)) { + error = true; + int3 ci = chebyshev::doublyperiodic::Grid(Box(L), n).getCell(qi); + System::log( + "Spreading a particle at %g %g %g (cell %d %d %d)", qi.x, qi.y, + qi.z, ci.x, ci.y, ci.z); + System::log("Difference in cell %d %d %d: Found %g " + "%g, expected %g %g (error %g)", + i, j, k, fr[id].real(), fr[id].imag(), + fijk.real(), fijk.imag(), + thrust::norm(err)); + return not error; + } + } + } + } + return not error; +} + +TEST(SpreadingTest, RandomPlacements) { + std::random_device r; + std::default_random_engine e1(r()); + std::uniform_real_distribution uniform(-0.5, 0.5); + real3 L = {1, 1, 1}; + int3 n = {16, 16, 16}; + // I check an even and an odd support + int supp = 5; + int ntest = 5; + for (int test = 0; test < ntest; test++) { + real3 qi = make_real3(uniform(e1), uniform(e1), uniform(e1)) * L; + ASSERT_TRUE(runSpreadingTest(qi, L, n, supp)); + } + supp = 6; + for (int test = 0; test < ntest; test++) { + real3 qi = make_real3(uniform(e1), uniform(e1), uniform(e1)) * L; + ASSERT_TRUE(runSpreadingTest(qi, L, n, supp)); + } +} + +TEST(SpreadingTest, DangerousPlacementsOddSupport) { + real3 L = {1, 1, 1}; + int3 n = {16, 16, 16}; + int supp = 5; + ASSERT_TRUE(runSpreadingTest(make_real3(0, 0, -L.z * 0.5), L, n, supp)); + ASSERT_TRUE(runSpreadingTest(make_real3(0, 0, L.z * 0.5), L, n, supp)); + ASSERT_TRUE( + runSpreadingTest(make_real3(0, 0, -L.z * 0.5 + 0.001), L, n, supp)); +} + +TEST(SpreadingTest, DangerousPlacementsEvenSupport) { + real3 L = {1, 1, 1}; + int3 n = {16, 16, 16}; + int supp = 6; + ASSERT_TRUE(runSpreadingTest(make_real3(0, 0, -L.z * 0.5), L, n, supp)); + ASSERT_TRUE(runSpreadingTest(make_real3(0, 0, L.z * 0.5), L, n, supp)); + ASSERT_TRUE( + runSpreadingTest(make_real3(0, 0, -L.z * 0.5 + 0.001), L, n, supp)); +} + +TEST(SpreadingTest, WholeDomain) { + real3 L = {1, 1, 1}; + ASSERT_TRUE( + runSpreadingTest(make_real3(0, 0, 0.001), L, {128, 128, 128}, 128)); +} + +void runInterpolationTest(real3 qi, real3 L, int3 n, int supp) { + complex Fi = complex{1, 0}; + real sigma = 0.1; + thrust::device_vector pos(1); + thrust::device_vector forces(1); + pos[0] = qi; + forces[0] = Fi; + thrust::device_vector fr = + spreadParticles(pos, forces, n, sigma, supp, L); + auto res = interpolateField(pos, fr, n, sigma, supp, L); + real l = L.x; + // The double integral of a Gaussian of width sigma in a cubic domain of size + // l + real solution = (9278850.0 * pow(M_PI, 1.5) * pow(erf(0.5 * l / sigma), 3)) / + (2301620723.0 * pow(sigma, 3)); + real JS = complex(res[0]).real(); + real error = abs(JS - solution) / solution; + ASSERT_LE(error, 1e-11); +} +// Checks that JS1=\int \delta_a(\vec{r})^2 d\vec{r}, i.e interpolating +// after spreading returns the inverse of the volume. +// Here \delta_a is the spreading kernel +TEST(SpreadingTest, InterpolateAfterSpreadYieldsInverseVolume) { + int3 n = {128, 128, 128}; + real l = 1; + real3 L = {l, l, l}; + int supp = n.x; + real3 qi = make_real3(0, 0, 0) * L; + runInterpolationTest(qi, L, n, supp); +} diff --git a/test/misc/ibm/utils.cuh b/test/misc/ibm/utils.cuh index 9b9264f0..c57620ff 100644 --- a/test/misc/ibm/utils.cuh +++ b/test/misc/ibm/utils.cuh @@ -1,46 +1,68 @@ #pragma once #include -#include "misc/IBM.cuh" #include "misc/ChevyshevUtils.cuh" -#include +#include "misc/IBM.cuh" +#include + using namespace uammd; using complex = thrust::complex; using real = uammd::real; -struct Gaussian{ +struct Gaussian { int3 support; - Gaussian(real width, real h, real H, real nz, int supportxy): - nz(nz){ - this-> prefactor = 1.0/(width*sqrt(2*M_PI)); - this-> tau = -1.0/(2.0*width*width); - this->rmax = supportxy*h*0.5; - support.x = supportxy; + Gaussian(real width, real h, real H, real nz, int supportxy) : nz(nz) { + this->prefactor = 1.0 / (width * sqrt(2 * M_PI)); + this->tau = -1.0 / (2.0 * width * width); + this->rmax = supportxy * h * 0.5; + support.x = supportxy + 1; support.y = support.x; this->Htot = H; - real smax = std::max((H*0.5-rmax), -H*0.5); - int czmax = ceil((nz-1)*(acos(smax/(0.5*Htot))/real(M_PI))); - support.z = 2*czmax+1; + real smax = std::max((H * 0.5 - rmax), -H * 0.5); + int czmax = ceil((nz - 1) * (acos(smax / (0.5 * Htot)) / real(M_PI))); + support.z = 2 * czmax + 1; } - inline __host__ __device__ int3 getMaxSupport() const{ + inline __host__ __device__ int3 getMaxSupport() const { return make_int3(support.x, support.y, support.z); } - inline __host__ __device__ int3 getSupport(real3 pos, int3 cell) const{ - real bound = Htot*real(0.5); - real ztop = thrust::min(pos.z+rmax, bound); - real zbot = thrust::max(pos.z-rmax, -bound); - int czb = int((nz-1)*(acos(ztop/bound)/real(M_PI))); - int czt = int((nz-1)*(acos(zbot/bound)/real(M_PI))); - int sz = 2*thrust::max(cell.z - czb, czt - cell.z)+1; + inline __host__ __device__ int3 getSupport(real3 pos, int3 cell) const { + real bound = Htot * real(0.5); + real ztop = thrust::min(pos.z + rmax, bound); + real zbot = thrust::max(pos.z - rmax, -bound); + int czb = int((nz - 1) * (acos(ztop / bound) / real(M_PI))); + int czt = int((nz - 1) * (acos(zbot / bound) / real(M_PI))); + int sz = 2 * thrust::max(cell.z - czb, czt - cell.z) + 1; return make_int3(support.x, support.y, sz); } - __host__ __device__ real phi(real r, real3 pos) const{ + inline __host__ __device__ real phiX(real r, real3 pos) const { + real val = 0; + if (abs(r) <= rmax * real(1.0001)) { + val = prefactor * exp(tau * r * r); + } + return val; + } + + inline __host__ __device__ real phiY(real r, real3 pos) const { + real val = 0; + if (abs(r) <= rmax * real(1.0001)) { + val = prefactor * exp(tau * r * r); + } + return val; + } + + __host__ __device__ real phiZ(real r, real3 pos) const { real val = 0; - if(abs(r)<=rmax*real(1.0001)){ - val = prefactor*exp(tau*r*r); + if (abs(r) <= rmax * real(1.0001)) { + // val = prefactor*exp(tau*r*r); + real bot_rimg = -this->Htot - real(2.0) * pos.z + r; + real rimg = fabs(bot_rimg); + real phi_img = + rimg >= rmax ? real(0.0) : prefactor * exp(tau * rimg * rimg); + real phi = prefactor * exp(tau * r * r); + val = phi - phi_img; } return val; } @@ -53,42 +75,49 @@ private: real Htot; }; -auto spreadParticles(std::vector pos, std::vector values, - int3 n, real sigma, int supp, real3 L){ - std::vector fr(n.x*n.y*n.z, complex()); - auto h = L/make_real3(n); + +// Spreads a group of particles onto a grid. +// The domain is such that some location r is r\in +-L +auto spreadParticles(thrust::device_vector &pos, + thrust::device_vector &values, int3 n, real sigma, + int supp, real3 L) { + L.z *= 2; + auto h = L / make_real3(n); auto kernel = std::make_shared(sigma, h.x, L.z, n.z, supp); using Grid = chebyshev::doublyperiodic::Grid; Grid grid(Box(L), n); IBM ibm(kernel, grid); - thrust::device_vector d_pos = pos; - auto pos_ptr = thrust::raw_pointer_cast(d_pos.data()); - thrust::device_vector d_values = values; - auto values_ptr = thrust::raw_pointer_cast(d_values.data()); - thrust::device_vector d_fr = fr; + auto pos_ptr = thrust::raw_pointer_cast(pos.data()); + auto values_ptr = thrust::raw_pointer_cast(values.data()); + thrust::device_vector d_fr(n.x * n.y * n.z); + thrust::fill(d_fr.begin(), d_fr.end(), complex()); auto fr_ptr = thrust::raw_pointer_cast(d_fr.data()); - ibm.spread(pos_ptr, (real2*)values_ptr, (real2*)fr_ptr, 1); + ibm.spread(pos_ptr, (real2 *)values_ptr, (real2 *)fr_ptr, pos.size()); + std::vector fr(n.x * n.y * n.z, complex()); thrust::copy(d_fr.begin(), d_fr.end(), fr.begin()); return fr; } -auto interpolateField(std::vector pos, std::vector field, - int3 n, real sigma, int supp, real3 L){ - std::vector values(pos.size(), complex()); - auto h = L/make_real3(n); +// Interpolates a discrete field into the locations of a group of particles. +// The domain is such that some location r is r\in +-L +auto interpolateField(thrust::device_vector &pos, + thrust::device_vector &field, int3 n, real sigma, + int supp, real3 L) { + L.z *= 2; + auto h = L / make_real3(n); auto kernel = std::make_shared(sigma, h.x, L.z, n.z, supp); using Grid = chebyshev::doublyperiodic::Grid; Grid grid(Box(L), n); IBM ibm(kernel, grid); - thrust::device_vector d_pos = pos; - auto pos_ptr = thrust::raw_pointer_cast(d_pos.data()); - thrust::device_vector d_values = values; + thrust::device_vector d_values(pos.size()); + thrust::fill(d_values.begin(), d_values.end(), complex{}); + auto pos_ptr = thrust::raw_pointer_cast(pos.data()); auto values_ptr = thrust::raw_pointer_cast(d_values.data()); - thrust::device_vector d_fr = field; - auto fr_ptr = thrust::raw_pointer_cast(d_fr.data()); + auto fr_ptr = thrust::raw_pointer_cast(field.data()); chebyshev::doublyperiodic::QuadratureWeights qw(L.z, h.x, h.y, n.z); IBM_ns::DefaultWeightCompute wc; - ibm.gather(pos_ptr, (real2*)values_ptr, (real2*)fr_ptr, qw, wc, 1); + ibm.gather(pos_ptr, (real2 *)values_ptr, (real2 *)fr_ptr, qw, wc, pos.size()); + std::vector values(pos.size(), complex()); thrust::copy(d_values.begin(), d_values.end(), values.begin()); return values; } diff --git a/test/misc/lanczos/CMakeLists.txt b/test/misc/lanczos/CMakeLists.txt index 1703d996..a7abd5f2 100644 --- a/test/misc/lanczos/CMakeLists.txt +++ b/test/misc/lanczos/CMakeLists.txt @@ -1,73 +1,11 @@ -cmake_minimum_required(VERSION 3.14) -project(lanczos_test) -enable_language(CUDA) -set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} "${CMAKE_SOURCE_DIR}/cmake/") -set(CMAKE_BUILD_TYPE Debug) -#set( CMAKE_VERBOSE_MAKEFILE on ) -#add_compile_definitions(PUBLIC MAXLOGLEVEL=15) -#GoogleTest requires at least C++14 -set(CMAKE_CXX_STANDARD 14) -set(CMAKE_CUDA_STANDARD 14) -set(CMAKE_CUDA_STANDARD_REQUIRED ON) -set(CMAKE_CUDA_SEPARABLE_COMPILATION OFF) -set(CUDA_ARCHITECTURES OFF) - -set(UAMMD_INCLUDE ../../../src ../../../src/third_party) - - -include(FetchContent) -FetchContent_Declare( - googletest - GIT_REPOSITORY https://github.com/google/googletest.git - GIT_TAG release-1.12.1 -) -# For Windows: Prevent overriding the parent project's compiler/linker settings -set(gtest_force_shared_crt ON CACHE BOOL "" FORCE) -FetchContent_MakeAvailable(googletest) -enable_testing() - -find_package(BLAS) -if(BLAS_FOUND) - message("mkl environment detected") - add_compile_definitions(PUBLIC USE_MKL) - link_libraries( - BLAS::BLAS - ) -else() - unset(BLA_VENDOR) - find_package(LAPACK REQUIRED) - find_package(LAPACKE REQUIRED) - find_package(BLAS REQUIRED) - link_libraries(${LAPACK_LIBRARIES} ${LAPACKE_LIBRARIES}) -endif() - -find_path(BLAS_INCLUDE_DIRS mkl.h - $ENV{CONDA_PREFIX}/include - /usr/include - /usr/local/include - $ENV{MKLROOT}/include - $ENV{BLAS_HOME}/include -) - -include_directories(${UAMMD_INCLUDE} ${BLAS_INCLUDE_DIRS}) -add_compile_definitions(PUBLIC DOUBLE_PRECISION) -link_libraries(${CUDA_LIBRARY} cufft cublas cusolver) - -add_executable( - lanczos_test - lanczos_test.cu -) - +add_executable(lanczos test_lanczos.cu) +target_include_directories(lanczos PUBLIC ${CMAKE_CURRENT_SOURCE_DIR}) +target_link_libraries(lanczos cufft cublas) +set_target_properties(lanczos PROPERTIES RUNTIME_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/bin") target_link_libraries( - lanczos_test + lanczos GTest::gtest_main GTest::gmock_main ) - include(GoogleTest) -gtest_discover_tests(lanczos_test) - -IF (CMAKE_BUILD_TYPE MATCHES "Debug") - set(CMAKE_CUDA_FLAGS "-g -G -lineinfo -src-in-ptx") - set(CMAKE_CXX_FLAGS "-g -Wall -Wextra -pedantic") -ENDIF() +gtest_discover_tests(lanczos) diff --git a/test/misc/lanczos/lanczos_test.cu b/test/misc/lanczos/test_lanczos.cu similarity index 100% rename from test/misc/lanczos/lanczos_test.cu rename to test/misc/lanczos/test_lanczos.cu