CUDA C++ 20

CMakeLists.txt

@@ -52,7 +52,7 @@ if(SPH_EXA_WITH_CUDA)
     if(CMAKE_CUDA_COMPILER)
         enable_language(CUDA)
         find_package(CUDAToolkit)
-        set(CMAKE_CUDA_STANDARD 17)
+        set(CMAKE_CUDA_STANDARD 20)
     else()
         message(STATUS "No CUDA support")
         set(SPH_EXA_WITH_CUDA OFF)
@@ -64,7 +64,7 @@ if(SPH_EXA_WITH_HIP)
     if(CMAKE_HIP_COMPILER AND NOT CMAKE_CUDA_COMPILER)
         enable_language(HIP)
         find_package(hip)
-        set(CMAKE_HIP_STANDARD 17)
+        set(CMAKE_HIP_STANDARD 20)
     else()
         message(STATUS "No HIP support")
         set(SPH_EXA_WITH_HIP OFF)

ci/cuda12/run_debug_sm90.yml

@@ -24,7 +24,6 @@ sph-run-cuda12-sm90-debug-1rank:
     - sbin/hydro/sph_tests
     - sbin/performance/octree_perf
     - sbin/performance/peers_perf
-    - sbin/performance/scan_perf
     - sbin/ryoanji/cpu_unit_tests/ryoanji_cpu_unit_tests
     - sbin/ryoanji/unit_tests/ryoanji_unit_tests
     # TODO:slow    - sbin/performance/hilbert_perf

ci/cuda12/run_release_sm90.yml

@@ -24,7 +24,6 @@ sph-run-cuda12-sm90-release-1rank:
     - sbin/hydro/sph_tests
     - sbin/performance/octree_perf
     - sbin/performance/peers_perf
-    - sbin/performance/scan_perf
     - sbin/ryoanji/cpu_unit_tests/ryoanji_cpu_unit_tests
     - sbin/ryoanji/unit_tests/ryoanji_unit_tests
     # TODO:slow    - sbin/performance/hilbert_perf

domain/CMakeLists.txt

@@ -40,7 +40,7 @@ if(CSTONE_WITH_CUDA)
     if(CMAKE_CUDA_COMPILER)
         enable_language(CUDA)
         find_package(CUDAToolkit)
-        set(CMAKE_CUDA_STANDARD 17)
+        set(CMAKE_CUDA_STANDARD 20)
     else()
         message(STATUS "No CUDA support")
         set(CSTONE_WITH_CUDA OFF)
@@ -52,7 +52,7 @@ if(CSTONE_WITH_HIP)
     if(CMAKE_HIP_COMPILER)
         enable_language(HIP)
         find_package(hip)
-        set(CMAKE_HIP_STANDARD 17)
+        set(CMAKE_HIP_STANDARD 20)
     else()
         message(STATUS "No HIP support")
         set(CSTONE_WITH_HIP OFF)

domain/LICENSE

@@ -1,7 +1,6 @@
 MIT License
 
-Copyright (c) 2021 CSCS, ETH Zurich
-              2021 University of Basel
+Copyright (c) 2024 CSCS, ETH Zurich
 
 Permission is hereby granted, free of charge, to any person obtaining a copy
 of this software and associated documentation files (the "Software"), to deal

domain/include/cstone/cuda/annotation.hpp

@@ -1,26 +1,10 @@
 /*
- * MIT License
+ * Cornerstone octree
  *
- * Copyright (c) 2021 CSCS, ETH Zurich
- *               2021 University of Basel
+ * Copyright (c) 2024 CSCS, ETH Zurich
  *
- * Permission is hereby granted, free of charge, to any person obtaining a copy
- * of this software and associated documentation files (the "Software"), to deal
- * in the Software without restriction, including without limitation the rights
- * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
- * copies of the Software, and to permit persons to whom the Software is
- * furnished to do so, subject to the following conditions:
- *
- * The above copyright notice and this permission notice shall be included in all
- * copies or substantial portions of the Software.
- *
- * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
- * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
- * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
- * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
- * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
- * SOFTWARE.
+ * Please, refer to the LICENSE file in the root directory.
+ * SPDX-License-Identifier: MIT License
  */
 
 /*! @file

domain/include/cstone/cuda/cub.hpp

@@ -1,3 +1,11 @@
+/*
+ * Cornerstone octree
+ *
+ * Copyright (c) 2024 CSCS, ETH Zurich
+ *
+ * Please, refer to the LICENSE file in the root directory.
+ * SPDX-License-Identifier: MIT License
+ */
 
 #pragma once
 

domain/include/cstone/cuda/cuda_runtime.hpp

