From 2e0888f7055021544f11feec9928db9266bf02ca Mon Sep 17 00:00:00 2001
From: Matt Borland <matt@mattborland.com>
Date: Thu, 9 Jan 2025 14:13:27 -0500
Subject: [PATCH] Update NVCC SHA256 testing

---
 test/nvcc_jamfile        |  2 +-
 test/test_sha256_nvcc.cu | 16 +++++++++-------
 2 files changed, 10 insertions(+), 8 deletions(-)

diff --git a/test/nvcc_jamfile b/test/nvcc_jamfile
index c0a5d679..bafcaf2b 100644
--- a/test/nvcc_jamfile
+++ b/test/nvcc_jamfile
@@ -11,7 +11,7 @@ project : requirements
 
 run test_sha1_nvcc.cu ;
 run test_sha224_nvcc.cu ;
-#run test_sha256_nvcc.cu ;
+run test_sha256_nvcc.cu ;
 #run test_sha384_nvcc.cu ;
 #run test_sha512_nvcc.cu ;
 #run test_sha512_224_nvcc.cu ;
diff --git a/test/test_sha256_nvcc.cu b/test/test_sha256_nvcc.cu
index d20d1414..f798b222 100644
--- a/test/test_sha256_nvcc.cu
+++ b/test/test_sha256_nvcc.cu
@@ -3,7 +3,8 @@
 //  Boost Software License, Version 1.0. (See accompanying file
 //  LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt)
 
-#include <boost/crypt/hash/sha256.hpp>
+#include <cuda_runtime.h>
+#include <boost/crypt2/hash/sha256.hpp>
 #include "cuda_managed_ptr.hpp"
 #include "stopwatch.hpp"
 #include "generate_random_strings.hpp"
@@ -11,10 +12,9 @@
 #include <iomanip>
 #include <exception>
 #include <memory>
+#include <span>
 
-#include <cuda_runtime.h>
-
-using digest_type = boost::crypt::sha256_hasher::return_type;
+using digest_type = cuda::std::array<cuda::std::byte, 32>;
 
 // The kernel function
 __global__ void cuda_test(char** in, digest_type* out, int numElements)
@@ -23,7 +23,8 @@ __global__ void cuda_test(char** in, digest_type* out, int numElements)
 
     if (i < numElements)
     {
-        out[i] = boost::crypt::sha256(in[i]);
+        auto in_span {cuda::std::span(in[i], 64)};
+        out[i] = boost::crypt::sha256(in_span);
     }
 }
 
@@ -59,7 +60,7 @@ int main()
 
         // Launch the Vector Add CUDA Kernel
         int threadsPerBlock = 256;
-        int blocksPerGrid =(numElements + threadsPerBlock - 1) / threadsPerBlock;
+        int blocksPerGrid = (numElements + threadsPerBlock - 1) / threadsPerBlock;
         std::cout << "CUDA kernel launch with " << blocksPerGrid << " blocks of " << threadsPerBlock << " threads" << std::endl;
 
         watch w;
@@ -80,7 +81,8 @@ int main()
         w.reset();
         for(int i = 0; i < numElements; ++i)
         {
-           results.emplace_back(boost::crypt::sha256(input_vector1[i]));
+            std::span<char> in(input_vector1[i], elementSize);
+            results.emplace_back(boost::crypt::sha256(in));
         }
         double t = w.elapsed();