From a3f1bd335153c70e2290f9e041b6ed3aba27ab97 Mon Sep 17 00:00:00 2001
From: Matteo Concas <matteo.concas@cern.ch>
Date: Tue, 13 Aug 2024 17:40:38 +0200
Subject: [PATCH 1/2] Make it double

---
 GPU/Common/test/testSMatrixImp.cu | 25 ++++++++++++-------------
 1 file changed, 12 insertions(+), 13 deletions(-)

diff --git a/GPU/Common/test/testSMatrixImp.cu b/GPU/Common/test/testSMatrixImp.cu
index 1e1ba9b8cf6dc..c273a7ce11c8c 100644
--- a/GPU/Common/test/testSMatrixImp.cu
+++ b/GPU/Common/test/testSMatrixImp.cu
@@ -27,10 +27,12 @@
 #include <Math/SMatrix.h>
 #include <random>
 
-using MatSym3DGPU = o2::math_utils::SMatrixGPU<float, 3, 3, o2::math_utils::MatRepSymGPU<float, 3>>;
-using MatSym3D = ROOT::Math::SMatrix<float, 3, 3, ROOT::Math::MatRepSym<float, 3>>;
-using Mat3DGPU = o2::math_utils::SMatrixGPU<float, 3, 3, o2::math_utils::MatRepStdGPU<float, 3, 3>>;
-using Mat3D = ROOT::Math::SMatrix<float, 3, 3, ROOT::Math::MatRepStd<float, 3, 3>>;
+using MatSym3DGPU = o2::math_utils::SMatrixGPU<double, 3, 3, o2::math_utils::MatRepSymGPU<double, 3>>;
+using MatSym3D = ROOT::Math::SMatrix<double, 3, 3, ROOT::Math::MatRepSym<double, 3>>;
+using Mat3DGPU = o2::math_utils::SMatrixGPU<double, 3, 3, o2::math_utils::MatRepStdGPU<double, 3, 3>>;
+using Mat3D = ROOT::Math::SMatrix<double, 3, 3, ROOT::Math::MatRepStd<double, 3, 3>>;
+
+static constexpr double tolerance = 1e-8;
 
 #define GPU_CHECK(call)                                                                      \
   do {                                                                                       \
@@ -130,7 +132,7 @@ GPUg() void copyMatrixKernelArray(
 
 // Function to compare two matrices element-wise with a specified tolerance
 template <typename MatrixType>
-void compareMatricesElementWise(const MatrixType& mat1, const MatrixType& mat2, float tolerance)
+void compareMatricesElementWise(const MatrixType& mat1, const MatrixType& mat2, double tolerance)
 {
   auto tol = boost::test_tools::tolerance(tolerance);
 
@@ -228,7 +230,7 @@ struct GPUSMatrixImplFixtureSolo {
   {
     std::random_device rd;
     std::mt19937 gen(rd());
-    std::uniform_real_distribution<float> dis(1.0, 10.0);
+    std::uniform_real_distribution<double> dis(1.0, 10.0);
 
     // Initialize host matrices with random values
     for (int i = 0; i < 3; ++i) {
@@ -255,7 +257,6 @@ struct GPUSMatrixImplFixtureSolo {
 
 BOOST_FIXTURE_TEST_CASE(MatrixInversion, GPUSMatrixImplFixtureSolo)
 {
-  float tolerance = 0.00001f;
   const int nBlocks{1}, nThreads{1};
   GPUBenchmark benchmark("Single symmetric matrix inversion (" + std::to_string(nBlocks) + " blocks, " + std::to_string(nThreads) + " threads)");
   benchmark.start();
@@ -290,7 +291,7 @@ BOOST_FIXTURE_TEST_CASE(MatrixInversion, GPUSMatrixImplFixtureSolo)
   identity(2, 2) = 1;
   auto operation = SMatrix_h * SMatrix_original_h;
   Mat3D result;
-  ROOT::Math::Assign<float, 3, 3, decltype(operation), ROOT::Math::MatRepStd<float, 3, 3>, ROOT::Math::MatRepStd<float, 3, 3>>::Evaluate(result, operation);
+  ROOT::Math::Assign<double, 3, 3, decltype(operation), ROOT::Math::MatRepStd<double, 3, 3>, ROOT::Math::MatRepStd<double, 3, 3>>::Evaluate(result, operation);
   compareMatricesElementWise(result, identity, tolerance);
 }
 
@@ -307,7 +308,7 @@ struct GPUSMatrixImplFixtureDuo {
   {
     std::random_device rd;
     std::mt19937 gen(rd());
-    std::uniform_real_distribution<float> dis(1.0, 10.0);
+    std::uniform_real_distribution<double> dis(1.0, 10.0);
 
     // Initialize host matrices with random values
     for (int i = 0; i < 3; ++i) {
@@ -384,7 +385,7 @@ struct GPUSmatrixImplFixtureSoloArray {
   {
     std::random_device rd;
     std::mt19937 gen(rd());
-    std::uniform_real_distribution<float> dis(1.0, 10.0);
+    std::uniform_real_distribution<double> dis(1.0, 10.0);
 
     // Initialize host matrices with random values
     for (size_t iMatrix{0}; iMatrix < D; ++iMatrix) {
@@ -414,9 +415,7 @@ struct GPUSmatrixImplFixtureSoloArray {
 
 BOOST_FIXTURE_TEST_CASE(MatrixInversionArray, GPUSmatrixImplFixtureSoloArray<1'000'000>)
 {
-  float tolerance = 0.00001f;
-  const int nBlocks{20}, nThreads{512};
-
+   const int nBlocks{20}, nThreads{512};
   GPUBenchmark benchmark("Array of 1'000'000 symmetric matrices inversion (" + std::to_string(nBlocks) + " blocks, " + std::to_string(nThreads) + " threads)");
   benchmark.start();
   gpu::invertMatrixKernelArray<MatSym3DGPU><<<nBlocks, nThreads>>>(static_cast<MatSym3DGPU*>(SMatrixSymArray_d.get()), 1'000'000);

From 297a95e61fa13da9ffde6c5c9e75048e4de3af0e Mon Sep 17 00:00:00 2001
From: ALICE Builder <alibuild@users.noreply.github.com>
Date: Tue, 13 Aug 2024 17:51:06 +0200
Subject: [PATCH 2/2] Please consider the following formatting changes (#94)

---
 GPU/Common/test/testSMatrixImp.cu | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/GPU/Common/test/testSMatrixImp.cu b/GPU/Common/test/testSMatrixImp.cu
index c273a7ce11c8c..4be45c145ce60 100644
--- a/GPU/Common/test/testSMatrixImp.cu
+++ b/GPU/Common/test/testSMatrixImp.cu
@@ -415,7 +415,7 @@ struct GPUSmatrixImplFixtureSoloArray {
 
 BOOST_FIXTURE_TEST_CASE(MatrixInversionArray, GPUSmatrixImplFixtureSoloArray<1'000'000>)
 {
-   const int nBlocks{20}, nThreads{512};
+  const int nBlocks{20}, nThreads{512};
   GPUBenchmark benchmark("Array of 1'000'000 symmetric matrices inversion (" + std::to_string(nBlocks) + " blocks, " + std::to_string(nThreads) + " threads)");
   benchmark.start();
   gpu::invertMatrixKernelArray<MatSym3DGPU><<<nBlocks, nThreads>>>(static_cast<MatSym3DGPU*>(SMatrixSymArray_d.get()), 1'000'000);