@@ -1,3 +1,12 @@
+/*
+ * Cornerstone octree
+ *
+ * Copyright (c) 2024 CSCS, ETH Zurich
+ *
+ * Please, refer to the LICENSE file in the root directory.
+ * SPDX-License-Identifier: MIT License
+ */
+
 /*! @file
  * @brief CUDA/HIP runtime API compatiblity wrapper
  *

domain/include/cstone/cuda/cuda_stubs.h

@@ -1,26 +1,10 @@
 /*
- * MIT License
+ * Cornerstone octree
  *
- * Copyright (c) 2022 Politechnical University of Catalonia UPC
- *               2022 University of Basel
+ * Copyright (c) 2024 CSCS, ETH Zurich
  *
- * Permission is hereby granted, free of charge, to any person obtaining a copy
- * of this software and associated documentation files (the "Software"), to deal
- * in the Software without restriction, including without limitation the rights
- * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
- * copies of the Software, and to permit persons to whom the Software is
- * furnished to do so, subject to the following conditions:
- *
- * The above copyright notice and this permission notice shall be included in all
- * copies or substantial portions of the Software.
- *
- * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
- * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
- * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
- * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
- * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
- * SOFTWARE.
+ * Please, refer to the LICENSE file in the root directory.
+ * SPDX-License-Identifier: MIT License
  */
 
 /*! @file
@@ -56,7 +40,7 @@ void memcpyD2D(const T* src, std::size_t n, T* dest);
 
 void syncGpu();
 
-/*! @brief detection trait to determine whether a template parameter is an instance of thrust::device_vector
+/*! @brief detection trait to determine whether a template parameter is a device vector
  *
  * @tparam Vector the Vector type to check
  *

domain/include/cstone/cuda/cuda_utils.cuh

@@ -1,3 +1,12 @@
+/*
+ * Cornerstone octree
+ *
+ * Copyright (c) 2024 CSCS, ETH Zurich
+ *
+ * Please, refer to the LICENSE file in the root directory.
+ * SPDX-License-Identifier: MIT License
+ */
+
 /*! @file
  * @brief CUDA runtime API wrapper for compatiblity with CPU code
  *

domain/include/cstone/cuda/cuda_utils.hpp

@@ -1,26 +1,10 @@
 /*
- * MIT License
+ * Cornerstone octree
  *
- * Copyright (c) 2021 CSCS, ETH Zurich
- *               2021 University of Basel
+ * Copyright (c) 2024 CSCS, ETH Zurich
  *
- * Permission is hereby granted, free of charge, to any person obtaining a copy
- * of this software and associated documentation files (the "Software"), to deal
- * in the Software without restriction, including without limitation the rights
- * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
- * copies of the Software, and to permit persons to whom the Software is
- * furnished to do so, subject to the following conditions:
- *
- * The above copyright notice and this permission notice shall be included in all
- * copies or substantial portions of the Software.
- *
- * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
- * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
- * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
- * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
- * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
- * SOFTWARE.
+ * Please, refer to the LICENSE file in the root directory.
+ * SPDX-License-Identifier: MIT License
  */
 
 #pragma once

domain/include/cstone/cuda/device_vector.cu

@@ -1,3 +1,12 @@
+/*
+ * Cornerstone octree
+ *
+ * Copyright (c) 2024 CSCS, ETH Zurich
+ *
+ * Please, refer to the LICENSE file in the root directory.
+ * SPDX-License-Identifier: MIT License
+ */
+
 /*! @file
  * @brief  Encapsulation around a std::vector like GPU-resident device vector, suitable for use in .cpp files
  *

domain/include/cstone/cuda/device_vector.h

@@ -1,3 +1,12 @@
+/*
+ * Cornerstone octree
+ *
+ * Copyright (c) 2024 CSCS, ETH Zurich
+ *
+ * Please, refer to the LICENSE file in the root directory.
+ * SPDX-License-Identifier: MIT License
+ */
+
 /*! @file
  * @brief  Encapsulation around a std::vector like GPU-resident device vector, suitable for use in .cpp files
  *

domain/include/cstone/cuda/errorcheck.cuh

@@ -1,25 +1,10 @@
 /*
- * MIT License
+ * Cornerstone octree
  *
- * Copyright (c) 2022 CSCS, ETH Zurich
+ * Copyright (c) 2024 CSCS, ETH Zurich
  *
- * Permission is hereby granted, free of charge, to any person obtaining a copy
- * of this software and associated documentation files (the "Software"), to deal
- * in the Software without restriction, including without limitation the rights
- * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
- * copies of the Software, and to permit persons to whom the Software is
- * furnished to do so, subject to the following conditions:
- *
- * The above copyright notice and this permission notice shall be included in all
- * copies or substantial portions of the Software.
- *
- * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
- * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
- * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
- * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
- * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
- * SOFTWARE.
+ * Please, refer to the LICENSE file in the root directory.
+ * SPDX-License-Identifier: MIT License
  */
 
 #pragma once

domain/include/cstone/cuda/gpu_config.cuh

