From 67fc871536dae40a3f364f6b549859f6eb11f758 Mon Sep 17 00:00:00 2001 From: Ruairi Moran Date: Fri, 2 Aug 2024 18:48:03 +0100 Subject: [PATCH 1/3] added min and max methods and tests, updated changelog --- CHANGELOG.md | 10 +++++++- include/tensor.cuh | 58 +++++++++++++++++++++++++++++++++++++++++++--- test/testTensor.cu | 36 +++++++++++++++++++++++++++- 3 files changed, 99 insertions(+), 5 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index 88d0a63..3096100 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -6,6 +6,15 @@ The format is based on [Keep a Changelog](https://keepachangelog.com/en/1.1.0/), and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0.html). + +## v1.1.0 - 03-08-2024 + +### Added + +- Implementation and test of methods `.max()` and `.min()` for any tensor. + @@ -21,7 +30,6 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0 - Using a function `numBlocks` instead of the macro `DIM2BLOCKS` - Using `TEMPLATE_WITH_TYPE_T` and `TEMPLATE_CONSTRAINT_REQUIRES_FPX` for the code to run on both C++17 and C++20 - diff --git a/include/tensor.cuh b/include/tensor.cuh index fb50e4a..a2698cb 100644 --- a/include/tensor.cuh +++ b/include/tensor.cuh @@ -377,6 +377,19 @@ public: */ T sumAbs() const; + /** + * Maximum of absolute of all elements. + * Equivalent to inf-norm, max(|x_i|) for all i. + * @return max as same data type + */ + T max() const; + + /** + * Minimum of absolute of all elements, min(|x_i|) for all i. + * @return min as same data type + */ + T min() const; + /** * Solves for the least squares solution of A \ b. * A is this tensor and b is the provided tensor. @@ -405,7 +418,7 @@ public: DTensor &operator=(const DTensor &other); - T operator()(size_t i, size_t j = 0, size_t k = 0); + T operator()(size_t i, size_t j = 0, size_t k = 0) const; DTensor &operator*=(T scalar); @@ -605,7 +618,6 @@ inline float DTensor::normF() const { return the_norm; } - template<> inline float DTensor::sumAbs() const { float sumAbsAllElements; @@ -622,6 +634,46 @@ inline double DTensor::sumAbs() const { return sumAbsAllElements; } +template<> +inline float DTensor::max() const { + int idx; + float hostDst; + gpuErrChk(cublasIsamax(Session::getInstance().cuBlasHandle(), m_numRows * m_numCols * m_numMats, m_d_data, 1, + &idx)); + gpuErrChk(cudaMemcpy(&hostDst, m_d_data + idx - 1, sizeof(float), cudaMemcpyDeviceToHost)); + return std::signbit(hostDst) ? -hostDst : hostDst; +} + +template<> +inline double DTensor::max() const { + int idx; + double hostDst; + gpuErrChk(cublasIdamax(Session::getInstance().cuBlasHandle(), m_numRows * m_numCols * m_numMats, m_d_data, 1, + &idx)); + gpuErrChk(cudaMemcpy(&hostDst, m_d_data + idx - 1, sizeof(double), cudaMemcpyDeviceToHost)); + return std::signbit(hostDst) ? -hostDst : hostDst; +} + +template<> +inline float DTensor::min() const { + int idx; + float hostDst; + gpuErrChk(cublasIsamin(Session::getInstance().cuBlasHandle(), m_numRows * m_numCols * m_numMats, m_d_data, 1, + &idx)); + gpuErrChk(cudaMemcpy(&hostDst, m_d_data + idx - 1, sizeof(float), cudaMemcpyDeviceToHost)); + return std::signbit(hostDst) ? -hostDst : hostDst; +} + +template<> +inline double DTensor::min() const { + int idx; + double hostDst; + gpuErrChk(cublasIdamin(Session::getInstance().cuBlasHandle(), m_numRows * m_numCols * m_numMats, m_d_data, 1, + &idx)); + gpuErrChk(cudaMemcpy(&hostDst, m_d_data + idx - 1, sizeof(double), cudaMemcpyDeviceToHost)); + return std::signbit(hostDst) ? -hostDst : hostDst; +} + template inline bool DTensor::allocateOnDevice(size_t size, bool zero) { if (size <= 0) return false; @@ -772,7 +824,7 @@ inline DTensor &DTensor::operator-=(const DTensor &rhs) } template -inline T DTensor::operator()(size_t i, size_t j, size_t k) { +inline T DTensor::operator()(size_t i, size_t j, size_t k) const { T hostDst; size_t offset = i + m_numRows * (j + m_numCols * k); gpuErrChk(cudaMemcpy(&hostDst, m_d_data + offset, sizeof(T), cudaMemcpyDeviceToHost)); diff --git a/test/testTensor.cu b/test/testTensor.cu index 3d1a7b4..df81014 100644 --- a/test/testTensor.cu +++ b/test/testTensor.cu @@ -352,11 +352,45 @@ void tensorSumAbs() { EXPECT_NEAR(112, tenz.sumAbs(), PRECISION_HIGH); // from MATLAB } -TEST_F(TensorTest, tensorNormFtensorSumAbs) { +TEST_F(TensorTest, tensorSumAbs) { tensorSumAbs(); tensorSumAbs(); } +/* --------------------------------------- + * Tensor: max of all elements + * --------------------------------------- */ + +TEMPLATE_WITH_TYPE_T +void tensorMax() { + std::vector data = TENSOR_DATA_234AMB; + DTensor tenz(data, 2, 3, 4); + T m = tenz.max(); + EXPECT_EQ(27, m); +} + +TEST_F(TensorTest, tensorMax) { + tensorMax(); + tensorMax(); +} + +/* --------------------------------------- + * Tensor: min of all elements + * --------------------------------------- */ + +TEMPLATE_WITH_TYPE_T +void tensorMin() { + std::vector data = TENSOR_DATA_234AMB; + DTensor tenz(data, 2, 3, 4); + T m = tenz.min(); + EXPECT_EQ(0, m); +} + +TEST_F(TensorTest, tensorMin) { + tensorMin(); + tensorMin(); +} + /* --------------------------------------- * Tensor operator() to access element * e.g., t(2, 3, 4) From 6c370171a4b5716e57f73215f60337c7298580b9 Mon Sep 17 00:00:00 2001 From: Ruairi Moran Date: Fri, 2 Aug 2024 19:14:49 +0100 Subject: [PATCH 2/3] rename --- CHANGELOG.md | 2 +- include/tensor.cuh | 16 ++++++++-------- test/testTensor.cu | 8 ++++---- 3 files changed, 13 insertions(+), 13 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index 3096100..2ba5842 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -13,7 +13,7 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0 ### Added -- Implementation and test of methods `.max()` and `.min()` for any tensor. +- Implementation and test of methods `.maxAbs()` and `.minAbs()` for any tensor.