Merge pull request #16 from cppalliance/CUDA_MD5

Add NVCC testing to MD5 and SHA1

.github/workflows/cuda.yml

@@ -0,0 +1,77 @@
+# Copyright 2024 Matt Borland
+# Distributed under the Boost Software License, Version 1.0.
+# (See accompanying file LICENSE_1_0.txt or copy at http://boost.org/LICENSE_1_0.txt)
+
+name: cuda
+on:
+  push:
+    branches:
+      - master
+      - develop
+      - feature/**
+  pull_request:
+  release:
+    types: [published, created, edited]
+
+concurrency:
+  group: ${{ github.head_ref || github.run_id }}
+  cancel-in-progress: true
+
+jobs:
+  cuda-cmake-test:
+    strategy:
+      fail-fast: false
+
+    runs-on: gpu-runner-1
+
+    steps:
+      - uses: Jimver/[email protected]
+        id: cuda-toolkit
+        with:
+          cuda: '12.5.0'
+          method: 'network'
+          sub-packages: '["nvcc"]'
+
+      - name: Output CUDA information
+        run: |
+          echo "Installed cuda version is: ${{steps.cuda-toolkit.outputs.cuda}}"+
+          echo "Cuda install location: ${{steps.cuda-toolkit.outputs.CUDA_PATH}}"
+          nvcc -V
+      - uses: actions/checkout@v4
+
+      - name: Install Packages
+        run: |
+          sudo apt-get install -y cmake make
+      - name: Setup Boost
+        run: |
+          echo GITHUB_REPOSITORY: $GITHUB_REPOSITORY
+          LIBRARY=${GITHUB_REPOSITORY#*/}
+          echo LIBRARY: $LIBRARY
+          echo "LIBRARY=$LIBRARY" >> $GITHUB_ENV
+          echo GITHUB_BASE_REF: $GITHUB_BASE_REF
+          echo GITHUB_REF: $GITHUB_REF
+          REF=${GITHUB_BASE_REF:-$GITHUB_REF}
+          REF=${REF#refs/heads/}
+          echo REF: $REF
+          BOOST_BRANCH=develop && [ "$REF" == "master" ] && BOOST_BRANCH=master || true
+          echo BOOST_BRANCH: $BOOST_BRANCH
+          cd ..
+          git clone -b $BOOST_BRANCH --depth 1 https://github.com/boostorg/boost.git boost-root
+          cd boost-root
+          mkdir -p libs/$LIBRARY
+          cp -r $GITHUB_WORKSPACE/* libs/$LIBRARY
+          git submodule update --init tools/boostdep
+          python3 tools/boostdep/depinst/depinst.py --git_args "--jobs 3" $LIBRARY
+      - name: Configure
+        run: |
+          cd ../boost-root
+          mkdir __build__ && cd __build__
+          cmake -DBOOST_INCLUDE_LIBRARIES=$LIBRARY -DBUILD_TESTING=ON -DCMAKE_CUDA_COMPILER=/usr/local/cuda/bin/nvcc -DBOOST_CRYPT_ENABLE_CUDA=1 -DCMAKE_CUDA_ARCHITECTURES=70 -DCUDA_TOOLKIT_ROOT_DIR=/usr/local/cuda-12.5 ..
+      - name: Build tests
+        run: |
+          cd ../boost-root/__build__
+          cmake --build . --target tests -j $(nproc)
+      - name: Run tests
+        run: |
+          cd ../boost-root/__build__
+          ctest --output-on-failure --no-tests=error

include/boost/crypt/hash/sha1.hpp

@@ -44,7 +44,7 @@ class sha1_hasher
     bool computed {};
     bool corrupted {};
 
-    constexpr auto sha1_process_message_block() -> void;
+    BOOST_CRYPT_GPU_ENABLED constexpr auto sha1_process_message_block() -> void;
 
     template <typename ForwardIter>
     BOOST_CRYPT_GPU_ENABLED constexpr auto sha1_update(ForwardIter data, boost::crypt::size_t size) noexcept -> hasher_state;
@@ -75,6 +75,7 @@ class sha1_hasher
 
 namespace detail {
 
+BOOST_CRYPT_GPU_ENABLED
 constexpr auto round1(boost::crypt::uint32_t& A,
                       boost::crypt::uint32_t& B,
                       boost::crypt::uint32_t& C,
@@ -90,6 +91,7 @@ constexpr auto round1(boost::crypt::uint32_t& A,
     A = temp;
 }
 
+BOOST_CRYPT_GPU_ENABLED
 constexpr auto round2(boost::crypt::uint32_t& A,
                       boost::crypt::uint32_t& B,
                       boost::crypt::uint32_t& C,
@@ -105,6 +107,7 @@ constexpr auto round2(boost::crypt::uint32_t& A,
     A = temp;
 }
 
+BOOST_CRYPT_GPU_ENABLED
 constexpr auto round3(boost::crypt::uint32_t& A,
                       boost::crypt::uint32_t& B,
                       boost::crypt::uint32_t& C,
@@ -120,6 +123,7 @@ constexpr auto round3(boost::crypt::uint32_t& A,
     A = temp;
 }
 
+BOOST_CRYPT_GPU_ENABLED
 constexpr auto round4(boost::crypt::uint32_t& A,
                       boost::crypt::uint32_t& B,
                       boost::crypt::uint32_t& C,

include/boost/crypt/utility/cstddef.hpp

@@ -16,8 +16,7 @@ namespace crypt {
 using size_t = unsigned long;
 using ptrdiff_t = long;
 using nullptr_t = void;
-using std::max_align_t = double;
-
+using max_align_t = double;
 
 } // namespace crypt
 } // namespace boost

include/boost/crypt/utility/iterator.hpp

@@ -13,13 +13,15 @@
 
 namespace boost {
 namespace crypt {
+namespace utility {
 
 template <typename Iter>
 struct iterator_traits : public cuda::std::iterator_traits<Iter> {};
 
 template <typename T>
 struct iterator_traits<T*> : public cuda::std::iterator_traits<T*> {};
 
+} // namespace utility
 } // namespace crypt
 } // namespace boost
 

test/CMakeLists.txt

@@ -7,6 +7,19 @@ include(BoostTestJamfile OPTIONAL RESULT_VARIABLE HAVE_BOOST_TEST)
 
 if(HAVE_BOOST_TEST)
 
-    boost_test_jamfile(FILE Jamfile LINK_LIBRARIES Boost::crypt Boost::core Boost::uuid)
+    if (BOOST_CRYPT_ENABLE_CUDA)
+
+        message(STATUS "Building boost.crypt with CUDA")
+        find_package(CUDA REQUIRED)
+        enable_language(CUDA)
+        set(CMAKE_CUDA_EXTENSIONS OFF)
+
+        boost_test_jamfile(FILE nvcc_jamfile LINK_LIBRARIES Boost::crypt ${CUDA_LIBRARIES} INCLUDE_DIRECTORIES ${CUDA_INCLUDE_DIRS} )
+
+    else ()
+
+        boost_test_jamfile(FILE Jamfile LINK_LIBRARIES Boost::crypt Boost::core Boost::uuid)
+
+    endif ()
 
 endif()

test/cuda_managed_ptr.hpp

@@ -0,0 +1,140 @@
+
+//  Copyright John Maddock 2016.
+//  Use, modification and distribution are subject to the
+//  Boost Software License, Version 1.0. (See accompanying file
+//  LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt)
+
+#ifndef BOOST_MATH_CUDA_MANAGED_PTR_HPP
+#define BOOST_MATH_CUDA_MANAGED_PTR_HPP
+
+#ifdef _MSC_VER
+#pragma once
+#endif
+
+#include <iostream>
+#include <cuda_runtime.h>
+
+class managed_holder_base
+{
+protected:
+   static int count;
+   managed_holder_base() { ++count; }
+   ~managed_holder_base()
+   {
+      if(0 == --count)
+         cudaDeviceSynchronize();
+   }
+};
+
+int managed_holder_base::count = 0;
+
+//
+// Reset the device and exit:
+// cudaDeviceReset causes the driver to clean up all state. While
+// not mandatory in normal operation, it is good practice.  It is also
+// needed to ensure correct operation when the application is being
+// profiled. Calling cudaDeviceReset causes all profile data to be
+// flushed before the application exits.
+//
+// We have a global instance of this class, plus instances for each
+// managed pointer.  Last one out the door switches the lights off.
+//
+class cudaResetter
+{
+   static int count;
+public:
+   cudaResetter() { ++count;  }
+   ~cudaResetter()
+   {
+      if(--count == 0)
+      {
+         cudaError_t err = cudaDeviceReset();
+         if(err != cudaSuccess)
+         {
+            std::cerr << "Failed to deinitialize the device! error=" << cudaGetErrorString(err) << std::endl;
+         }
+      }
+   }
+};
+
+int cudaResetter::count = 0;
+
+cudaResetter global_resetter;
+
+template <class T>
+class cuda_managed_ptr
+{
+   T* data;
+   static const cudaResetter resetter;
+   cuda_managed_ptr(const cuda_managed_ptr&) = delete;
+   cuda_managed_ptr& operator=(cuda_managed_ptr const&) = delete;
+   void free()
+   {
+      if(data)
+      {
+         cudaDeviceSynchronize();
+         cudaError_t err = cudaFree(data);
+         if(err != cudaSuccess)
+         {
+            std::cerr << "Failed to deinitialize the device! error=" << cudaGetErrorString(err) << std::endl;
+         }
+      }
+   }
+public:
+   cuda_managed_ptr() : data(0) {}
+   cuda_managed_ptr(std::size_t n)
+   {
+      cudaError_t err = cudaSuccess;
+      void *ptr;
+      err = cudaMallocManaged(&ptr, n * sizeof(T));
+      if(err != cudaSuccess)
+         throw std::runtime_error(cudaGetErrorString(err));
+      cudaDeviceSynchronize();
+      data = static_cast<T*>(ptr);
+   }
+   cuda_managed_ptr(cuda_managed_ptr&& o)
+   {
+      data = o.data;
+      o.data = 0;
+   }
+   cuda_managed_ptr& operator=(cuda_managed_ptr&& o)
+   {
+      free();
+      data = o.data;
+      o.data = 0;
+      return *this;
+   }
+   ~cuda_managed_ptr()
+   {
+      free();
+   }
+
+   class managed_holder : managed_holder_base
+   {
+      T* pdata;
+   public:
+      managed_holder(T* p) : managed_holder_base(), pdata(p) {}
+      managed_holder(const managed_holder& o) : managed_holder_base(), pdata(o.pdata) {}
+      operator T* () { return pdata; }
+      T& operator[] (std::size_t n) { return pdata[n]; }
+   };
+   class const_managed_holder : managed_holder_base
+   {
+      const T* pdata;
+   public:
+      const_managed_holder(T* p) : managed_holder_base(), pdata(p) {}
+      const_managed_holder(const managed_holder& o) : managed_holder_base(), pdata(o.pdata) {}
+      operator const T* () { return pdata; }
+      const T& operator[] (std::size_t n) { return pdata[n]; }
+   };
+
+   managed_holder get() { return managed_holder(data); }
+   const_managed_holder get()const { return data; }
+   T& operator[](std::size_t n) { return data[n]; }
+   const T& operator[](std::size_t n)const { return data[n]; }
+};
+
+template <class T>
+cudaResetter const cuda_managed_ptr<T>::resetter;
+
+#endif

test/nvcc_jamfile

@@ -0,0 +1,13 @@
+# Copyright 2024 Matt Borland
+# Distributed under the Boost Software License, Version 1.0.
+# https://www.boost.org/LICENSE_1_0.txt
+
+import testing ;
+import ../../config/checks/config : requires ;
+
+project : requirements
+    [ requires cxx14_decltype_auto cxx14_generic_lambdas cxx14_return_type_deduction cxx14_variable_templates cxx14_constexpr ]
+    ;
+
+run test_md5_nvcc.cu ;
+run test_sha1_nvcc.cu ;

test/stopwatch.hpp

@@ -0,0 +1,39 @@
+//  Copyright John Maddock 2016.
+//  Use, modification and distribution are subject to the
+//  Boost Software License, Version 1.0. (See accompanying file
+//  LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt)
+
+#ifndef BOOST_MATH_CUDA_STOPWATCH_HPP
+#define BOOST_MATH_CUDA_STOPWATCH_HPP
+
+#ifdef _MSC_VER
+#pragma once
+#endif
+
+#include <chrono>
+
+template <class Clock>
+struct stopwatch
+{
+    typedef typename Clock::duration duration;
+    stopwatch()
+    {
+        m_start = Clock::now();
+    }
+    double elapsed()
+    {
+        duration t = Clock::now() - m_start;
+        return std::chrono::duration_cast<std::chrono::duration<double>>(t).count();
+    }
+    void reset()
+    {
+        m_start = Clock::now();
+    }
+
+private:
+    typename Clock::time_point m_start;
+};
+
+typedef stopwatch<std::chrono::high_resolution_clock> watch;
+
+#endif