@@ -1,26 +1,10 @@
 /*
- * MIT License
+ * Cornerstone octree
  *
- * Copyright (c) 2022 CSCS, ETH Zurich
- *               2021 University of Basel
+ * Copyright (c) 2024 CSCS, ETH Zurich
  *
- * Permission is hereby granted, free of charge, to any person obtaining a copy
- * of this software and associated documentation files (the "Software"), to deal
- * in the Software without restriction, including without limitation the rights
- * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
- * copies of the Software, and to permit persons to whom the Software is
- * furnished to do so, subject to the following conditions:
- *
- * The above copyright notice and this permission notice shall be included in all
- * copies or substantial portions of the Software.
- *
- * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
- * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
- * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
- * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
- * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
- * SOFTWARE.
+ * Please, refer to the LICENSE file in the root directory.
+ * SPDX-License-Identifier: MIT License
  */
 
 /*! @file

domain/include/cstone/cuda/thrust_util.cuh

@@ -1,3 +1,12 @@
+/*
+ * Cornerstone octree
+ *
+ * Copyright (c) 2024 CSCS, ETH Zurich
+ *
+ * Please, refer to the LICENSE file in the root directory.
+ * SPDX-License-Identifier: MIT License
+ */
+
 /*! @file
  * @brief Utilities for Thrust device vectors for use in .cu translation units only
  */

domain/include/cstone/domain/assignment.hpp

@@ -1,7 +1,7 @@
 /*
  * Cornerstone octree
  *
- * Copyright (c) 2024 CSCS, ETH Zurich, University of Zurich, 2021 University of Basel
+ * Copyright (c) 2024 CSCS, ETH Zurich
  *
  * Please, refer to the LICENSE file in the root directory.
  * SPDX-License-Identifier: MIT License
@@ -82,18 +82,18 @@ class GlobalAssignment
         else { box_ = limitBoxShrinking(fittingBox, box_); }
 
         // compute SFC particle keys only for particles participating in tree build
-        gsl::span<KeyType> keyView(particleKeys + bufDesc.start, numParticles);
+        std::span<KeyType> keyView(particleKeys + bufDesc.start, numParticles);
         computeSfcKeys(x + bufDesc.start, y + bufDesc.start, z + bufDesc.start, sfcKindPointer(keyView.data()),
                        numParticles, box_);
 
         // sort keys and keep track of ordering for later use
-        reorderFunctor.setMapFromCodes(keyView.begin(), keyView.end());
+        reorderFunctor.setMapFromCodes(keyView);
 
-        updateOctreeGlobal(keyView.begin(), keyView.end(), bucketSize_, tree_, nodeCounts_);
+        updateOctreeGlobal<KeyType>(keyView, bucketSize_, tree_, nodeCounts_);
         if (firstCall_)
         {
             firstCall_ = false;
-            while (!updateOctreeGlobal(keyView.begin(), keyView.end(), bucketSize_, tree_, nodeCounts_))
+            while (!updateOctreeGlobal<KeyType>(keyView, bucketSize_, tree_, nodeCounts_))
                 ;
         }
 
@@ -142,7 +142,7 @@ class GlobalAssignment
 
         auto [newStart, newEnd] = domain_exchange::assignedEnvelope(bufDesc, numPresent(), numAssigned());
         LocalIndex envelopeSize = newEnd - newStart;
-        gsl::span<KeyType> keyView(keys + newStart, envelopeSize);
+        std::span<KeyType> keyView(keys + newStart, envelopeSize);
 
         auto recvStart = domain_exchange::receiveStart(bufDesc, numPresent(), numAssigned());
         auto numRecv   = numAssigned() - numPresent();
@@ -153,7 +153,7 @@ class GlobalAssignment
         reorderFunctor.extendMap(shifts, s1);
 
         // sort keys and keep track of the ordering
-        reorderFunctor.updateMap(keyView.begin(), keyView.end());
+        reorderFunctor.updateMap(keyView);
 
         return std::make_tuple(newStart, keyView.subspan(numSendDown(), numAssigned()));
     }
@@ -170,11 +170,11 @@ class GlobalAssignment
     }
 
     //! @brief read only visibility of the global octree leaves to the outside
-    gsl::span<const KeyType> treeLeaves() const { return tree_.treeLeaves(); }
+    std::span<const KeyType> treeLeaves() const { return tree_.treeLeaves(); }
     //! @brief the octree, including the internal part
     const Octree<KeyType>& octree() const { return tree_; }
     //! @brief read only visibility of the global octree leaf counts to the outside
-    gsl::span<const unsigned> nodeCounts() const { return nodeCounts_; }
+    std::span<const unsigned> nodeCounts() const { return nodeCounts_; }
     //! @brief the global coordinate bounding box
     const Box<T>& box() const { return box_; }
     //! @brief return the space filling curve rank assignment of the last call to @a assign()