diff --git a/library/src/reduction/CMakeLists.txt b/library/src/reduction/CMakeLists.txt
index 05f2d5e9..1d2949d9 100644
--- a/library/src/reduction/CMakeLists.txt
+++ b/library/src/reduction/CMakeLists.txt
@@ -30,12 +30,53 @@ get_target_property(composable_kernel_INCLUDES composable_kernel::device_reducti
 get_target_property(composable_kernel_INCLUDES composable_kernel::device_other_operations INTERFACE_INCLUDE_DIRECTORIES)
 set(HIPTENSOR_REDUCTION_SOURCES
    ${CMAKE_CURRENT_SOURCE_DIR}/hiptensor_reduction.cpp
-   # ${CMAKE_CURRENT_SOURCE_DIR}/reduction_cpu_reference.cpp
-   # ${CMAKE_CURRENT_SOURCE_DIR}/reduction_selection.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_cpu_reference.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_cpu_reference_instances.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution_1_1_f32_f32_instance.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution_2_1_f32_f32_instance.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution_2_2_f32_f32_instance.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution_3_1_f32_f32_instance.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution_3_2_f32_f32_instance.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution_3_3_f32_f32_instance.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution_4_1_f32_f32_instance.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution_4_2_f32_f32_instance.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution_4_3_f32_f32_instance.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution_4_4_f32_f32_instance.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution_5_1_f32_f32_instance.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution_5_2_f32_f32_instance.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution_5_3_f32_f32_instance.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution_5_4_f32_f32_instance.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution_5_5_f32_f32_instance.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution_6_1_f32_f32_instance.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution_6_2_f32_f32_instance.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution_6_3_f32_f32_instance.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution_6_4_f32_f32_instance.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution_6_5_f32_f32_instance.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution_6_6_f32_f32_instance.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution_1_1_f64_f64_instance.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution_2_1_f64_f64_instance.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution_2_2_f64_f64_instance.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution_3_1_f64_f64_instance.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution_3_2_f64_f64_instance.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution_3_3_f64_f64_instance.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution_4_1_f64_f64_instance.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution_4_2_f64_f64_instance.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution_4_3_f64_f64_instance.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution_4_4_f64_f64_instance.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution_5_1_f64_f64_instance.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution_5_2_f64_f64_instance.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution_5_3_f64_f64_instance.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution_5_4_f64_f64_instance.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution_5_5_f64_f64_instance.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution_6_1_f64_f64_instance.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution_6_2_f64_f64_instance.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution_6_3_f64_f64_instance.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution_6_4_f64_f64_instance.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution_6_5_f64_f64_instance.cpp
+   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution_6_6_f64_f64_instance.cpp
    ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution_instances.cpp
    ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution_registry.cpp
-   # ${CMAKE_CURRENT_SOURCE_DIR}/reduction_cpu_reference_instances.cpp
-   ${CMAKE_CURRENT_SOURCE_DIR}/reduction_solution.cpp
 )
 
 add_hiptensor_component(hiptensor_reduction ${HIPTENSOR_REDUCTION_SOURCES})
diff --git a/library/src/reduction/hiptensor_reduction.cpp b/library/src/reduction/hiptensor_reduction.cpp
index 540d6de5..7fe62943 100644
--- a/library/src/reduction/hiptensor_reduction.cpp
+++ b/library/src/reduction/hiptensor_reduction.cpp
@@ -185,8 +185,7 @@ hiptensorStatus_t hiptensorReduction(const hiptensorHandle_t*           handle,
 {
     using hiptensor::Logger;
     auto& logger = Logger::instance();
-    logger->setLogMask(0x1F);
-    char msg[2048];
+    char  msg[2048];
 
     snprintf(msg,
              sizeof(msg),
diff --git a/library/src/reduction/reduction_cpu_reference.cpp b/library/src/reduction/reduction_cpu_reference.cpp
new file mode 100644
index 00000000..4ffffa4f
--- /dev/null
+++ b/library/src/reduction/reduction_cpu_reference.cpp
@@ -0,0 +1,100 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include "reduction_cpu_reference.hpp"
+#include "reduction_cpu_reference_impl.hpp"
+#include "reduction_cpu_reference_instances.hpp"
+
+hiptensorStatus_t hiptensorReductionReference(const void*                        alpha,
+                                              const void*                        A,
+                                              const hiptensorTensorDescriptor_t* descA,
+                                              const int32_t                      modeA[],
+                                              const void*                        beta,
+                                              const void*                        C,
+                                              const hiptensorTensorDescriptor_t* descC,
+                                              const int32_t                      modeC[],
+                                              void*                              D,
+                                              const hiptensorTensorDescriptor_t* descD,
+                                              const int32_t                      modeD[],
+                                              hiptensorOperator_t                opReduce,
+                                              hiptensorComputeType_t             typeCompute,
+                                              hipStream_t                        stream)
+{
+    int  rankA        = descA->mLengths.size();
+    int  numReduceDim = descA->mLengths.size() - descD->mLengths.size();
+    auto ADataType    = descA->mType;
+    auto DDataType    = descD->mType;
+
+    auto& instances = hiptensor::ReductionCpuReferenceInstances::instance();
+    auto  solutionQ = instances->querySolutions(ADataType,
+                                               typeCompute,
+                                               DDataType,
+                                               rankA,
+                                               numReduceDim,
+                                               opReduce,
+                                               true, // @TODO hardcode
+                                               false); // @TODO hardcode
+
+    double alphaD;
+    if(alpha != nullptr)
+    {
+        alphaD = hiptensor::readVal<double>(alpha, typeCompute);
+    }
+    double betaD;
+    if(beta != nullptr)
+    {
+        betaD = hiptensor::readVal<double>(beta, typeCompute);
+    }
+
+    for(auto [_, pSolution] : solutionQ.solutions())
+    {
+        // Perform reduction with timing if LOG_LEVEL_PERF_TRACE
+        auto streamConfig        = StreamConfig{stream, false};
+        auto [isSupported, time] = (*pSolution)(descA->mLengths,
+                                                // @todo pass stride from descA
+                                                {},
+                                                {modeA, modeA + descA->mLengths.size()},
+                                                descC->mLengths,
+                                                {},
+                                                {modeC, modeC + descC->mLengths.size()},
+                                                alphaD,
+                                                betaD,
+                                                A,
+                                                D,
+                                                opReduce,
+                                                streamConfig);
+        if(isSupported)
+        {
+            if(time < 0)
+            {
+                return HIPTENSOR_STATUS_CK_ERROR;
+            }
+            return HIPTENSOR_STATUS_SUCCESS;
+        }
+    }
+
+    return HIPTENSOR_STATUS_INTERNAL_ERROR;
+}
diff --git a/library/src/reduction/reduction_cpu_reference.hpp b/library/src/reduction/reduction_cpu_reference.hpp
new file mode 100644
index 00000000..cc765270
--- /dev/null
+++ b/library/src/reduction/reduction_cpu_reference.hpp
@@ -0,0 +1,49 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#ifndef HIPTENSOR_REDUCTION_CPU_REFERENCE_HPP
+#define HIPTENSOR_REDUCTION_CPU_REFERENCE_HPP
+
+#include <hip/library_types.h>
+#include <vector>
+
+#include <hiptensor/hiptensor.hpp>
+
+hiptensorStatus_t hiptensorReductionReference(const void*                        alpha,
+                                              const void*                        A,
+                                              const hiptensorTensorDescriptor_t* descA,
+                                              const int32_t                      modeA[],
+                                              const void*                        beta,
+                                              const void*                        C,
+                                              const hiptensorTensorDescriptor_t* descC,
+                                              const int32_t                      modeC[],
+                                              void*                              D,
+                                              const hiptensorTensorDescriptor_t* descD,
+                                              const int32_t                      modeD[],
+                                              hiptensorOperator_t                opReduce,
+                                              hiptensorComputeType_t             typeCompute,
+                                              hipStream_t                        stream);
+#endif // HIPTENSOR_REDUCTION_CPU_REFERENCE_HPP
diff --git a/library/src/reduction/reduction_cpu_reference_impl.hpp b/library/src/reduction/reduction_cpu_reference_impl.hpp
new file mode 100644
index 00000000..11c0b2c5
--- /dev/null
+++ b/library/src/reduction/reduction_cpu_reference_impl.hpp
@@ -0,0 +1,139 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#ifndef HIPTENSOR_REDUCTION_CPU_REFERENCE_IMPL_HPP
+#define HIPTENSOR_REDUCTION_CPU_REFERENCE_IMPL_HPP
+
+// Std includes
+#include <array>
+#include <list>
+#include <numeric>
+#include <vector>
+
+// CK includes
+#include "ck/library/reference_tensor_operation/cpu/reference_reduce.hpp"
+
+#include "reduction_meta_traits.hpp"
+#include "reduction_solution.hpp"
+
+namespace hiptensor
+{
+
+    template <typename InDataType,
+              typename AccDataType,
+              typename OutDataType,
+              ck::index_t Rank,
+              ck::index_t NumReduceDim,
+              typename ReduceOperation,
+              typename InElementwiseOperation,
+              typename AccElementwiseOperation,
+              bool PropagateNan,
+              bool OutputIndex>
+    using ReferenceReduction = ck::tensor_operation::host::ReferenceReduce<InDataType,
+                                                                           AccDataType,
+                                                                           OutDataType,
+                                                                           Rank,
+                                                                           NumReduceDim,
+                                                                           ReduceOperation,
+                                                                           InElementwiseOperation,
+                                                                           AccElementwiseOperation,
+                                                                           PropagateNan,
+                                                                           OutputIndex>;
+
+    // Partial specialize for reference reduction
+    template <typename InDataType,
+              typename AccDataType,
+              typename OutDataType,
+              ck::index_t Rank,
+              ck::index_t NumReduceDim,
+              typename ReduceOperation,
+              typename InElementwiseOperation,
+              typename AccElementwiseOperation,
+              bool PropagateNan,
+              bool OutputIndex>
+    struct MetaTraits<ReferenceReduction<InDataType,
+                                         AccDataType,
+                                         OutDataType,
+                                         Rank,
+                                         NumReduceDim,
+                                         ReduceOperation,
+                                         InElementwiseOperation,
+                                         AccElementwiseOperation,
+                                         PropagateNan,
+                                         OutputIndex>>
+        : public MetaTraits<ck::tensor_operation::device::DeviceReduce<InDataType,
+                                                                       AccDataType,
+                                                                       OutDataType,
+                                                                       Rank,
+                                                                       NumReduceDim,
+                                                                       ReduceOperation,
+                                                                       InElementwiseOperation,
+                                                                       AccElementwiseOperation,
+                                                                       PropagateNan,
+                                                                       OutputIndex>>
+    {
+    };
+
+    template <typename InDataType,
+              typename AccDataType,
+              typename OutDataType,
+              int                 Rank,
+              int                 NumReduceDim,
+              hiptensorOperator_t opReduce,
+              bool                PropagateNan,
+              bool                OutputIndex>
+    auto enumerateReferenceSolutions()
+    {
+        constexpr auto ReduceOpId = convertHiptensorReduceOperatorToCk<opReduce>();
+
+        using ReduceOperation = typename ck::reduce_binary_operator<ReduceOpId>::opType;
+        using InElementwiseOperation =
+            typename ck::reduce_unary_operator<ReduceOpId, true, true>::InElementwiseOperation;
+        using AccElementwiseOperation =
+            typename ck::reduce_unary_operator<ReduceOpId, true, true>::AccElementwiseOperation;
+        using ReferenceOp = ReferenceReduction<InDataType,
+                                               AccDataType,
+                                               OutDataType,
+                                               Rank,
+                                               NumReduceDim,
+                                               ReduceOperation,
+                                               InElementwiseOperation,
+                                               AccElementwiseOperation,
+                                               PropagateNan,
+                                               OutputIndex>;
+
+        auto solution
+            = std::make_unique<ReductionSolutionImpl<ReferenceOp>>(std::make_unique<ReferenceOp>());
+
+        auto result = std::vector<std::unique_ptr<ReductionSolution>>();
+        result.push_back(std::move(solution));
+
+        return result;
+    }
+
+} // namespace hiptensor
+
+#endif // HIPTENSOR_REDUCTION_CPU_REFERENCE_IMPL_HPP
diff --git a/library/src/reduction/reduction_cpu_reference_instances.cpp b/library/src/reduction/reduction_cpu_reference_instances.cpp
new file mode 100644
index 00000000..f0af6bad
--- /dev/null
+++ b/library/src/reduction/reduction_cpu_reference_instances.cpp
@@ -0,0 +1,131 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include "reduction_cpu_reference_instances.hpp"
+#include "reduction_cpu_reference_impl.hpp"
+#include <hiptensor/internal/types.hpp>
+
+#define REG_CPU_SOLUTION(dim_count, reduced_dim_count, type)         \
+    registerSolutions(enumerateReferenceSolutions<type,              \
+                                                  type,              \
+                                                  type,              \
+                                                  dim_count,         \
+                                                  reduced_dim_count, \
+                                                  HIPTENSOR_OP_ADD,  \
+                                                  true,              \
+                                                  false>());         \
+    registerSolutions(enumerateReferenceSolutions<type,              \
+                                                  type,              \
+                                                  type,              \
+                                                  dim_count,         \
+                                                  reduced_dim_count, \
+                                                  HIPTENSOR_OP_MUL,  \
+                                                  true,              \
+                                                  false>());         \
+    registerSolutions(enumerateReferenceSolutions<type,              \
+                                                  type,              \
+                                                  type,              \
+                                                  dim_count,         \
+                                                  reduced_dim_count, \
+                                                  HIPTENSOR_OP_MIN,  \
+                                                  true,              \
+                                                  false>());         \
+    registerSolutions(enumerateReferenceSolutions<type,              \
+                                                  type,              \
+                                                  type,              \
+                                                  dim_count,         \
+                                                  reduced_dim_count, \
+                                                  HIPTENSOR_OP_MAX,  \
+                                                  true,              \
+                                                  false>());
+
+namespace hiptensor
+{
+    ReductionCpuReferenceInstances::ReductionCpuReferenceInstances()
+    {
+        // @todo Add all reference instances. How to stop the explosion of number of instances? Wait for CK
+        // registerSolutions(enumerateReferenceSolutions<hiptensor::bfloat16_t,
+        // hiptensor::bfloat16_t,
+        // hiptensor::bfloat16_t,
+        // 4,
+        // 2,
+        // HIPTENSOR_OP_ADD,
+        // true, // PropagateNan,
+        // false>()); // OutputIndex,
+        // registerSolutions(enumerateReferenceSolutions<hiptensor::float16_t,
+        // hiptensor::float16_t,
+        // hiptensor::float16_t,
+        // 4,
+        // 2,
+        // HIPTENSOR_OP_ADD,
+        // true, // PropagateNan,
+        // false>()); // OutputIndex,
+
+        REG_CPU_SOLUTION(1, 1, hiptensor::float32_t);
+        REG_CPU_SOLUTION(2, 1, hiptensor::float32_t);
+        REG_CPU_SOLUTION(2, 2, hiptensor::float32_t);
+        REG_CPU_SOLUTION(3, 1, hiptensor::float32_t);
+        REG_CPU_SOLUTION(3, 2, hiptensor::float32_t);
+        REG_CPU_SOLUTION(3, 3, hiptensor::float32_t);
+        REG_CPU_SOLUTION(4, 1, hiptensor::float32_t);
+        REG_CPU_SOLUTION(4, 2, hiptensor::float32_t);
+        REG_CPU_SOLUTION(4, 3, hiptensor::float32_t);
+        REG_CPU_SOLUTION(4, 4, hiptensor::float32_t);
+        REG_CPU_SOLUTION(5, 1, hiptensor::float32_t);
+        REG_CPU_SOLUTION(5, 2, hiptensor::float32_t);
+        REG_CPU_SOLUTION(5, 3, hiptensor::float32_t);
+        REG_CPU_SOLUTION(5, 4, hiptensor::float32_t);
+        REG_CPU_SOLUTION(5, 5, hiptensor::float32_t);
+        REG_CPU_SOLUTION(6, 1, hiptensor::float32_t);
+        REG_CPU_SOLUTION(6, 2, hiptensor::float32_t);
+        REG_CPU_SOLUTION(6, 3, hiptensor::float32_t);
+        REG_CPU_SOLUTION(6, 4, hiptensor::float32_t);
+        REG_CPU_SOLUTION(6, 5, hiptensor::float32_t);
+        REG_CPU_SOLUTION(6, 6, hiptensor::float32_t);
+
+        REG_CPU_SOLUTION(1, 1, hiptensor::float64_t);
+        REG_CPU_SOLUTION(2, 1, hiptensor::float64_t);
+        REG_CPU_SOLUTION(2, 2, hiptensor::float64_t);
+        REG_CPU_SOLUTION(3, 1, hiptensor::float64_t);
+        REG_CPU_SOLUTION(3, 2, hiptensor::float64_t);
+        REG_CPU_SOLUTION(3, 3, hiptensor::float64_t);
+        REG_CPU_SOLUTION(4, 1, hiptensor::float64_t);
+        REG_CPU_SOLUTION(4, 2, hiptensor::float64_t);
+        REG_CPU_SOLUTION(4, 3, hiptensor::float64_t);
+        REG_CPU_SOLUTION(4, 4, hiptensor::float64_t);
+        REG_CPU_SOLUTION(5, 1, hiptensor::float64_t);
+        REG_CPU_SOLUTION(5, 2, hiptensor::float64_t);
+        REG_CPU_SOLUTION(5, 3, hiptensor::float64_t);
+        REG_CPU_SOLUTION(5, 4, hiptensor::float64_t);
+        REG_CPU_SOLUTION(5, 5, hiptensor::float64_t);
+        REG_CPU_SOLUTION(6, 1, hiptensor::float64_t);
+        REG_CPU_SOLUTION(6, 2, hiptensor::float64_t);
+        REG_CPU_SOLUTION(6, 3, hiptensor::float64_t);
+        REG_CPU_SOLUTION(6, 4, hiptensor::float64_t);
+        REG_CPU_SOLUTION(6, 5, hiptensor::float64_t);
+        REG_CPU_SOLUTION(6, 6, hiptensor::float64_t);
+    }
+} // namespace hiptensor
diff --git a/library/src/reduction/reduction_cpu_reference_instances.hpp b/library/src/reduction/reduction_cpu_reference_instances.hpp
new file mode 100644
index 00000000..996395b1
--- /dev/null
+++ b/library/src/reduction/reduction_cpu_reference_instances.hpp
@@ -0,0 +1,63 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#ifndef HIPTENSOR_REDUCTION_CPU_REFERENCE_INSTANCES_HPP
+#define HIPTENSOR_REDUCTION_CPU_REFERENCE_INSTANCES_HPP
+
+#include <memory>
+
+#include "reduction_solution_registry.hpp"
+#include "singleton.hpp"
+
+namespace hiptensor
+{
+    class ReductionCpuReferenceInstances : public ReductionSolutionRegistry,
+                                           public LazySingleton<ReductionCpuReferenceInstances>
+    {
+    public:
+        // For static initialization
+        friend std::unique_ptr<ReductionCpuReferenceInstances>
+            std::make_unique<ReductionCpuReferenceInstances>();
+
+        ~ReductionCpuReferenceInstances() = default;
+
+    private:
+        void ReductionCpuReference2DInstances();
+        void ReductionCpuReference3DInstances();
+        void ReductionCpuReference4DInstances();
+        void ReductionCpuReference5DInstances();
+        void ReductionCpuReference6DInstances();
+        // Singleton: only one instance
+        ReductionCpuReferenceInstances();
+        ReductionCpuReferenceInstances(ReductionCpuReferenceInstances const&)            = delete;
+        ReductionCpuReferenceInstances(ReductionCpuReferenceInstances&&)                 = delete;
+        ReductionCpuReferenceInstances& operator=(ReductionCpuReferenceInstances const&) = delete;
+        ReductionCpuReferenceInstances& operator=(ReductionCpuReferenceInstances&&)      = delete;
+    };
+
+} // namespace hiptensor
+
+#endif // HIPTENSOR_REDUCTION_SOLUTION_INSTANCES_HPP
diff --git a/library/src/reduction/reduction_meta_traits.hpp b/library/src/reduction/reduction_meta_traits.hpp
index c54e8d74..aebdea75 100644
--- a/library/src/reduction/reduction_meta_traits.hpp
+++ b/library/src/reduction/reduction_meta_traits.hpp
@@ -61,10 +61,12 @@ namespace hiptensor
         constexpr static bool        TensorPropagateNan = PropagateNan;
         constexpr static bool        TensorOutputIndex  = OutputIndex;
 
-        using TensorInDataType      = InDataType;
-        using TensorAccDataType     = AccDataType;
-        using TensorOutDataType     = OutDataType;
-        using TensorReduceOperation = ReduceOperation;
+        using TensorInDataType       = InDataType;
+        using TensorAccDataType      = AccDataType;
+        using TensorOutDataType      = OutDataType;
+        using TensorReduceOperation  = ReduceOperation;
+        using TensorInElementwiseOp  = InElementwiseOp;
+        using TensorAccElementwiseOp = AccElementwiseOp;
         static_assert((std::is_same_v<TensorReduceOperation, typename ck::reduce::Add>)
                           || (std::is_same_v<TensorReduceOperation, typename ck::reduce::Mul>)
                           || (std::is_same_v<TensorReduceOperation, typename ck::reduce::Min>)
diff --git a/library/src/reduction/reduction_solution_1_1_f32_f32_instance.cpp b/library/src/reduction/reduction_solution_1_1_f32_f32_instance.cpp
new file mode 100644
index 00000000..89c0535d
--- /dev/null
+++ b/library/src/reduction/reduction_solution_1_1_f32_f32_instance.cpp
@@ -0,0 +1,70 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include "hiptensor/internal/types.hpp"
+#include "reduction_solution.hpp"
+#include "reduction_solution_instances.hpp"
+
+namespace hiptensor
+{
+    void ReductionSolutionInstances::genReductionSolution1x1F32F32Instances()
+    {
+
+        // add entries to mSolutionQuery
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      1,
+                                                      1,
+                                                      HIPTENSOR_OP_ADD,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      1,
+                                                      1,
+                                                      HIPTENSOR_OP_MUL,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      1,
+                                                      1,
+                                                      HIPTENSOR_OP_MIN,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      1,
+                                                      1,
+                                                      HIPTENSOR_OP_MAX,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+    }
+} // namespace hiptensor
diff --git a/library/src/reduction/reduction_solution_1_1_f64_f64_instance.cpp b/library/src/reduction/reduction_solution_1_1_f64_f64_instance.cpp
new file mode 100644
index 00000000..0a343589
--- /dev/null
+++ b/library/src/reduction/reduction_solution_1_1_f64_f64_instance.cpp
@@ -0,0 +1,69 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include "hiptensor/internal/types.hpp"
+#include "reduction_solution.hpp"
+#include "reduction_solution_instances.hpp"
+
+namespace hiptensor
+{
+    void ReductionSolutionInstances::genReductionSolution1x1F64F64Instances()
+    {
+        // add entries to mSolutionQuery
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      1,
+                                                      1,
+                                                      HIPTENSOR_OP_ADD,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      1,
+                                                      1,
+                                                      HIPTENSOR_OP_MUL,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      1,
+                                                      1,
+                                                      HIPTENSOR_OP_MIN,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      1,
+                                                      1,
+                                                      HIPTENSOR_OP_MAX,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+    }
+} // namespace hiptensor
diff --git a/library/src/reduction/reduction_solution_2_1_f32_f32_instance.cpp b/library/src/reduction/reduction_solution_2_1_f32_f32_instance.cpp
new file mode 100644
index 00000000..e69c36be
--- /dev/null
+++ b/library/src/reduction/reduction_solution_2_1_f32_f32_instance.cpp
@@ -0,0 +1,70 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include "hiptensor/internal/types.hpp"
+#include "reduction_solution.hpp"
+#include "reduction_solution_instances.hpp"
+
+namespace hiptensor
+{
+    void ReductionSolutionInstances::genReductionSolution2x1F32F32Instances()
+    {
+
+        // add entries to mSolutionQuery
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      2,
+                                                      1,
+                                                      HIPTENSOR_OP_ADD,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      2,
+                                                      1,
+                                                      HIPTENSOR_OP_MUL,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      2,
+                                                      1,
+                                                      HIPTENSOR_OP_MIN,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      2,
+                                                      1,
+                                                      HIPTENSOR_OP_MAX,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+    }
+} // namespace hiptensor
diff --git a/library/src/reduction/reduction_solution_2_1_f64_f64_instance.cpp b/library/src/reduction/reduction_solution_2_1_f64_f64_instance.cpp
new file mode 100644
index 00000000..30cc9707
--- /dev/null
+++ b/library/src/reduction/reduction_solution_2_1_f64_f64_instance.cpp
@@ -0,0 +1,69 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include "hiptensor/internal/types.hpp"
+#include "reduction_solution.hpp"
+#include "reduction_solution_instances.hpp"
+
+namespace hiptensor
+{
+    void ReductionSolutionInstances::genReductionSolution2x1F64F64Instances()
+    {
+        // add entries to mSolutionQuery
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      2,
+                                                      1,
+                                                      HIPTENSOR_OP_ADD,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      2,
+                                                      1,
+                                                      HIPTENSOR_OP_MUL,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      2,
+                                                      1,
+                                                      HIPTENSOR_OP_MIN,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      2,
+                                                      1,
+                                                      HIPTENSOR_OP_MAX,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+    }
+} // namespace hiptensor
diff --git a/library/src/reduction/reduction_solution_2_2_f32_f32_instance.cpp b/library/src/reduction/reduction_solution_2_2_f32_f32_instance.cpp
new file mode 100644
index 00000000..14aec415
--- /dev/null
+++ b/library/src/reduction/reduction_solution_2_2_f32_f32_instance.cpp
@@ -0,0 +1,70 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include "hiptensor/internal/types.hpp"
+#include "reduction_solution.hpp"
+#include "reduction_solution_instances.hpp"
+
+namespace hiptensor
+{
+    void ReductionSolutionInstances::genReductionSolution2x2F32F32Instances()
+    {
+
+        // add entries to mSolutionQuery
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      2,
+                                                      2,
+                                                      HIPTENSOR_OP_ADD,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      2,
+                                                      2,
+                                                      HIPTENSOR_OP_MUL,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      2,
+                                                      2,
+                                                      HIPTENSOR_OP_MIN,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      2,
+                                                      2,
+                                                      HIPTENSOR_OP_MAX,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+    }
+} // namespace hiptensor
diff --git a/library/src/reduction/reduction_solution_2_2_f64_f64_instance.cpp b/library/src/reduction/reduction_solution_2_2_f64_f64_instance.cpp
new file mode 100644
index 00000000..35f34b66
--- /dev/null
+++ b/library/src/reduction/reduction_solution_2_2_f64_f64_instance.cpp
@@ -0,0 +1,69 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include "hiptensor/internal/types.hpp"
+#include "reduction_solution.hpp"
+#include "reduction_solution_instances.hpp"
+
+namespace hiptensor
+{
+    void ReductionSolutionInstances::genReductionSolution2x2F64F64Instances()
+    {
+        // add entries to mSolutionQuery
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      2,
+                                                      2,
+                                                      HIPTENSOR_OP_ADD,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      2,
+                                                      2,
+                                                      HIPTENSOR_OP_MUL,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      2,
+                                                      2,
+                                                      HIPTENSOR_OP_MIN,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      2,
+                                                      2,
+                                                      HIPTENSOR_OP_MAX,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+    }
+} // namespace hiptensor
diff --git a/library/src/reduction/reduction_solution_3_1_f32_f32_instance.cpp b/library/src/reduction/reduction_solution_3_1_f32_f32_instance.cpp
new file mode 100644
index 00000000..fdc768fa
--- /dev/null
+++ b/library/src/reduction/reduction_solution_3_1_f32_f32_instance.cpp
@@ -0,0 +1,70 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include "hiptensor/internal/types.hpp"
+#include "reduction_solution.hpp"
+#include "reduction_solution_instances.hpp"
+
+namespace hiptensor
+{
+    void ReductionSolutionInstances::genReductionSolution3x1F32F32Instances()
+    {
+
+        // add entries to mSolutionQuery
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      3,
+                                                      1,
+                                                      HIPTENSOR_OP_ADD,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      3,
+                                                      1,
+                                                      HIPTENSOR_OP_MUL,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      3,
+                                                      1,
+                                                      HIPTENSOR_OP_MIN,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      3,
+                                                      1,
+                                                      HIPTENSOR_OP_MAX,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+    }
+} // namespace hiptensor
diff --git a/library/src/reduction/reduction_solution_3_1_f64_f64_instance.cpp b/library/src/reduction/reduction_solution_3_1_f64_f64_instance.cpp
new file mode 100644
index 00000000..7281c939
--- /dev/null
+++ b/library/src/reduction/reduction_solution_3_1_f64_f64_instance.cpp
@@ -0,0 +1,69 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include "hiptensor/internal/types.hpp"
+#include "reduction_solution.hpp"
+#include "reduction_solution_instances.hpp"
+
+namespace hiptensor
+{
+    void ReductionSolutionInstances::genReductionSolution3x1F64F64Instances()
+    {
+        // add entries to mSolutionQuery
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      3,
+                                                      1,
+                                                      HIPTENSOR_OP_ADD,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      3,
+                                                      1,
+                                                      HIPTENSOR_OP_MUL,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      3,
+                                                      1,
+                                                      HIPTENSOR_OP_MIN,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      3,
+                                                      1,
+                                                      HIPTENSOR_OP_MAX,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+    }
+} // namespace hiptensor
diff --git a/library/src/reduction/reduction_solution_3_2_f32_f32_instance.cpp b/library/src/reduction/reduction_solution_3_2_f32_f32_instance.cpp
new file mode 100644
index 00000000..5e8f5e15
--- /dev/null
+++ b/library/src/reduction/reduction_solution_3_2_f32_f32_instance.cpp
@@ -0,0 +1,70 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include "hiptensor/internal/types.hpp"
+#include "reduction_solution.hpp"
+#include "reduction_solution_instances.hpp"
+
+namespace hiptensor
+{
+    void ReductionSolutionInstances::genReductionSolution3x2F32F32Instances()
+    {
+
+        // add entries to mSolutionQuery
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      3,
+                                                      2,
+                                                      HIPTENSOR_OP_ADD,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      3,
+                                                      2,
+                                                      HIPTENSOR_OP_MUL,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      3,
+                                                      2,
+                                                      HIPTENSOR_OP_MIN,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      3,
+                                                      2,
+                                                      HIPTENSOR_OP_MAX,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+    }
+} // namespace hiptensor
diff --git a/library/src/reduction/reduction_solution_3_2_f64_f64_instance.cpp b/library/src/reduction/reduction_solution_3_2_f64_f64_instance.cpp
new file mode 100644
index 00000000..205ab06b
--- /dev/null
+++ b/library/src/reduction/reduction_solution_3_2_f64_f64_instance.cpp
@@ -0,0 +1,69 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include "hiptensor/internal/types.hpp"
+#include "reduction_solution.hpp"
+#include "reduction_solution_instances.hpp"
+
+namespace hiptensor
+{
+    void ReductionSolutionInstances::genReductionSolution3x2F64F64Instances()
+    {
+        // add entries to mSolutionQuery
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      3,
+                                                      2,
+                                                      HIPTENSOR_OP_ADD,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      3,
+                                                      2,
+                                                      HIPTENSOR_OP_MUL,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      3,
+                                                      2,
+                                                      HIPTENSOR_OP_MIN,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      3,
+                                                      2,
+                                                      HIPTENSOR_OP_MAX,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+    }
+} // namespace hiptensor
diff --git a/library/src/reduction/reduction_solution_3_3_f32_f32_instance.cpp b/library/src/reduction/reduction_solution_3_3_f32_f32_instance.cpp
new file mode 100644
index 00000000..868b35a0
--- /dev/null
+++ b/library/src/reduction/reduction_solution_3_3_f32_f32_instance.cpp
@@ -0,0 +1,70 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include "hiptensor/internal/types.hpp"
+#include "reduction_solution.hpp"
+#include "reduction_solution_instances.hpp"
+
+namespace hiptensor
+{
+    void ReductionSolutionInstances::genReductionSolution3x3F32F32Instances()
+    {
+
+        // add entries to mSolutionQuery
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      3,
+                                                      3,
+                                                      HIPTENSOR_OP_ADD,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      3,
+                                                      3,
+                                                      HIPTENSOR_OP_MUL,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      3,
+                                                      3,
+                                                      HIPTENSOR_OP_MIN,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      3,
+                                                      3,
+                                                      HIPTENSOR_OP_MAX,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+    }
+} // namespace hiptensor
diff --git a/library/src/reduction/reduction_solution_3_3_f64_f64_instance.cpp b/library/src/reduction/reduction_solution_3_3_f64_f64_instance.cpp
new file mode 100644
index 00000000..12bb4f89
--- /dev/null
+++ b/library/src/reduction/reduction_solution_3_3_f64_f64_instance.cpp
@@ -0,0 +1,69 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include "hiptensor/internal/types.hpp"
+#include "reduction_solution.hpp"
+#include "reduction_solution_instances.hpp"
+
+namespace hiptensor
+{
+    void ReductionSolutionInstances::genReductionSolution3x3F64F64Instances()
+    {
+        // add entries to mSolutionQuery
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      3,
+                                                      3,
+                                                      HIPTENSOR_OP_ADD,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      3,
+                                                      3,
+                                                      HIPTENSOR_OP_MUL,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      3,
+                                                      3,
+                                                      HIPTENSOR_OP_MIN,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      3,
+                                                      3,
+                                                      HIPTENSOR_OP_MAX,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+    }
+} // namespace hiptensor
diff --git a/library/src/reduction/reduction_solution_4_1_f32_f32_instance.cpp b/library/src/reduction/reduction_solution_4_1_f32_f32_instance.cpp
new file mode 100644
index 00000000..810ea926
--- /dev/null
+++ b/library/src/reduction/reduction_solution_4_1_f32_f32_instance.cpp
@@ -0,0 +1,70 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include "hiptensor/internal/types.hpp"
+#include "reduction_solution.hpp"
+#include "reduction_solution_instances.hpp"
+
+namespace hiptensor
+{
+    void ReductionSolutionInstances::genReductionSolution4x1F32F32Instances()
+    {
+
+        // add entries to mSolutionQuery
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      4,
+                                                      1,
+                                                      HIPTENSOR_OP_ADD,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      4,
+                                                      1,
+                                                      HIPTENSOR_OP_MUL,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      4,
+                                                      1,
+                                                      HIPTENSOR_OP_MIN,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      4,
+                                                      1,
+                                                      HIPTENSOR_OP_MAX,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+    }
+} // namespace hiptensor
diff --git a/library/src/reduction/reduction_solution_4_1_f64_f64_instance.cpp b/library/src/reduction/reduction_solution_4_1_f64_f64_instance.cpp
new file mode 100644
index 00000000..e011484e
--- /dev/null
+++ b/library/src/reduction/reduction_solution_4_1_f64_f64_instance.cpp
@@ -0,0 +1,69 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include "hiptensor/internal/types.hpp"
+#include "reduction_solution.hpp"
+#include "reduction_solution_instances.hpp"
+
+namespace hiptensor
+{
+    void ReductionSolutionInstances::genReductionSolution4x1F64F64Instances()
+    {
+        // add entries to mSolutionQuery
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      4,
+                                                      1,
+                                                      HIPTENSOR_OP_ADD,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      4,
+                                                      1,
+                                                      HIPTENSOR_OP_MUL,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      4,
+                                                      1,
+                                                      HIPTENSOR_OP_MIN,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      4,
+                                                      1,
+                                                      HIPTENSOR_OP_MAX,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+    }
+} // namespace hiptensor
diff --git a/library/src/reduction/reduction_solution_4_2_f32_f32_instance.cpp b/library/src/reduction/reduction_solution_4_2_f32_f32_instance.cpp
new file mode 100644
index 00000000..42c988da
--- /dev/null
+++ b/library/src/reduction/reduction_solution_4_2_f32_f32_instance.cpp
@@ -0,0 +1,70 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include "hiptensor/internal/types.hpp"
+#include "reduction_solution.hpp"
+#include "reduction_solution_instances.hpp"
+
+namespace hiptensor
+{
+    void ReductionSolutionInstances::genReductionSolution4x2F32F32Instances()
+    {
+
+        // add entries to mSolutionQuery
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      4,
+                                                      2,
+                                                      HIPTENSOR_OP_ADD,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      4,
+                                                      2,
+                                                      HIPTENSOR_OP_MUL,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      4,
+                                                      2,
+                                                      HIPTENSOR_OP_MIN,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      4,
+                                                      2,
+                                                      HIPTENSOR_OP_MAX,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+    }
+} // namespace hiptensor
diff --git a/library/src/reduction/reduction_solution_4_2_f64_f64_instance.cpp b/library/src/reduction/reduction_solution_4_2_f64_f64_instance.cpp
new file mode 100644
index 00000000..7c23b4c7
--- /dev/null
+++ b/library/src/reduction/reduction_solution_4_2_f64_f64_instance.cpp
@@ -0,0 +1,69 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include "hiptensor/internal/types.hpp"
+#include "reduction_solution.hpp"
+#include "reduction_solution_instances.hpp"
+
+namespace hiptensor
+{
+    void ReductionSolutionInstances::genReductionSolution4x2F64F64Instances()
+    {
+        // add entries to mSolutionQuery
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      4,
+                                                      2,
+                                                      HIPTENSOR_OP_ADD,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      4,
+                                                      2,
+                                                      HIPTENSOR_OP_MUL,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      4,
+                                                      2,
+                                                      HIPTENSOR_OP_MIN,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      4,
+                                                      2,
+                                                      HIPTENSOR_OP_MAX,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+    }
+} // namespace hiptensor
diff --git a/library/src/reduction/reduction_solution_4_3_f32_f32_instance.cpp b/library/src/reduction/reduction_solution_4_3_f32_f32_instance.cpp
new file mode 100644
index 00000000..a18b078a
--- /dev/null
+++ b/library/src/reduction/reduction_solution_4_3_f32_f32_instance.cpp
@@ -0,0 +1,70 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include "hiptensor/internal/types.hpp"
+#include "reduction_solution.hpp"
+#include "reduction_solution_instances.hpp"
+
+namespace hiptensor
+{
+    void ReductionSolutionInstances::genReductionSolution4x3F32F32Instances()
+    {
+
+        // add entries to mSolutionQuery
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      4,
+                                                      3,
+                                                      HIPTENSOR_OP_ADD,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      4,
+                                                      3,
+                                                      HIPTENSOR_OP_MUL,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      4,
+                                                      3,
+                                                      HIPTENSOR_OP_MIN,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      4,
+                                                      3,
+                                                      HIPTENSOR_OP_MAX,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+    }
+} // namespace hiptensor
diff --git a/library/src/reduction/reduction_solution_4_3_f64_f64_instance.cpp b/library/src/reduction/reduction_solution_4_3_f64_f64_instance.cpp
new file mode 100644
index 00000000..845a710f
--- /dev/null
+++ b/library/src/reduction/reduction_solution_4_3_f64_f64_instance.cpp
@@ -0,0 +1,69 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include "hiptensor/internal/types.hpp"
+#include "reduction_solution.hpp"
+#include "reduction_solution_instances.hpp"
+
+namespace hiptensor
+{
+    void ReductionSolutionInstances::genReductionSolution4x3F64F64Instances()
+    {
+        // add entries to mSolutionQuery
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      4,
+                                                      3,
+                                                      HIPTENSOR_OP_ADD,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      4,
+                                                      3,
+                                                      HIPTENSOR_OP_MUL,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      4,
+                                                      3,
+                                                      HIPTENSOR_OP_MIN,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      4,
+                                                      3,
+                                                      HIPTENSOR_OP_MAX,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+    }
+} // namespace hiptensor
diff --git a/library/src/reduction/reduction_solution_4_4_f32_f32_instance.cpp b/library/src/reduction/reduction_solution_4_4_f32_f32_instance.cpp
new file mode 100644
index 00000000..969feedb
--- /dev/null
+++ b/library/src/reduction/reduction_solution_4_4_f32_f32_instance.cpp
@@ -0,0 +1,70 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include "hiptensor/internal/types.hpp"
+#include "reduction_solution.hpp"
+#include "reduction_solution_instances.hpp"
+
+namespace hiptensor
+{
+    void ReductionSolutionInstances::genReductionSolution4x4F32F32Instances()
+    {
+
+        // add entries to mSolutionQuery
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      4,
+                                                      4,
+                                                      HIPTENSOR_OP_ADD,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      4,
+                                                      4,
+                                                      HIPTENSOR_OP_MUL,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      4,
+                                                      4,
+                                                      HIPTENSOR_OP_MIN,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      4,
+                                                      4,
+                                                      HIPTENSOR_OP_MAX,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+    }
+} // namespace hiptensor
diff --git a/library/src/reduction/reduction_solution_4_4_f64_f64_instance.cpp b/library/src/reduction/reduction_solution_4_4_f64_f64_instance.cpp
new file mode 100644
index 00000000..c0ae9289
--- /dev/null
+++ b/library/src/reduction/reduction_solution_4_4_f64_f64_instance.cpp
@@ -0,0 +1,69 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include "hiptensor/internal/types.hpp"
+#include "reduction_solution.hpp"
+#include "reduction_solution_instances.hpp"
+
+namespace hiptensor
+{
+    void ReductionSolutionInstances::genReductionSolution4x4F64F64Instances()
+    {
+        // add entries to mSolutionQuery
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      4,
+                                                      4,
+                                                      HIPTENSOR_OP_ADD,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      4,
+                                                      4,
+                                                      HIPTENSOR_OP_MUL,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      4,
+                                                      4,
+                                                      HIPTENSOR_OP_MIN,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      4,
+                                                      4,
+                                                      HIPTENSOR_OP_MAX,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+    }
+} // namespace hiptensor
diff --git a/library/src/reduction/reduction_solution_5_1_f32_f32_instance.cpp b/library/src/reduction/reduction_solution_5_1_f32_f32_instance.cpp
new file mode 100644
index 00000000..90b3bc64
--- /dev/null
+++ b/library/src/reduction/reduction_solution_5_1_f32_f32_instance.cpp
@@ -0,0 +1,70 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include "hiptensor/internal/types.hpp"
+#include "reduction_solution.hpp"
+#include "reduction_solution_instances.hpp"
+
+namespace hiptensor
+{
+    void ReductionSolutionInstances::genReductionSolution5x1F32F32Instances()
+    {
+
+        // add entries to mSolutionQuery
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      5,
+                                                      1,
+                                                      HIPTENSOR_OP_ADD,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      5,
+                                                      1,
+                                                      HIPTENSOR_OP_MUL,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      5,
+                                                      1,
+                                                      HIPTENSOR_OP_MIN,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      5,
+                                                      1,
+                                                      HIPTENSOR_OP_MAX,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+    }
+} // namespace hiptensor
diff --git a/library/src/reduction/reduction_solution_5_1_f64_f64_instance.cpp b/library/src/reduction/reduction_solution_5_1_f64_f64_instance.cpp
new file mode 100644
index 00000000..6e84173f
--- /dev/null
+++ b/library/src/reduction/reduction_solution_5_1_f64_f64_instance.cpp
@@ -0,0 +1,69 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include "hiptensor/internal/types.hpp"
+#include "reduction_solution.hpp"
+#include "reduction_solution_instances.hpp"
+
+namespace hiptensor
+{
+    void ReductionSolutionInstances::genReductionSolution5x1F64F64Instances()
+    {
+        // add entries to mSolutionQuery
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      5,
+                                                      1,
+                                                      HIPTENSOR_OP_ADD,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      5,
+                                                      1,
+                                                      HIPTENSOR_OP_MUL,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      5,
+                                                      1,
+                                                      HIPTENSOR_OP_MIN,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      5,
+                                                      1,
+                                                      HIPTENSOR_OP_MAX,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+    }
+} // namespace hiptensor
diff --git a/library/src/reduction/reduction_solution_5_2_f32_f32_instance.cpp b/library/src/reduction/reduction_solution_5_2_f32_f32_instance.cpp
new file mode 100644
index 00000000..45f2641e
--- /dev/null
+++ b/library/src/reduction/reduction_solution_5_2_f32_f32_instance.cpp
@@ -0,0 +1,70 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include "hiptensor/internal/types.hpp"
+#include "reduction_solution.hpp"
+#include "reduction_solution_instances.hpp"
+
+namespace hiptensor
+{
+    void ReductionSolutionInstances::genReductionSolution5x2F32F32Instances()
+    {
+
+        // add entries to mSolutionQuery
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      5,
+                                                      2,
+                                                      HIPTENSOR_OP_ADD,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      5,
+                                                      2,
+                                                      HIPTENSOR_OP_MUL,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      5,
+                                                      2,
+                                                      HIPTENSOR_OP_MIN,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      5,
+                                                      2,
+                                                      HIPTENSOR_OP_MAX,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+    }
+} // namespace hiptensor
diff --git a/library/src/reduction/reduction_solution_5_2_f64_f64_instance.cpp b/library/src/reduction/reduction_solution_5_2_f64_f64_instance.cpp
new file mode 100644
index 00000000..45d59e92
--- /dev/null
+++ b/library/src/reduction/reduction_solution_5_2_f64_f64_instance.cpp
@@ -0,0 +1,69 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include "hiptensor/internal/types.hpp"
+#include "reduction_solution.hpp"
+#include "reduction_solution_instances.hpp"
+
+namespace hiptensor
+{
+    void ReductionSolutionInstances::genReductionSolution5x2F64F64Instances()
+    {
+        // add entries to mSolutionQuery
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      5,
+                                                      2,
+                                                      HIPTENSOR_OP_ADD,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      5,
+                                                      2,
+                                                      HIPTENSOR_OP_MUL,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      5,
+                                                      2,
+                                                      HIPTENSOR_OP_MIN,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      5,
+                                                      2,
+                                                      HIPTENSOR_OP_MAX,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+    }
+} // namespace hiptensor
diff --git a/library/src/reduction/reduction_solution_5_3_f32_f32_instance.cpp b/library/src/reduction/reduction_solution_5_3_f32_f32_instance.cpp
new file mode 100644
index 00000000..605676c8
--- /dev/null
+++ b/library/src/reduction/reduction_solution_5_3_f32_f32_instance.cpp
@@ -0,0 +1,70 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include "hiptensor/internal/types.hpp"
+#include "reduction_solution.hpp"
+#include "reduction_solution_instances.hpp"
+
+namespace hiptensor
+{
+    void ReductionSolutionInstances::genReductionSolution5x3F32F32Instances()
+    {
+
+        // add entries to mSolutionQuery
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      5,
+                                                      3,
+                                                      HIPTENSOR_OP_ADD,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      5,
+                                                      3,
+                                                      HIPTENSOR_OP_MUL,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      5,
+                                                      3,
+                                                      HIPTENSOR_OP_MIN,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      5,
+                                                      3,
+                                                      HIPTENSOR_OP_MAX,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+    }
+} // namespace hiptensor
diff --git a/library/src/reduction/reduction_solution_5_3_f64_f64_instance.cpp b/library/src/reduction/reduction_solution_5_3_f64_f64_instance.cpp
new file mode 100644
index 00000000..561b943f
--- /dev/null
+++ b/library/src/reduction/reduction_solution_5_3_f64_f64_instance.cpp
@@ -0,0 +1,69 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include "hiptensor/internal/types.hpp"
+#include "reduction_solution.hpp"
+#include "reduction_solution_instances.hpp"
+
+namespace hiptensor
+{
+    void ReductionSolutionInstances::genReductionSolution5x3F64F64Instances()
+    {
+        // add entries to mSolutionQuery
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      5,
+                                                      3,
+                                                      HIPTENSOR_OP_ADD,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      5,
+                                                      3,
+                                                      HIPTENSOR_OP_MUL,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      5,
+                                                      3,
+                                                      HIPTENSOR_OP_MIN,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      5,
+                                                      3,
+                                                      HIPTENSOR_OP_MAX,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+    }
+} // namespace hiptensor
diff --git a/library/src/reduction/reduction_solution_5_4_f32_f32_instance.cpp b/library/src/reduction/reduction_solution_5_4_f32_f32_instance.cpp
new file mode 100644
index 00000000..09dd7ffe
--- /dev/null
+++ b/library/src/reduction/reduction_solution_5_4_f32_f32_instance.cpp
@@ -0,0 +1,70 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include "hiptensor/internal/types.hpp"
+#include "reduction_solution.hpp"
+#include "reduction_solution_instances.hpp"
+
+namespace hiptensor
+{
+    void ReductionSolutionInstances::genReductionSolution5x4F32F32Instances()
+    {
+
+        // add entries to mSolutionQuery
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      5,
+                                                      4,
+                                                      HIPTENSOR_OP_ADD,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      5,
+                                                      4,
+                                                      HIPTENSOR_OP_MUL,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      5,
+                                                      4,
+                                                      HIPTENSOR_OP_MIN,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      5,
+                                                      4,
+                                                      HIPTENSOR_OP_MAX,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+    }
+} // namespace hiptensor
diff --git a/library/src/reduction/reduction_solution_5_4_f64_f64_instance.cpp b/library/src/reduction/reduction_solution_5_4_f64_f64_instance.cpp
new file mode 100644
index 00000000..2e9a0e28
--- /dev/null
+++ b/library/src/reduction/reduction_solution_5_4_f64_f64_instance.cpp
@@ -0,0 +1,69 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include "hiptensor/internal/types.hpp"
+#include "reduction_solution.hpp"
+#include "reduction_solution_instances.hpp"
+
+namespace hiptensor
+{
+    void ReductionSolutionInstances::genReductionSolution5x4F64F64Instances()
+    {
+        // add entries to mSolutionQuery
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      5,
+                                                      4,
+                                                      HIPTENSOR_OP_ADD,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      5,
+                                                      4,
+                                                      HIPTENSOR_OP_MUL,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      5,
+                                                      4,
+                                                      HIPTENSOR_OP_MIN,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      5,
+                                                      4,
+                                                      HIPTENSOR_OP_MAX,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+    }
+} // namespace hiptensor
diff --git a/library/src/reduction/reduction_solution_5_5_f32_f32_instance.cpp b/library/src/reduction/reduction_solution_5_5_f32_f32_instance.cpp
new file mode 100644
index 00000000..b2fc0423
--- /dev/null
+++ b/library/src/reduction/reduction_solution_5_5_f32_f32_instance.cpp
@@ -0,0 +1,70 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include "hiptensor/internal/types.hpp"
+#include "reduction_solution.hpp"
+#include "reduction_solution_instances.hpp"
+
+namespace hiptensor
+{
+    void ReductionSolutionInstances::genReductionSolution5x5F32F32Instances()
+    {
+
+        // add entries to mSolutionQuery
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      5,
+                                                      5,
+                                                      HIPTENSOR_OP_ADD,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      5,
+                                                      5,
+                                                      HIPTENSOR_OP_MUL,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      5,
+                                                      5,
+                                                      HIPTENSOR_OP_MIN,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      5,
+                                                      5,
+                                                      HIPTENSOR_OP_MAX,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+    }
+} // namespace hiptensor
diff --git a/library/src/reduction/reduction_solution_5_5_f64_f64_instance.cpp b/library/src/reduction/reduction_solution_5_5_f64_f64_instance.cpp
new file mode 100644
index 00000000..acab9126
--- /dev/null
+++ b/library/src/reduction/reduction_solution_5_5_f64_f64_instance.cpp
@@ -0,0 +1,69 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include "hiptensor/internal/types.hpp"
+#include "reduction_solution.hpp"
+#include "reduction_solution_instances.hpp"
+
+namespace hiptensor
+{
+    void ReductionSolutionInstances::genReductionSolution5x5F64F64Instances()
+    {
+        // add entries to mSolutionQuery
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      5,
+                                                      5,
+                                                      HIPTENSOR_OP_ADD,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      5,
+                                                      5,
+                                                      HIPTENSOR_OP_MUL,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      5,
+                                                      5,
+                                                      HIPTENSOR_OP_MIN,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      5,
+                                                      5,
+                                                      HIPTENSOR_OP_MAX,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+    }
+} // namespace hiptensor
diff --git a/library/src/reduction/reduction_solution_6_1_f32_f32_instance.cpp b/library/src/reduction/reduction_solution_6_1_f32_f32_instance.cpp
new file mode 100644
index 00000000..acedaf4e
--- /dev/null
+++ b/library/src/reduction/reduction_solution_6_1_f32_f32_instance.cpp
@@ -0,0 +1,70 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include "hiptensor/internal/types.hpp"
+#include "reduction_solution.hpp"
+#include "reduction_solution_instances.hpp"
+
+namespace hiptensor
+{
+    void ReductionSolutionInstances::genReductionSolution6x1F32F32Instances()
+    {
+
+        // add entries to mSolutionQuery
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      6,
+                                                      1,
+                                                      HIPTENSOR_OP_ADD,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      6,
+                                                      1,
+                                                      HIPTENSOR_OP_MUL,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      6,
+                                                      1,
+                                                      HIPTENSOR_OP_MIN,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      6,
+                                                      1,
+                                                      HIPTENSOR_OP_MAX,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+    }
+} // namespace hiptensor
diff --git a/library/src/reduction/reduction_solution_6_1_f64_f64_instance.cpp b/library/src/reduction/reduction_solution_6_1_f64_f64_instance.cpp
new file mode 100644
index 00000000..e282ed89
--- /dev/null
+++ b/library/src/reduction/reduction_solution_6_1_f64_f64_instance.cpp
@@ -0,0 +1,69 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include "hiptensor/internal/types.hpp"
+#include "reduction_solution.hpp"
+#include "reduction_solution_instances.hpp"
+
+namespace hiptensor
+{
+    void ReductionSolutionInstances::genReductionSolution6x1F64F64Instances()
+    {
+        // add entries to mSolutionQuery
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      6,
+                                                      1,
+                                                      HIPTENSOR_OP_ADD,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      6,
+                                                      1,
+                                                      HIPTENSOR_OP_MUL,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      6,
+                                                      1,
+                                                      HIPTENSOR_OP_MIN,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      6,
+                                                      1,
+                                                      HIPTENSOR_OP_MAX,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+    }
+} // namespace hiptensor
diff --git a/library/src/reduction/reduction_solution_6_2_f32_f32_instance.cpp b/library/src/reduction/reduction_solution_6_2_f32_f32_instance.cpp
new file mode 100644
index 00000000..13a3acd3
--- /dev/null
+++ b/library/src/reduction/reduction_solution_6_2_f32_f32_instance.cpp
@@ -0,0 +1,70 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include "hiptensor/internal/types.hpp"
+#include "reduction_solution.hpp"
+#include "reduction_solution_instances.hpp"
+
+namespace hiptensor
+{
+    void ReductionSolutionInstances::genReductionSolution6x2F32F32Instances()
+    {
+
+        // add entries to mSolutionQuery
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      6,
+                                                      2,
+                                                      HIPTENSOR_OP_ADD,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      6,
+                                                      2,
+                                                      HIPTENSOR_OP_MUL,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      6,
+                                                      2,
+                                                      HIPTENSOR_OP_MIN,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      6,
+                                                      2,
+                                                      HIPTENSOR_OP_MAX,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+    }
+} // namespace hiptensor
diff --git a/library/src/reduction/reduction_solution_6_2_f64_f64_instance.cpp b/library/src/reduction/reduction_solution_6_2_f64_f64_instance.cpp
new file mode 100644
index 00000000..3cbf8aed
--- /dev/null
+++ b/library/src/reduction/reduction_solution_6_2_f64_f64_instance.cpp
@@ -0,0 +1,69 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include "hiptensor/internal/types.hpp"
+#include "reduction_solution.hpp"
+#include "reduction_solution_instances.hpp"
+
+namespace hiptensor
+{
+    void ReductionSolutionInstances::genReductionSolution6x2F64F64Instances()
+    {
+        // add entries to mSolutionQuery
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      6,
+                                                      2,
+                                                      HIPTENSOR_OP_ADD,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      6,
+                                                      2,
+                                                      HIPTENSOR_OP_MUL,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      6,
+                                                      2,
+                                                      HIPTENSOR_OP_MIN,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      6,
+                                                      2,
+                                                      HIPTENSOR_OP_MAX,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+    }
+} // namespace hiptensor
diff --git a/library/src/reduction/reduction_solution_6_3_f32_f32_instance.cpp b/library/src/reduction/reduction_solution_6_3_f32_f32_instance.cpp
new file mode 100644
index 00000000..8bb2c3e3
--- /dev/null
+++ b/library/src/reduction/reduction_solution_6_3_f32_f32_instance.cpp
@@ -0,0 +1,70 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include "hiptensor/internal/types.hpp"
+#include "reduction_solution.hpp"
+#include "reduction_solution_instances.hpp"
+
+namespace hiptensor
+{
+    void ReductionSolutionInstances::genReductionSolution6x3F32F32Instances()
+    {
+
+        // add entries to mSolutionQuery
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      6,
+                                                      3,
+                                                      HIPTENSOR_OP_ADD,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      6,
+                                                      3,
+                                                      HIPTENSOR_OP_MUL,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      6,
+                                                      3,
+                                                      HIPTENSOR_OP_MIN,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      6,
+                                                      3,
+                                                      HIPTENSOR_OP_MAX,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+    }
+} // namespace hiptensor
diff --git a/library/src/reduction/reduction_solution_6_3_f64_f64_instance.cpp b/library/src/reduction/reduction_solution_6_3_f64_f64_instance.cpp
new file mode 100644
index 00000000..701479fe
--- /dev/null
+++ b/library/src/reduction/reduction_solution_6_3_f64_f64_instance.cpp
@@ -0,0 +1,69 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include "hiptensor/internal/types.hpp"
+#include "reduction_solution.hpp"
+#include "reduction_solution_instances.hpp"
+
+namespace hiptensor
+{
+    void ReductionSolutionInstances::genReductionSolution6x3F64F64Instances()
+    {
+        // add entries to mSolutionQuery
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      6,
+                                                      3,
+                                                      HIPTENSOR_OP_ADD,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      6,
+                                                      3,
+                                                      HIPTENSOR_OP_MUL,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      6,
+                                                      3,
+                                                      HIPTENSOR_OP_MIN,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      6,
+                                                      3,
+                                                      HIPTENSOR_OP_MAX,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+    }
+} // namespace hiptensor
diff --git a/library/src/reduction/reduction_solution_6_4_f32_f32_instance.cpp b/library/src/reduction/reduction_solution_6_4_f32_f32_instance.cpp
new file mode 100644
index 00000000..3c0d024a
--- /dev/null
+++ b/library/src/reduction/reduction_solution_6_4_f32_f32_instance.cpp
@@ -0,0 +1,70 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include "hiptensor/internal/types.hpp"
+#include "reduction_solution.hpp"
+#include "reduction_solution_instances.hpp"
+
+namespace hiptensor
+{
+    void ReductionSolutionInstances::genReductionSolution6x4F32F32Instances()
+    {
+
+        // add entries to mSolutionQuery
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      6,
+                                                      4,
+                                                      HIPTENSOR_OP_ADD,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      6,
+                                                      4,
+                                                      HIPTENSOR_OP_MUL,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      6,
+                                                      4,
+                                                      HIPTENSOR_OP_MIN,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      6,
+                                                      4,
+                                                      HIPTENSOR_OP_MAX,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+    }
+} // namespace hiptensor
diff --git a/library/src/reduction/reduction_solution_6_4_f64_f64_instance.cpp b/library/src/reduction/reduction_solution_6_4_f64_f64_instance.cpp
new file mode 100644
index 00000000..314067e9
--- /dev/null
+++ b/library/src/reduction/reduction_solution_6_4_f64_f64_instance.cpp
@@ -0,0 +1,69 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include "hiptensor/internal/types.hpp"
+#include "reduction_solution.hpp"
+#include "reduction_solution_instances.hpp"
+
+namespace hiptensor
+{
+    void ReductionSolutionInstances::genReductionSolution6x4F64F64Instances()
+    {
+        // add entries to mSolutionQuery
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      6,
+                                                      4,
+                                                      HIPTENSOR_OP_ADD,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      6,
+                                                      4,
+                                                      HIPTENSOR_OP_MUL,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      6,
+                                                      4,
+                                                      HIPTENSOR_OP_MIN,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      6,
+                                                      4,
+                                                      HIPTENSOR_OP_MAX,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+    }
+} // namespace hiptensor
diff --git a/library/src/reduction/reduction_solution_6_5_f32_f32_instance.cpp b/library/src/reduction/reduction_solution_6_5_f32_f32_instance.cpp
new file mode 100644
index 00000000..f38cd786
--- /dev/null
+++ b/library/src/reduction/reduction_solution_6_5_f32_f32_instance.cpp
@@ -0,0 +1,70 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include "hiptensor/internal/types.hpp"
+#include "reduction_solution.hpp"
+#include "reduction_solution_instances.hpp"
+
+namespace hiptensor
+{
+    void ReductionSolutionInstances::genReductionSolution6x5F32F32Instances()
+    {
+
+        // add entries to mSolutionQuery
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      6,
+                                                      5,
+                                                      HIPTENSOR_OP_ADD,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      6,
+                                                      5,
+                                                      HIPTENSOR_OP_MUL,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      6,
+                                                      5,
+                                                      HIPTENSOR_OP_MIN,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      6,
+                                                      5,
+                                                      HIPTENSOR_OP_MAX,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+    }
+} // namespace hiptensor
diff --git a/library/src/reduction/reduction_solution_6_5_f64_f64_instance.cpp b/library/src/reduction/reduction_solution_6_5_f64_f64_instance.cpp
new file mode 100644
index 00000000..5ae0b7f5
--- /dev/null
+++ b/library/src/reduction/reduction_solution_6_5_f64_f64_instance.cpp
@@ -0,0 +1,69 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include "hiptensor/internal/types.hpp"
+#include "reduction_solution.hpp"
+#include "reduction_solution_instances.hpp"
+
+namespace hiptensor
+{
+    void ReductionSolutionInstances::genReductionSolution6x5F64F64Instances()
+    {
+        // add entries to mSolutionQuery
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      6,
+                                                      5,
+                                                      HIPTENSOR_OP_ADD,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      6,
+                                                      5,
+                                                      HIPTENSOR_OP_MUL,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      6,
+                                                      5,
+                                                      HIPTENSOR_OP_MIN,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      6,
+                                                      5,
+                                                      HIPTENSOR_OP_MAX,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+    }
+} // namespace hiptensor
diff --git a/library/src/reduction/reduction_solution_6_6_f32_f32_instance.cpp b/library/src/reduction/reduction_solution_6_6_f32_f32_instance.cpp
new file mode 100644
index 00000000..f95473a2
--- /dev/null
+++ b/library/src/reduction/reduction_solution_6_6_f32_f32_instance.cpp
@@ -0,0 +1,70 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include "hiptensor/internal/types.hpp"
+#include "reduction_solution.hpp"
+#include "reduction_solution_instances.hpp"
+
+namespace hiptensor
+{
+    void ReductionSolutionInstances::genReductionSolution6x6F32F32Instances()
+    {
+
+        // add entries to mSolutionQuery
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      6,
+                                                      6,
+                                                      HIPTENSOR_OP_ADD,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      6,
+                                                      6,
+                                                      HIPTENSOR_OP_MUL,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      6,
+                                                      6,
+                                                      HIPTENSOR_OP_MIN,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      hiptensor::float32_t,
+                                                      6,
+                                                      6,
+                                                      HIPTENSOR_OP_MAX,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+    }
+} // namespace hiptensor
diff --git a/library/src/reduction/reduction_solution_6_6_f64_f64_instance.cpp b/library/src/reduction/reduction_solution_6_6_f64_f64_instance.cpp
new file mode 100644
index 00000000..027925b6
--- /dev/null
+++ b/library/src/reduction/reduction_solution_6_6_f64_f64_instance.cpp
@@ -0,0 +1,69 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include "hiptensor/internal/types.hpp"
+#include "reduction_solution.hpp"
+#include "reduction_solution_instances.hpp"
+
+namespace hiptensor
+{
+    void ReductionSolutionInstances::genReductionSolution6x6F64F64Instances()
+    {
+        // add entries to mSolutionQuery
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      6,
+                                                      6,
+                                                      HIPTENSOR_OP_ADD,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      6,
+                                                      6,
+                                                      HIPTENSOR_OP_MUL,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      6,
+                                                      6,
+                                                      HIPTENSOR_OP_MIN,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      hiptensor::float64_t,
+                                                      6,
+                                                      6,
+                                                      HIPTENSOR_OP_MAX,
+                                                      true, // PropagateNan,
+                                                      false>()); // OutputIndex,
+    }
+} // namespace hiptensor
diff --git a/library/src/reduction/reduction_solution_impl.hpp b/library/src/reduction/reduction_solution_impl.hpp
index 04685013..9f44bbfb 100644
--- a/library/src/reduction/reduction_solution_impl.hpp
+++ b/library/src/reduction/reduction_solution_impl.hpp
@@ -100,12 +100,17 @@ namespace hiptensor
                 return 0;
             }
 
-            std::array<ck::index_t, Traits::TensorRank>                              arrInLengths;
-            std::array<ck::index_t, Traits::TensorRank>                              arrInStrides;
-            std::array<ck::index_t, Traits::TensorRank - Traits::TensorNumReduceDim> arrOutLengths;
-            std::array<ck::index_t, Traits::TensorRank - Traits::TensorNumReduceDim> arrOutStrides;
-            std::array<ck::index_t, Traits::TensorNumReduceDim>                      reduceDims;
-            auto                                                                     toCKArr
+            static_assert(Traits::TensorRank >= Traits::TensorNumReduceDim, "TensorRank must be greater than or equal to TensorNumReduceDim");
+            constexpr ck::index_t OutputDim
+                = (Traits::TensorRank - Traits::TensorNumReduceDim)
+                      ? (Traits::TensorRank - Traits::TensorNumReduceDim)
+                      : 1;
+            std::array<ck::index_t, Traits::TensorRank>         arrInLengths;
+            std::array<ck::index_t, Traits::TensorRank>         arrInStrides;
+            std::array<ck::index_t, OutputDim>                  arrOutLengths;
+            std::array<ck::index_t, OutputDim>                  arrOutStrides;
+            std::array<ck::index_t, Traits::TensorNumReduceDim> reduceDims;
+            auto                                                toCKArr
                 = [](auto const& v, auto& a) { std::copy(v.cbegin(), v.cend(), a.begin()); };
             auto findReduceModes
                 = [](const std::vector<int32_t>& modeA, const std::vector<int32_t> modeD) {
@@ -113,7 +118,7 @@ namespace hiptensor
                       for(int i = 0; i < modeA.size(); i++)
                       {
                           if(auto it = std::find(modeD.cbegin(), modeD.cend(), modeA[i]);
-                             it != modeD.cend())
+                             it == modeD.cend())
                           {
                               reduceModes.push_back(i);
                           }
@@ -131,36 +136,10 @@ namespace hiptensor
             // toCKArr( c_strides.empty() ? hiptensor::stridesFromLengths(c_lengths, HIPTENSOR_DATA_LAYOUT_COL_MAJOR) : c_strides, arrOutStrides);
             toCKArr(findReduceModes(a_modes, c_modes), reduceDims);
 
-            // TODO test
-            auto pa = [](std::string msg, auto& a) {
-                std::cout << msg << "\n";
-                for(auto item : a)
-                {
-                    std::cout << item << " ";
-                }
-                std::cout << "\n";
-            };
-            std::cout << "alpha, beta: " << alpha << ", " << beta << "\n";
-            pa("arrInLengths", arrInLengths);
-            pa("arrInStrides", arrInStrides);
-            pa("arrOutLengths", arrOutLengths);
-            pa("arrOutStrides", arrOutStrides);
-            pa("reduceDims", reduceDims);
-            // end test
-
-            // TODO hardcode ReduceOpId
-            constexpr ck::ReduceTensorOp ReduceOpId = ck::ReduceTensorOp::ADD;
-            using InElementwiseOperation =
-                typename ck::reduce_unary_operator<ReduceOpId, true, true>::InElementwiseOperation;
-            using AccElementwiseOperation =
-                typename ck::reduce_unary_operator<ReduceOpId, true, true>::AccElementwiseOperation;
-            InElementwiseOperation  in_elementwise_op;
-            AccElementwiseOperation acc_elementwise_op;
-
-            std::tie(in_elementwise_op, acc_elementwise_op)
-                = ck::reduce_unary_operator<ReduceOpId, true, true>::GetElementwiseOperator(
-                    hiptensor::elementsFromLengths(a_lengths)
-                    / hiptensor::elementsFromLengths(c_lengths));
+            auto [in_elementwise_op, acc_elementwise_op]
+                = reductionUnaryOperators(opReduce,
+                                          hiptensor::elementsFromLengths(a_lengths)
+                                              / hiptensor::elementsFromLengths(c_lengths));
 
             Base::mInvokerArgPtr = std::move(deviceOp->MakeArgumentPointer(arrInLengths,
                                                                            arrInStrides,
@@ -209,9 +188,9 @@ namespace hiptensor
         constexpr auto ReduceOpId = convertHiptensorReduceOperatorToCk<opReduce>();
 
         using ReduceOperation = typename ck::reduce_binary_operator<ReduceOpId>::opType;
-        using InElementwiseOp =
+        using InElementwiseOperation =
             typename ck::reduce_unary_operator<ReduceOpId, true, true>::InElementwiseOperation;
-        using AccElementwiseOp =
+        using AccElementwiseOperation =
             typename ck::reduce_unary_operator<ReduceOpId, true, true>::AccElementwiseOperation;
 
         using DeviceOp    = ck::tensor_operation::device::DeviceReduce<InDataType,
@@ -220,8 +199,8 @@ namespace hiptensor
                                                                     Rank,
                                                                     NumReduceDim,
                                                                     ReduceOperation,
-                                                                    InElementwiseOp,
-                                                                    AccElementwiseOp,
+                                                                    InElementwiseOperation,
+                                                                    AccElementwiseOperation,
                                                                     PropagateNan,
                                                                     OutputIndex>;
         using DeviceOpPtr = ck::tensor_operation::device::DeviceReducePtr<InDataType,
@@ -230,8 +209,8 @@ namespace hiptensor
                                                                           Rank,
                                                                           NumReduceDim,
                                                                           ReduceOperation,
-                                                                          InElementwiseOp,
-                                                                          AccElementwiseOp,
+                                                                          InElementwiseOperation,
+                                                                          AccElementwiseOperation,
                                                                           PropagateNan,
                                                                           OutputIndex>;
 
@@ -243,8 +222,8 @@ namespace hiptensor
             Rank,
             NumReduceDim,
             ReduceOperation,
-            InElementwiseOp,
-            AccElementwiseOp,
+            InElementwiseOperation,
+            AccElementwiseOperation,
             PropagateNan,
             OutputIndex>(opPtrs);
 
diff --git a/library/src/reduction/reduction_solution_instances.cpp b/library/src/reduction/reduction_solution_instances.cpp
index 1f17bbb1..40cbc604 100644
--- a/library/src/reduction/reduction_solution_instances.cpp
+++ b/library/src/reduction/reduction_solution_instances.cpp
@@ -37,135 +37,51 @@ namespace hiptensor
     {
 
         // add entries to mSolutionQuery
-#if 0
-        registerSolutions(enumerateReductionSolutions<hiptensor::float16_t,
-                                                      hiptensor::float16_t,
-                                                      hiptensor::float16_t,
-                                                      4,
-                                                      2,
-                                                      HIPTENSOR_OP_ADD,
-                                                      true, // PropagateNan,
-                                                      false>()); // OutputIndex,
-        registerSolutions(enumerateReductionSolutions<hiptensor::float16_t,
-                                                      hiptensor::float16_t,
-                                                      hiptensor::float16_t,
-                                                      4,
-                                                      2,
-                                                      HIPTENSOR_OP_MUL,
-                                                      true, // PropagateNan,
-                                                      false>()); // OutputIndex,
-        registerSolutions(enumerateReductionSolutions<hiptensor::float16_t,
-                                                      hiptensor::float16_t,
-                                                      hiptensor::float16_t,
-                                                      4,
-                                                      2,
-                                                      HIPTENSOR_OP_MIN,
-                                                      true, // PropagateNan,
-                                                      false>()); // OutputIndex,
-        registerSolutions(enumerateReductionSolutions<hiptensor::float16_t,
-                                                      hiptensor::float16_t,
-                                                      hiptensor::float16_t,
-                                                      4,
-                                                      2,
-                                                      HIPTENSOR_OP_MAX,
-                                                      true, // PropagateNan,
-                                                      false>()); // OutputIndex,
-        registerSolutions(enumerateReductionSolutions<hiptensor::bfloat16_t,
-                                                      hiptensor::bfloat16_t,
-                                                      hiptensor::bfloat16_t,
-                                                      4,
-                                                      2,
-                                                      HIPTENSOR_OP_ADD,
-                                                      true, // PropagateNan,
-                                                      false>()); // OutputIndex,
-        registerSolutions(enumerateReductionSolutions<hiptensor::bfloat16_t,
-                                                      hiptensor::bfloat16_t,
-                                                      hiptensor::bfloat16_t,
-                                                      4,
-                                                      2,
-                                                      HIPTENSOR_OP_MUL,
-                                                      true, // PropagateNan,
-                                                      false>()); // OutputIndex,
-        registerSolutions(enumerateReductionSolutions<hiptensor::bfloat16_t,
-                                                      hiptensor::bfloat16_t,
-                                                      hiptensor::bfloat16_t,
-                                                      4,
-                                                      2,
-                                                      HIPTENSOR_OP_MIN,
-                                                      true, // PropagateNan,
-                                                      false>()); // OutputIndex,
-        registerSolutions(enumerateReductionSolutions<hiptensor::bfloat16_t,
-                                                      hiptensor::bfloat16_t,
-                                                      hiptensor::bfloat16_t,
-                                                      4,
-                                                      2,
-                                                      HIPTENSOR_OP_MAX,
-                                                      true, // PropagateNan,
-                                                      false>()); // OutputIndex,
-#endif
-        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
-                                                      hiptensor::float32_t,
-                                                      hiptensor::float32_t,
-                                                      4,
-                                                      2,
-                                                      HIPTENSOR_OP_ADD,
-                                                      true, // PropagateNan,
-                                                      false>()); // OutputIndex,
-        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
-                                                      hiptensor::float32_t,
-                                                      hiptensor::float32_t,
-                                                      4,
-                                                      2,
-                                                      HIPTENSOR_OP_MUL,
-                                                      true, // PropagateNan,
-                                                      false>()); // OutputIndex,
-        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
-                                                      hiptensor::float32_t,
-                                                      hiptensor::float32_t,
-                                                      4,
-                                                      2,
-                                                      HIPTENSOR_OP_MIN,
-                                                      true, // PropagateNan,
-                                                      false>()); // OutputIndex,
-        registerSolutions(enumerateReductionSolutions<hiptensor::float32_t,
-                                                      hiptensor::float32_t,
-                                                      hiptensor::float32_t,
-                                                      4,
-                                                      2,
-                                                      HIPTENSOR_OP_MAX,
-                                                      true, // PropagateNan,
-                                                      false>()); // OutputIndex,
-        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
-                                                      hiptensor::float64_t,
-                                                      hiptensor::float64_t,
-                                                      4,
-                                                      2,
-                                                      HIPTENSOR_OP_ADD,
-                                                      true, // PropagateNan,
-                                                      false>()); // OutputIndex,
-        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
-                                                      hiptensor::float64_t,
-                                                      hiptensor::float64_t,
-                                                      4,
-                                                      2,
-                                                      HIPTENSOR_OP_MUL,
-                                                      true, // PropagateNan,
-                                                      false>()); // OutputIndex,
-        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
-                                                      hiptensor::float64_t,
-                                                      hiptensor::float64_t,
-                                                      4,
-                                                      2,
-                                                      HIPTENSOR_OP_MIN,
-                                                      true, // PropagateNan,
-                                                      false>()); // OutputIndex,
-        registerSolutions(enumerateReductionSolutions<hiptensor::float64_t,
-                                                      hiptensor::float64_t,
-                                                      hiptensor::float64_t,
-                                                      4,
-                                                      2,
-                                                      HIPTENSOR_OP_MAX,
-                                                      true, // PropagateNan,
-                                                      false>()); // OutputIndex,
+
+        // @todo Add all reference instances. How to stop the explosion of number of instances? Wait for CK
+        // genReductionSolutionF16F16Instances();
+        // genReductionSolutionBf16Bf16Instances();
+        genReductionSolution1x1F32F32Instances();
+        genReductionSolution2x1F32F32Instances();
+        genReductionSolution2x2F32F32Instances();
+        genReductionSolution3x1F32F32Instances();
+        genReductionSolution3x2F32F32Instances();
+        genReductionSolution3x3F32F32Instances();
+        genReductionSolution4x1F32F32Instances();
+        genReductionSolution4x2F32F32Instances();
+        genReductionSolution4x3F32F32Instances();
+        genReductionSolution4x4F32F32Instances();
+        genReductionSolution5x1F32F32Instances();
+        genReductionSolution5x2F32F32Instances();
+        genReductionSolution5x3F32F32Instances();
+        genReductionSolution5x4F32F32Instances();
+        genReductionSolution5x5F32F32Instances();
+        genReductionSolution6x1F32F32Instances();
+        genReductionSolution6x2F32F32Instances();
+        genReductionSolution6x3F32F32Instances();
+        genReductionSolution6x4F32F32Instances();
+        genReductionSolution6x5F32F32Instances();
+        genReductionSolution6x6F32F32Instances();
+        genReductionSolution1x1F64F64Instances();
+        genReductionSolution2x1F64F64Instances();
+        genReductionSolution2x2F64F64Instances();
+        genReductionSolution3x1F64F64Instances();
+        genReductionSolution3x2F64F64Instances();
+        genReductionSolution3x3F64F64Instances();
+        genReductionSolution4x1F64F64Instances();
+        genReductionSolution4x2F64F64Instances();
+        genReductionSolution4x3F64F64Instances();
+        genReductionSolution4x4F64F64Instances();
+        genReductionSolution5x1F64F64Instances();
+        genReductionSolution5x2F64F64Instances();
+        genReductionSolution5x3F64F64Instances();
+        genReductionSolution5x4F64F64Instances();
+        genReductionSolution5x5F64F64Instances();
+        genReductionSolution6x1F64F64Instances();
+        genReductionSolution6x2F64F64Instances();
+        genReductionSolution6x3F64F64Instances();
+        genReductionSolution6x4F64F64Instances();
+        genReductionSolution6x5F64F64Instances();
+        genReductionSolution6x6F64F64Instances();
     }
 } // namespace hiptensor
diff --git a/library/src/reduction/reduction_solution_instances.hpp b/library/src/reduction/reduction_solution_instances.hpp
index df06902f..15ae4863 100644
--- a/library/src/reduction/reduction_solution_instances.hpp
+++ b/library/src/reduction/reduction_solution_instances.hpp
@@ -51,6 +51,51 @@ namespace hiptensor
         ReductionSolutionInstances(ReductionSolutionInstances&&)                 = delete;
         ReductionSolutionInstances& operator=(ReductionSolutionInstances const&) = delete;
         ReductionSolutionInstances& operator=(ReductionSolutionInstances&&)      = delete;
+
+        // void genReductionSolutionF16F16Instances();
+        // void genReductionSolutionBf16Bf16Instances();
+        void genReductionSolution1x1F32F32Instances();
+        void genReductionSolution2x1F32F32Instances();
+        void genReductionSolution2x2F32F32Instances();
+        void genReductionSolution3x1F32F32Instances();
+        void genReductionSolution3x2F32F32Instances();
+        void genReductionSolution3x3F32F32Instances();
+        void genReductionSolution4x1F32F32Instances();
+        void genReductionSolution4x2F32F32Instances();
+        void genReductionSolution4x3F32F32Instances();
+        void genReductionSolution4x4F32F32Instances();
+        void genReductionSolution5x1F32F32Instances();
+        void genReductionSolution5x2F32F32Instances();
+        void genReductionSolution5x3F32F32Instances();
+        void genReductionSolution5x4F32F32Instances();
+        void genReductionSolution5x5F32F32Instances();
+        void genReductionSolution6x1F32F32Instances();
+        void genReductionSolution6x2F32F32Instances();
+        void genReductionSolution6x3F32F32Instances();
+        void genReductionSolution6x4F32F32Instances();
+        void genReductionSolution6x5F32F32Instances();
+        void genReductionSolution6x6F32F32Instances();
+        void genReductionSolution1x1F64F64Instances();
+        void genReductionSolution2x1F64F64Instances();
+        void genReductionSolution2x2F64F64Instances();
+        void genReductionSolution3x1F64F64Instances();
+        void genReductionSolution3x2F64F64Instances();
+        void genReductionSolution3x3F64F64Instances();
+        void genReductionSolution4x1F64F64Instances();
+        void genReductionSolution4x2F64F64Instances();
+        void genReductionSolution4x3F64F64Instances();
+        void genReductionSolution4x4F64F64Instances();
+        void genReductionSolution5x1F64F64Instances();
+        void genReductionSolution5x2F64F64Instances();
+        void genReductionSolution5x3F64F64Instances();
+        void genReductionSolution5x4F64F64Instances();
+        void genReductionSolution5x5F64F64Instances();
+        void genReductionSolution6x1F64F64Instances();
+        void genReductionSolution6x2F64F64Instances();
+        void genReductionSolution6x3F64F64Instances();
+        void genReductionSolution6x4F64F64Instances();
+        void genReductionSolution6x5F64F64Instances();
+        void genReductionSolution6x6F64F64Instances();
     };
 
 } // namespace hiptensor
diff --git a/library/src/reduction/reduction_solution_params_impl.hpp b/library/src/reduction/reduction_solution_params_impl.hpp
index 6abdc8a8..b1c13851 100644
--- a/library/src/reduction/reduction_solution_params_impl.hpp
+++ b/library/src/reduction/reduction_solution_params_impl.hpp
@@ -117,7 +117,7 @@ namespace hiptensor
             }
             else
             {
-                // MetaTraitsT::TensorReduceOperation has only 4 options
+                // MetaTraitsT::TensorReduceOperation has only 4 options : ADD, MUL, MIN, MAX
                 return HIPTENSOR_OP_MAX;
             }
         }
diff --git a/library/src/reduction/reduction_types.hpp b/library/src/reduction/reduction_types.hpp
index 1a55fce9..e4c3dc1d 100644
--- a/library/src/reduction/reduction_types.hpp
+++ b/library/src/reduction/reduction_types.hpp
@@ -27,6 +27,7 @@
 #ifndef HIPTENSOR_REDUCTION_TYPES_HPP
 #define HIPTENSOR_REDUCTION_TYPES_HPP
 
+#include <ck/tensor_operation/gpu/device/reduction_operator_mapping.hpp>
 #include <ck/utility/reduction_enums.hpp>
 #include <hiptensor/hiptensor_types.hpp>
 
@@ -61,5 +62,29 @@ namespace hiptensor
                                                                             : HIPTENSOR_OP_MAX;
         return reduceOpId;
     }
+    inline auto reductionUnaryOperators(hiptensorOperator_t opReduce, int32_t reduceLength)
+    {
+        if(opReduce == HIPTENSOR_OP_ADD)
+        {
+            return ck::reduce_unary_operator<ck::ReduceTensorOp::ADD, true, true>::
+                GetElementwiseOperator(reduceLength);
+        }
+        else if(opReduce == HIPTENSOR_OP_MUL)
+        {
+            return ck::reduce_unary_operator<ck::ReduceTensorOp::MUL, true, true>::
+                GetElementwiseOperator(reduceLength);
+        }
+        else if(opReduce == HIPTENSOR_OP_MIN)
+        {
+            return ck::reduce_unary_operator<ck::ReduceTensorOp::MIN, true, true>::
+                GetElementwiseOperator(reduceLength);
+        }
+        else
+        {
+            return ck::reduce_unary_operator<ck::ReduceTensorOp::MAX, true, true>::
+                GetElementwiseOperator(reduceLength);
+        }
+    }
+
 } // namespace hiptensor
 #endif // HIPTENSOR_REDUCTION_TYPES_HPP
diff --git a/samples/03_reduction/reduction.cpp b/samples/03_reduction/reduction.cpp
index b2c5e82c..05992129 100644
--- a/samples/03_reduction/reduction.cpp
+++ b/samples/03_reduction/reduction.cpp
@@ -62,7 +62,7 @@ int main()
      **********************/
 
     std::vector<int32_t> modeA{'m', 'h', 'k', 'v'};
-    std::vector<int32_t> modeC{'k', 'm'};
+    std::vector<int32_t> modeC{'k', 'v'};
     int32_t              nmodeA = modeA.size();
     int32_t              nmodeC = modeC.size();
 
@@ -207,7 +207,6 @@ int main()
         hiptensorPrintElementsToFile(tensorC, C, elementsC, ", ");
         tensorC.close();
     }
-
 #endif
 
     CHECK_HIPTENSOR_ERROR(hiptensorDestroy(handle));
diff --git a/test/03_reduction/CMakeLists.txt b/test/03_reduction/CMakeLists.txt
new file mode 100644
index 00000000..efd941db
--- /dev/null
+++ b/test/03_reduction/CMakeLists.txt
@@ -0,0 +1,61 @@
+###############################################################################
+ #
+ # MIT License
+ #
+ # Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ #
+ # 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.
+ #
+ ###############################################################################
+set(ReductionCommonSources ${HIPTENSOR_COMMON_TEST_SOURCES}
+                             ${CMAKE_CURRENT_SOURCE_DIR}/reduction_resource.cpp
+                             ${CMAKE_CURRENT_SOURCE_DIR}/reduction_test.cpp)
+
+#  tests
+set (ReductionTestSources ${ReductionCommonSources}
+    ${CMAKE_CURRENT_SOURCE_DIR}/reduction_cpu_impl_test.cpp
+    )
+set (ReductionConfig ${CMAKE_CURRENT_SOURCE_DIR}/configs/reduction_params.yaml)
+add_hiptensor_test(reduction_cpu_impl_test ${ReductionConfig} ${ReductionTestSources})
+
+set (ReductionRank2TestSources ${ReductionCommonSources}
+                                 ${CMAKE_CURRENT_SOURCE_DIR}/rank2_reduction_test.cpp)
+set (ReductionRank2TestConfig  ${CMAKE_CURRENT_SOURCE_DIR}/configs/rank2_test_params.yaml)
+add_hiptensor_test(rank2_reduction_test ${ReductionRank2TestConfig}  ${ReductionRank2TestSources})
+
+set (ReductionRank3TestSources ${ReductionCommonSources}
+                                 ${CMAKE_CURRENT_SOURCE_DIR}/rank3_reduction_test.cpp)
+set (ReductionRank3TestConfig  ${CMAKE_CURRENT_SOURCE_DIR}/configs/rank3_test_params.yaml)
+add_hiptensor_test(rank3_reduction_test ${ReductionRank3TestConfig}  ${ReductionRank3TestSources})
+
+set (ReductionRank4TestSources ${ReductionCommonSources}
+                                 ${CMAKE_CURRENT_SOURCE_DIR}/rank4_reduction_test.cpp)
+set (ReductionRank4TestConfig  ${CMAKE_CURRENT_SOURCE_DIR}/configs/rank4_test_params.yaml)
+add_hiptensor_test(rank4_reduction_test ${ReductionRank4TestConfig}  ${ReductionRank4TestSources})
+
+set (ReductionRank5TestSources ${ReductionCommonSources}
+                                 ${CMAKE_CURRENT_SOURCE_DIR}/rank5_reduction_test.cpp)
+set (ReductionRank5TestConfig  ${CMAKE_CURRENT_SOURCE_DIR}/configs/rank5_test_params.yaml)
+add_hiptensor_test(rank5_reduction_test ${ReductionRank5TestConfig}  ${ReductionRank5TestSources})
+
+set (ReductionRank6TestSources ${ReductionCommonSources}
+                                 ${CMAKE_CURRENT_SOURCE_DIR}/rank6_reduction_test.cpp)
+set (ReductionRank6TestConfig  ${CMAKE_CURRENT_SOURCE_DIR}/configs/rank6_test_params.yaml)
+add_hiptensor_test(rank6_reduction_test ${ReductionRank6TestConfig}  ${ReductionRank6TestSources})
+
diff --git a/test/03_reduction/configs/rank2_test_params.yaml b/test/03_reduction/configs/rank2_test_params.yaml
new file mode 100644
index 00000000..24fe7de5
--- /dev/null
+++ b/test/03_reduction/configs/rank2_test_params.yaml
@@ -0,0 +1,29 @@
+---
+Log Level:       [ HIPTENSOR_LOG_LEVEL_ERROR, HIPTENSOR_LOG_LEVEL_PERF_TRACE ]
+Tensor Data Types:
+  # - [ HIP_R_16F, HIP_R_16F]
+  # - [ HIP_R_16F, HIP_R_32F]
+  # - [ HIP_R_16BF, HIP_R_16BF]
+  # - [ HIP_R_16BF, HIP_R_32F]
+  - [ HIP_R_64F, HIP_R_64F]
+  - [ HIP_R_32F, HIP_R_32F]
+Alphas:
+  - 1.0
+  - 2.3
+Betas:
+  - 0.0
+  # - 1.0 # failed if beta != 0
+  # - 2.3
+Lengths:
+  - [ 13, 11]
+  - [ 6, 4]
+Operators:
+  - HIPTENSOR_OP_ADD
+  - HIPTENSOR_OP_MUL
+  - HIPTENSOR_OP_MIN
+  - HIPTENSOR_OP_MAX
+Reduced Dims:
+  # - [0, 1, 2, 3] # failed if 3 exists
+  - [0]
+  - [1]
+...
diff --git a/test/03_reduction/configs/rank3_test_params.yaml b/test/03_reduction/configs/rank3_test_params.yaml
new file mode 100644
index 00000000..1ba00e21
--- /dev/null
+++ b/test/03_reduction/configs/rank3_test_params.yaml
@@ -0,0 +1,29 @@
+---
+Log Level:       [ HIPTENSOR_LOG_LEVEL_ERROR, HIPTENSOR_LOG_LEVEL_PERF_TRACE ]
+Tensor Data Types:
+  # - [ HIP_R_16F, HIP_R_16F]
+  # - [ HIP_R_16F, HIP_R_32F]
+  # - [ HIP_R_16BF, HIP_R_16BF]
+  # - [ HIP_R_16BF, HIP_R_32F]
+  - [ HIP_R_64F, HIP_R_64F]
+  - [ HIP_R_32F, HIP_R_32F]
+Alphas:
+  - 1.0
+  - 2.3
+Betas:
+  - 0.0
+  # - 1.0 # failed if beta != 0
+  # - 2.3
+Lengths:
+  - [ 11, 8, 5]
+  - [ 5, 6, 4]
+Operators:
+  - HIPTENSOR_OP_ADD
+  - HIPTENSOR_OP_MUL
+  - HIPTENSOR_OP_MIN
+  - HIPTENSOR_OP_MAX
+Reduced Dims:
+  # - [0, 1, 2] # failed if 3 exists
+  - [0, 1]
+  - [0]
+...
diff --git a/test/03_reduction/configs/rank4_test_params.yaml b/test/03_reduction/configs/rank4_test_params.yaml
new file mode 100644
index 00000000..4d8930d5
--- /dev/null
+++ b/test/03_reduction/configs/rank4_test_params.yaml
@@ -0,0 +1,33 @@
+---
+Log Level:       [ HIPTENSOR_LOG_LEVEL_ERROR, HIPTENSOR_LOG_LEVEL_PERF_TRACE ]
+Tensor Data Types:
+  # - [ HIP_R_16F, HIP_R_16F]
+  # - [ HIP_R_16F, HIP_R_32F]
+  # - [ HIP_R_16BF, HIP_R_16BF]
+  # - [ HIP_R_16BF, HIP_R_32F]
+  - [ HIP_R_64F, HIP_R_64F]
+  - [ HIP_R_32F, HIP_R_32F]
+Alphas:
+  - 1.0
+  - 2.3
+Betas:
+  - 0.0
+  # - 1.0 # failed if beta != 0
+  # - 2.3
+Lengths:
+  - [ 13, 11, 8, 5]
+  - [ 3, 5, 6, 4]
+Operators:
+  - HIPTENSOR_OP_ADD
+  - HIPTENSOR_OP_MUL
+  - HIPTENSOR_OP_MIN
+  - HIPTENSOR_OP_MAX
+Reduced Dims:
+  # - [0, 1, 2, 3] # failed if 3 exists
+  - [0, 1, 2]
+  - [0, 1]
+  - [0]
+  - [1, 2]
+  - [0, 1]
+  # - [1, 0] # failed
+...
diff --git a/test/03_reduction/configs/rank5_test_params.yaml b/test/03_reduction/configs/rank5_test_params.yaml
new file mode 100644
index 00000000..7fc416ef
--- /dev/null
+++ b/test/03_reduction/configs/rank5_test_params.yaml
@@ -0,0 +1,33 @@
+---
+Log Level:       [ HIPTENSOR_LOG_LEVEL_ERROR, HIPTENSOR_LOG_LEVEL_PERF_TRACE ]
+Tensor Data Types:
+  # - [ HIP_R_16F, HIP_R_16F]
+  # - [ HIP_R_16F, HIP_R_32F]
+  # - [ HIP_R_16BF, HIP_R_16BF]
+  # - [ HIP_R_16BF, HIP_R_32F]
+  - [ HIP_R_64F, HIP_R_64F]
+  - [ HIP_R_32F, HIP_R_32F]
+Alphas:
+  - 1.0
+  - 2.3
+Betas:
+  - 0.0
+  # - 1.0 # failed if beta != 0
+  # - 2.3
+Lengths:
+  - [ 13, 11, 8, 5, 4]
+  - [ 3, 5, 6, 4, 5]
+Operators:
+  - HIPTENSOR_OP_ADD
+  - HIPTENSOR_OP_MUL
+  - HIPTENSOR_OP_MIN
+  - HIPTENSOR_OP_MAX
+Reduced Dims:
+  - [0, 1, 2, 3]
+  - [0, 1, 2]
+  - [0, 1]
+  - [0]
+  - [1, 2]
+  - [0, 1]
+  # - [1, 0] # failed
+...
diff --git a/test/03_reduction/configs/rank6_test_params.yaml b/test/03_reduction/configs/rank6_test_params.yaml
new file mode 100644
index 00000000..583d9009
--- /dev/null
+++ b/test/03_reduction/configs/rank6_test_params.yaml
@@ -0,0 +1,33 @@
+---
+Log Level:       [ HIPTENSOR_LOG_LEVEL_ERROR, HIPTENSOR_LOG_LEVEL_PERF_TRACE ]
+Tensor Data Types:
+  # - [ HIP_R_16F, HIP_R_16F]
+  # - [ HIP_R_16F, HIP_R_32F]
+  # - [ HIP_R_16BF, HIP_R_16BF]
+  # - [ HIP_R_16BF, HIP_R_32F]
+  - [ HIP_R_64F, HIP_R_64F]
+  - [ HIP_R_32F, HIP_R_32F]
+Alphas:
+  - 1.0
+  - 2.3
+Betas:
+  - 0.0
+  # - 1.0 # failed if beta != 0
+  # - 2.3
+Lengths:
+  - [ 7, 2, 8, 5, 4, 3]
+  - [ 3, 5, 6, 4, 5, 3]
+Operators:
+  - HIPTENSOR_OP_ADD
+  - HIPTENSOR_OP_MUL
+  - HIPTENSOR_OP_MIN
+  - HIPTENSOR_OP_MAX
+Reduced Dims:
+  - [0, 1, 2, 3, 4]
+  - [0, 1, 2, 3]
+  - [0, 1, 2]
+  - [0, 1]
+  - [0]
+  - [1, 2]
+  - [0, 1]
+...
diff --git a/test/03_reduction/configs/reduction_params.yaml b/test/03_reduction/configs/reduction_params.yaml
new file mode 100644
index 00000000..adda799a
--- /dev/null
+++ b/test/03_reduction/configs/reduction_params.yaml
@@ -0,0 +1,27 @@
+---
+Log Level:       [ HIPTENSOR_LOG_LEVEL_ERROR, HIPTENSOR_LOG_LEVEL_PERF_TRACE ]
+Tensor Data Types:
+  - [ HIP_R_16F, HIP_R_16F]
+  - [ HIP_R_16F, HIP_R_32F]
+  - [ HIP_R_16BF, HIP_R_16BF]
+  - [ HIP_R_16BF, HIP_R_32F]
+  - [ HIP_R_32F, HIP_R_32F]
+  - [ HIP_R_64F, HIP_R_64F]
+Alphas:
+  - 1.0
+  - 2.3
+Betas:
+  - 1.0
+  - 2.3
+Lengths:
+  - [ 1, 1]
+  - [ 5, 2]
+  - [ 3, 4]
+  - [ 15, 12]
+  - [ 23, 11]
+Operators:
+  - HIPTENSOR_OP_ADD
+Reduced Dims:
+  - [0, 1]
+  - [1, 0]
+...
diff --git a/test/03_reduction/rank2_reduction_test.cpp b/test/03_reduction/rank2_reduction_test.cpp
new file mode 100644
index 00000000..7e15ee94
--- /dev/null
+++ b/test/03_reduction/rank2_reduction_test.cpp
@@ -0,0 +1,49 @@
+
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2021-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include <hiptensor/hiptensor.hpp>
+#include <hiptensor/hiptensor_types.hpp>
+
+#include "reduction_test.hpp"
+#include "reduction_test_helpers.hpp"
+
+class Rank2ReductionTest : public hiptensor::ReductionTest
+{
+};
+
+TEST_P(Rank2ReductionTest, RunKernel)
+{
+    static bool ranWarmup = false;
+    if(!ranWarmup)
+    {
+        this->Warmup();
+        ranWarmup = true;
+    }
+    this->RunKernel();
+}
+
+INSTANTIATE_TEST_SUITE_P(ReductionTests, Rank2ReductionTest, load_config_helper());
diff --git a/test/03_reduction/rank3_reduction_test.cpp b/test/03_reduction/rank3_reduction_test.cpp
new file mode 100644
index 00000000..a83764e3
--- /dev/null
+++ b/test/03_reduction/rank3_reduction_test.cpp
@@ -0,0 +1,49 @@
+
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2021-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include <hiptensor/hiptensor.hpp>
+#include <hiptensor/hiptensor_types.hpp>
+
+#include "reduction_test.hpp"
+#include "reduction_test_helpers.hpp"
+
+class Rank3ReductionTest : public hiptensor::ReductionTest
+{
+};
+
+TEST_P(Rank3ReductionTest, RunKernel)
+{
+    static bool ranWarmup = false;
+    if(!ranWarmup)
+    {
+        this->Warmup();
+        ranWarmup = true;
+    }
+    this->RunKernel();
+}
+
+INSTANTIATE_TEST_SUITE_P(ReductionTests, Rank3ReductionTest, load_config_helper());
diff --git a/test/03_reduction/rank4_reduction_test.cpp b/test/03_reduction/rank4_reduction_test.cpp
new file mode 100644
index 00000000..5e000ac0
--- /dev/null
+++ b/test/03_reduction/rank4_reduction_test.cpp
@@ -0,0 +1,49 @@
+
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2021-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include <hiptensor/hiptensor.hpp>
+#include <hiptensor/hiptensor_types.hpp>
+
+#include "reduction_test.hpp"
+#include "reduction_test_helpers.hpp"
+
+class Rank4ReductionTest : public hiptensor::ReductionTest
+{
+};
+
+TEST_P(Rank4ReductionTest, RunKernel)
+{
+    static bool ranWarmup = false;
+    if(!ranWarmup)
+    {
+        this->Warmup();
+        ranWarmup = true;
+    }
+    this->RunKernel();
+}
+
+INSTANTIATE_TEST_SUITE_P(ReductionTests, Rank4ReductionTest, load_config_helper());
diff --git a/test/03_reduction/rank5_reduction_test.cpp b/test/03_reduction/rank5_reduction_test.cpp
new file mode 100644
index 00000000..96edcbdf
--- /dev/null
+++ b/test/03_reduction/rank5_reduction_test.cpp
@@ -0,0 +1,49 @@
+
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2021-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include <hiptensor/hiptensor.hpp>
+#include <hiptensor/hiptensor_types.hpp>
+
+#include "reduction_test.hpp"
+#include "reduction_test_helpers.hpp"
+
+class Rank5ReductionTest : public hiptensor::ReductionTest
+{
+};
+
+TEST_P(Rank5ReductionTest, RunKernel)
+{
+    static bool ranWarmup = false;
+    if(!ranWarmup)
+    {
+        this->Warmup();
+        ranWarmup = true;
+    }
+    this->RunKernel();
+}
+
+INSTANTIATE_TEST_SUITE_P(ReductionTests, Rank5ReductionTest, load_config_helper());
diff --git a/test/03_reduction/rank6_reduction_test.cpp b/test/03_reduction/rank6_reduction_test.cpp
new file mode 100644
index 00000000..0940b9bb
--- /dev/null
+++ b/test/03_reduction/rank6_reduction_test.cpp
@@ -0,0 +1,49 @@
+
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2021-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#include <hiptensor/hiptensor.hpp>
+#include <hiptensor/hiptensor_types.hpp>
+
+#include "reduction_test.hpp"
+#include "reduction_test_helpers.hpp"
+
+class Rank6ReductionTest : public hiptensor::ReductionTest
+{
+};
+
+TEST_P(Rank6ReductionTest, RunKernel)
+{
+    static bool ranWarmup = false;
+    if(!ranWarmup)
+    {
+        this->Warmup();
+        ranWarmup = true;
+    }
+    this->RunKernel();
+}
+
+INSTANTIATE_TEST_SUITE_P(ReductionTests, Rank6ReductionTest, load_config_helper());
diff --git a/test/03_reduction/reduction_cpu_impl_test.cpp b/test/03_reduction/reduction_cpu_impl_test.cpp
new file mode 100644
index 00000000..90308e5e
--- /dev/null
+++ b/test/03_reduction/reduction_cpu_impl_test.cpp
@@ -0,0 +1,178 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2021-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+#include <hiptensor/hiptensor.hpp>
+
+#include "data_types.hpp"
+#include "logger.hpp"
+#include "reduction/reduction_cpu_reference.hpp"
+#include "reduction_test.hpp"
+#include "utils.hpp"
+#include "llvm/hiptensor_options.hpp"
+
+template <typename floatTypeA, typename floatTypeC, typename floatTypeCompute>
+auto reduceWithCpu(hipDataType typeA, hipDataType typeC, hiptensorComputeType_t typeCompute)
+{
+    floatTypeCompute alpha = (floatTypeCompute)1.2f;
+    floatTypeCompute beta  = (floatTypeCompute)2.1f;
+
+    std::vector<int32_t> modeA{'m', 'h', 'k', 'v'};
+    std::vector<int32_t> modeC{'k', 'v'};
+    int32_t              nmodeA = modeA.size();
+    int32_t              nmodeC = modeC.size();
+
+    std::unordered_map<int32_t, int64_t> extent;
+    extent['m'] = 3;
+    extent['h'] = 5;
+    extent['k'] = 6;
+    extent['v'] = 4;
+
+    std::vector<int64_t> extentC;
+    for(auto mode : modeC)
+        extentC.push_back(extent[mode]);
+    std::vector<int64_t> extentA;
+    for(auto mode : modeA)
+        extentA.push_back(extent[mode]);
+
+    size_t elementsA = 1;
+    for(auto mode : modeA)
+        elementsA *= extent[mode];
+    size_t elementsC = 1;
+    for(auto mode : modeC)
+        elementsC *= extent[mode];
+
+    size_t sizeA = sizeof(floatTypeA) * elementsA;
+    size_t sizeC = sizeof(floatTypeC) * elementsC;
+
+    std::vector<floatTypeA> aArray(elementsA);
+    std::vector<floatTypeC> cArray(elementsC, 1);
+    std::iota(aArray.begin(), aArray.end(), 0);
+
+#if HIPTENSOR_DATA_LAYOUT_COL_MAJOR
+    std::vector<floatTypeC> referenceArray
+        = {3026.1, 3044.1, 3062.1, 3080.1, 3098.1, 3116.1, 3134.1, 3152.1,
+           3170.1, 3188.1, 3206.1, 3224.1, 3242.1, 3260.1, 3278.1, 3296.1,
+           3314.1, 3332.1, 3350.1, 3368.1, 3386.1, 3404.1, 3422.1, 3440.1};
+#else // HIPTENSOR_DATA_LAYOUT_COL_MAJOR
+    std::vector<floatTypeC> referenceArray
+        = {115.5f,  1600.5f, 3085.5f, 4570.5f, 363.f,  1848.f, 3333.f, 4818.f,
+           610.5f,  2095.5f, 3580.5f, 5065.5f, 858.f,  2343.f, 3828.f, 5313.f,
+           1105.5f, 2590.5f, 4075.5f, 5560.5f, 1353.f, 2838.f, 4323.f, 5808.f};
+#endif // HIPTENSOR_DATA_LAYOUT_COL_MAJOR
+
+    hiptensorHandle_t* handle;
+    CHECK_HIPTENSOR_ERROR(hiptensorCreate(&handle));
+
+    hiptensorTensorDescriptor_t descA;
+    CHECK_HIPTENSOR_ERROR(hiptensorInitTensorDescriptor(
+        handle, &descA, nmodeA, extentA.data(), NULL /* stride */, typeA, HIPTENSOR_OP_IDENTITY));
+
+    hiptensorTensorDescriptor_t descC;
+    CHECK_HIPTENSOR_ERROR(hiptensorInitTensorDescriptor(
+        handle, &descC, nmodeC, extentC.data(), NULL /* stride */, typeC, HIPTENSOR_OP_IDENTITY));
+
+    const hiptensorOperator_t opReduce = HIPTENSOR_OP_ADD;
+
+    uint64_t worksize = 0;
+    CHECK_HIPTENSOR_ERROR(hiptensorReductionGetWorkspaceSize(handle,
+                                                             aArray.data(),
+                                                             &descA,
+                                                             modeA.data(),
+                                                             cArray.data(),
+                                                             &descC,
+                                                             modeC.data(),
+                                                             cArray.data(),
+                                                             &descC,
+                                                             modeC.data(),
+                                                             opReduce,
+                                                             typeCompute,
+                                                             &worksize));
+
+    double alphaValue{};
+    double betaValue{};
+    hiptensor::writeVal(&alphaValue, typeCompute, {typeCompute, alpha});
+    hiptensor::writeVal(&betaValue, typeCompute, {typeCompute, beta});
+    CHECK_HIPTENSOR_ERROR(hiptensorReductionReference((const void*)&alphaValue,
+                                                      aArray.data(),
+                                                      &descA,
+                                                      modeA.data(),
+                                                      (const void*)&betaValue,
+                                                      cArray.data(),
+                                                      &descC,
+                                                      modeC.data(),
+                                                      cArray.data(),
+                                                      &descC,
+                                                      modeC.data(),
+                                                      opReduce,
+                                                      typeCompute,
+                                                      0 /* stream */));
+
+    return compareEqual(referenceArray.data(), cArray.data(), cArray.size(), typeCompute);
+}
+
+TEST(ReductionCpuImplTest, CompareF32ResultWithReference)
+{
+    using floatTypeA       = hiptensor::float32_t;
+    using floatTypeC       = hiptensor::float32_t;
+    using floatTypeCompute = hiptensor::float32_t;
+
+    hipDataType            typeA       = HIP_R_32F;
+    hipDataType            typeC       = HIP_R_32F;
+    hiptensorComputeType_t typeCompute = HIPTENSOR_COMPUTE_32F;
+
+    auto [result, maxRelativeError]
+        = reduceWithCpu<floatTypeA, floatTypeC, floatTypeCompute>(typeA, typeC, typeCompute);
+    EXPECT_TRUE(result) << "max_relative_error: " << maxRelativeError;
+}
+
+TEST(ReductionCpuImplTest, CompareF64ResultWithReference)
+{
+    using floatTypeA       = hiptensor::float64_t;
+    using floatTypeC       = hiptensor::float64_t;
+    using floatTypeCompute = hiptensor::float64_t;
+
+    hipDataType            typeA       = HIP_R_64F;
+    hipDataType            typeC       = HIP_R_64F;
+    hiptensorComputeType_t typeCompute = HIPTENSOR_COMPUTE_64F;
+
+    auto [result, maxRelativeError]
+        = reduceWithCpu<floatTypeA, floatTypeC, floatTypeCompute>(typeA, typeC, typeCompute);
+    EXPECT_TRUE(result) << "max_relative_error: " << maxRelativeError;
+}
+
+// TEST(ReductionCpuImplTest, CompareF16ResultWithReference)
+// {
+// typedef _Float16 floatTypeA;
+// typedef _Float16 floatTypeC;
+// typedef _Float16 floatTypeCompute;
+//
+// hipDataType typeA       = HIP_R_16F;
+// hipDataType typeC       = HIP_R_16F;
+// hiptensorComputeType_t typeCompute = HIPTENSOR_COMPUTE_16F;
+//
+// auto [result, maxRelativeError]
+// = reduceWithCpu<floatTypeA, floatTypeC, floatTypeCompute>(typeA, typeC, typeCompute);
+// EXPECT_TRUE(result) << "max_relative_error: " << maxRelativeError;
+// }
diff --git a/test/03_reduction/reduction_resource.cpp b/test/03_reduction/reduction_resource.cpp
new file mode 100644
index 00000000..22e1fb32
--- /dev/null
+++ b/test/03_reduction/reduction_resource.cpp
@@ -0,0 +1,210 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2021-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#ifndef HIPTENSOR_REDUCTION_RESOURCE_IMPL_HPP
+#define HIPTENSOR_REDUCTION_RESOURCE_IMPL_HPP
+
+#include "reduction_resource.hpp"
+#include "data_types.hpp"
+#include "utils.hpp"
+
+namespace hiptensor
+{
+
+    ReductionResource::ReductionResource()
+        : HipResource()
+        , mDeviceA(Base::allocDevice(0))
+        , mDeviceC(Base::allocDevice(0))
+        , mHostA(Base::allocHost(0))
+        , mHostC(Base::allocHost(0))
+        , mCurrentDataType(HIP_R_32F)
+        , mCurrentMatrixAElement(0)
+        , mCurrentAllocByteA(0)
+        , mCurrentMatrixCElement(0)
+        , mCurrentAllocByteC(0)
+    {
+    }
+
+    ReductionResource::ReductionResource(ReductionResource&& rhs)
+        : HipResource()
+        , mDeviceA(std::move(rhs.mDeviceA))
+        , mDeviceC(std::move(rhs.mDeviceC))
+        , mHostA(std::move(rhs.mHostA))
+        , mHostC(std::move(rhs.mHostC))
+        , mCurrentDataType(rhs.mCurrentDataType)
+        , mCurrentMatrixAElement(rhs.mCurrentMatrixAElement)
+        , mCurrentAllocByteA(rhs.mCurrentAllocByteA)
+        , mCurrentMatrixCElement(rhs.mCurrentMatrixCElement)
+        , mCurrentAllocByteC(rhs.mCurrentAllocByteC)
+    {
+    }
+
+    void ReductionResource::setupStorage(ProblemDims const& dimSizes,
+                                         ProblemDims const& outputSizes,
+                                         hipDataType        dataType)
+    {
+        // check buffer for A
+        auto requiredElementCountA = getProduct(dimSizes);
+        auto requiredMemorySizeA   = requiredElementCountA * hipDataTypeSize(dataType);
+
+        bool needFillDataA = false;
+        if(requiredMemorySizeA > mCurrentAllocByteA)
+        {
+            Base::reallocDeviceHostPair(mDeviceA, mHostA, requiredMemorySizeA);
+            mCurrentAllocByteA = requiredMemorySizeA;
+            needFillDataA      = true;
+        }
+        if(mCurrentDataType != dataType || mCurrentMatrixAElement < requiredElementCountA)
+        {
+            needFillDataA = true;
+        }
+        mCurrentMatrixAElement = requiredElementCountA;
+
+        // check buffer for C
+        auto requiredElementCountC = getProduct(outputSizes);
+        auto requiredMemorySizeC   = requiredElementCountC * hipDataTypeSize(dataType);
+        auto needFillDataC         = false;
+        if(requiredMemorySizeC > mCurrentAllocByteC)
+        {
+            Base::reallocDeviceHostPair(mDeviceC, mHostC, requiredMemorySizeC);
+            Base::reallocDeviceHostPair(mDeviceReference, mHostReference, requiredMemorySizeC);
+            mCurrentAllocByteC = requiredMemorySizeC;
+            needFillDataC      = true;
+        }
+        if(mCurrentDataType != dataType || mCurrentMatrixCElement < requiredElementCountC)
+        {
+            needFillDataC = true;
+        }
+        mCurrentMatrixCElement = requiredElementCountC;
+
+        mCurrentDataType = dataType;
+
+        if(needFillDataA)
+        {
+            fillRand(hostA(), deviceA(), getCurrentMatrixAMemorySize());
+        }
+
+        if(needFillDataC)
+        {
+            fillRand(hostC(), deviceC(), getCurrentMatrixCMemorySize());
+        }
+    }
+
+    void ReductionResource::reset()
+    {
+        Base::reallocDeviceHostPair(mDeviceA, mHostA, 0);
+        Base::reallocDeviceHostPair(mDeviceC, mHostC, 0);
+        Base::reallocDeviceHostPair(mDeviceReference, mHostReference, 0);
+        mCurrentDataType       = HIP_R_32F;
+        mCurrentMatrixAElement = 0;
+        mCurrentAllocByteA     = 0;
+        mCurrentMatrixCElement = 0;
+        mCurrentAllocByteC     = 0;
+    }
+
+    void ReductionResource::fillRand(HostPtrT& hostBuf, DevicePtrT& deviceBuf, size_t elementCount)
+    {
+        if(mCurrentDataType == HIP_R_16F)
+        {
+            fillLaunchKernel<float16_t>((float16_t*)deviceBuf.get(), elementCount);
+        }
+        else if(mCurrentDataType == HIP_R_16BF)
+        {
+            fillLaunchKernel<bfloat16_t>((bfloat16_t*)deviceBuf.get(), elementCount);
+        }
+        else if(mCurrentDataType == HIP_R_32F)
+        {
+            fillLaunchKernel<float32_t>((float32_t*)deviceBuf.get(), elementCount);
+        }
+        else if(mCurrentDataType == HIP_R_64F)
+        {
+            fillLaunchKernel<float64_t>((float64_t*)deviceBuf.get(), elementCount);
+        }
+        Base::copyData(hostBuf, deviceBuf, elementCount);
+    }
+
+    void ReductionResource::copyCToHost()
+    {
+        Base::copyData(hostC(), deviceC(), getCurrentMatrixCMemorySize());
+    }
+
+    void ReductionResource::copyReferenceToDevice()
+    {
+        Base::copyData(deviceReference(), hostReference(), getCurrentMatrixCMemorySize());
+    }
+
+    size_t ReductionResource::getCurrentMatrixAElement() const
+    {
+        return mCurrentMatrixAElement;
+    }
+
+    size_t ReductionResource::getCurrentMatrixAMemorySize() const
+    {
+        return mCurrentMatrixAElement * hipDataTypeSize(mCurrentDataType);
+    }
+
+    size_t ReductionResource::getCurrentMatrixCElement() const
+    {
+        return mCurrentMatrixCElement;
+    }
+
+    size_t ReductionResource::getCurrentMatrixCMemorySize() const
+    {
+        return mCurrentMatrixCElement * hipDataTypeSize(mCurrentDataType);
+    }
+
+    auto ReductionResource::hostA() -> HostPtrT&
+    {
+        return mHostA;
+    }
+
+    auto ReductionResource::hostC() -> HostPtrT&
+    {
+        return mHostC;
+    }
+
+    auto ReductionResource::hostReference() -> HostPtrT&
+    {
+        return mHostReference;
+    }
+
+    auto ReductionResource::deviceA() -> DevicePtrT&
+    {
+        return mDeviceA;
+    }
+
+    auto ReductionResource::deviceC() -> DevicePtrT&
+    {
+        return mDeviceC;
+    }
+
+    auto ReductionResource::deviceReference() -> DevicePtrT&
+    {
+        return mDeviceReference;
+    }
+} // namespace hiptensor
+
+#endif // HIPTENSOR_REDUCTION_RESOURCE_IMPL_HPP
diff --git a/test/03_reduction/reduction_resource.hpp b/test/03_reduction/reduction_resource.hpp
new file mode 100644
index 00000000..362fc0bd
--- /dev/null
+++ b/test/03_reduction/reduction_resource.hpp
@@ -0,0 +1,111 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2021-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#ifndef HIPTENSOR_REDUCTION_RESOURCE_HPP
+#define HIPTENSOR_REDUCTION_RESOURCE_HPP
+
+#include <memory>
+#include <tuple>
+
+#include "hip_resource.hpp"
+#include "singleton.hpp"
+
+// ReductionResource class is intended to manage a shared pool of resources for
+// testing hiptensor reduction kernels on the GPU.
+//
+// It minimizes the memory handling overhead for launching thousands of GPU
+// kernels by allowing re-use of existing memory allocations. Memory is only
+// re-allocated as necessary to satisfy minimum size requirements.
+//
+// The interface indicates memory ownership by this class and shall only be
+// used to access for read/write purposes.
+//
+// Currently uses HIP as the backend for device allocation.
+
+namespace hiptensor
+{
+
+    struct ReductionResource : public HipResource, public LazySingleton<ReductionResource>
+    {
+        // For static initialization
+        friend std::unique_ptr<ReductionResource> std::make_unique<ReductionResource>();
+
+        using Base = HipResource;
+
+    public:
+        using DevicePtrT = Base::DevicePtrT;
+        using HostPtrT   = Base::HostPtrT;
+
+        // N, C, W, H
+        using ProblemDims = std::vector<std::size_t>;
+
+    private: // No public instantiation except make_unique.
+             // No copy
+        ReductionResource();
+        ReductionResource(const ReductionResource&)            = delete;
+        ReductionResource& operator=(const ReductionResource&) = delete;
+
+    public:
+        ReductionResource(ReductionResource&&);
+        virtual ~ReductionResource() = default;
+
+        void setupStorage(ProblemDims const& dimSizes,
+                          ProblemDims const& outputSizes,
+                          hipDataType        dataType);
+        void fillRand(HostPtrT& hostBuf, DevicePtrT& deviceBuf, size_t elementCount);
+        void copyCToHost();
+        void copyReferenceToDevice();
+
+        HostPtrT& hostA();
+        HostPtrT& hostC();
+        HostPtrT& hostReference();
+
+        DevicePtrT& deviceA();
+        DevicePtrT& deviceC();
+        DevicePtrT& deviceReference();
+
+        size_t getCurrentMatrixAElement() const;
+        size_t getCurrentMatrixAMemorySize() const;
+        size_t getCurrentMatrixCElement() const;
+        size_t getCurrentMatrixCMemorySize() const;
+        void   reset() final;
+
+    protected:
+        DevicePtrT mDeviceA, mDeviceC, mDeviceReference;
+        HostPtrT   mHostA, mHostC, mHostReference;
+
+        hipDataType mCurrentDataType; /**< Type size of element of A/C */
+
+        size_t mCurrentMatrixAElement; /**< Element count of A */
+        size_t mCurrentAllocByteA; /**< Allocated size of memory */
+
+        size_t mCurrentMatrixCElement; /**< Element count of C */
+        size_t mCurrentAllocByteC; /**< Allocated size of memory */
+    };
+
+} // namespace hiptensor
+
+#endif // HIPTENSOR_REDUCTION_RESOURCE_HPP
diff --git a/test/03_reduction/reduction_test.cpp b/test/03_reduction/reduction_test.cpp
new file mode 100644
index 00000000..e831935f
--- /dev/null
+++ b/test/03_reduction/reduction_test.cpp
@@ -0,0 +1,419 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2021-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+#include <hiptensor/hiptensor.hpp>
+
+#include "data_types.hpp"
+#include "logger.hpp"
+#include "reduction/reduction_cpu_reference.hpp"
+#include "reduction_test.hpp"
+#include "utils.hpp"
+#include "llvm/hiptensor_options.hpp"
+
+namespace hiptensor
+{
+    /*static*/ std::stringstream ReductionTest::sAPILogBuff = std::stringstream();
+
+    static void logMessage(int32_t logLevel, const char* funcName /*=""*/, const char* msg /*=""*/)
+    {
+        ReductionTest::sAPILogBuff << msg;
+    }
+
+    ReductionTest::ReductionTest()
+        : Base()
+    {
+        reset();
+        // Handle our own outputs
+        hiptensorLoggerOpenFile("/dev/null");
+        hiptensorLoggerSetCallback(logMessage);
+    }
+
+    // Kernel run checks. Virtual as different Reduction kernels have different requirements
+    // True = run test
+    // False = skip test
+    bool ReductionTest::checkDevice(hipDataType            datatype,
+                                    hiptensorComputeType_t computeDataType) const
+    {
+        return ((datatype == HIP_R_32F || computeDataType == HIP_R_32F) && isF32Supported())
+               || ((datatype == HIP_R_64F || computeDataType == HIP_R_64F) && isF64Supported());
+    }
+
+    bool ReductionTest::checkSizes() const
+    {
+        return true;
+    }
+
+    void ReductionTest::reset()
+    {
+        handle = nullptr;
+
+        mRepeats          = 1u;
+        mRunFlag          = true;
+        mValidationResult = false;
+        mMaxRelativeError = 0.0;
+    }
+
+    ReductionResource* ReductionTest::getResource() const
+    {
+        return DataStorage::instance().get();
+    }
+
+    void ReductionTest::SetUp()
+    {
+        // reset API log buffer
+        sAPILogBuff.str(std::string());
+
+        auto param      = Base::GetParam();
+        auto testType   = std::get<0>(param);
+        auto logLevel   = std::get<1>(param);
+        auto lengths    = std::get<2>(param);
+        auto outputDims = std::get<3>(param);
+        auto alpha      = std::get<4>(param);
+        auto beta       = std::get<5>(param);
+        auto op         = std::get<6>(param);
+
+        EXPECT_TRUE((lengths.size() > 1) && (lengths.size() <= 6));
+        EXPECT_TRUE((outputDims.size() >= 1) && (outputDims.size() <= 6));
+
+        EXPECT_TRUE((op == HIPTENSOR_OP_ADD) || (op == HIPTENSOR_OP_MUL) || (op == HIPTENSOR_OP_MAX)
+                    || (op == HIPTENSOR_OP_MIN));
+
+        EXPECT_EQ(testType.size(), 2); // HIP_R_16F or HIP_R_32F
+        auto acDataType      = testType[0];
+        auto computeDataType = convertToComputeType(testType[1]);
+        EXPECT_TRUE(
+            (((acDataType == HIP_R_16F) || (acDataType == HIP_R_16BF) || (acDataType == HIP_R_32F))
+             && computeDataType == HIPTENSOR_COMPUTE_32F)
+            || (acDataType == HIP_R_64F && computeDataType == HIPTENSOR_COMPUTE_64F));
+
+        mRunFlag &= checkDevice(acDataType, computeDataType);
+        mRunFlag &= lengths.size() > outputDims.size();
+
+        if(!mRunFlag)
+        {
+            GTEST_SKIP();
+        }
+        else
+        {
+            std::vector<size_t> outputLengths;
+            for(auto dim : outputDims)
+            {
+                outputLengths.push_back(lengths[dim]);
+            }
+            getResource()->setupStorage(lengths, outputLengths, acDataType);
+
+            // set mPrintElements to true to print element
+            mPrintElements = false;
+        }
+    }
+
+    void ReductionTest::reportResults(std::ostream& stream,
+                                      hipDataType   dataType,
+                                      bool          omitSkipped,
+                                      bool          omitFailed,
+                                      bool          omitPassed) const
+    {
+        // Conditionally print outputs
+        if((mRunFlag || !omitSkipped) && (mValidationResult || !omitFailed)
+           && (!mValidationResult || !omitPassed))
+        {
+            stream << ReductionTest::sAPILogBuff.str();
+
+            if(mPrintElements)
+            {
+                auto resource = getResource();
+
+                size_t elementsA = resource->getCurrentMatrixAElement();
+                size_t elementsC = resource->getCurrentMatrixCElement();
+
+                if(dataType == HIP_R_16BF)
+                {
+                    stream << "Tensor A elements (" << elementsA << "):\n";
+                    hiptensorPrintArrayElements<bfloat16_t>(
+                        stream, (bfloat16_t*)resource->hostA().get(), elementsA);
+                    stream << std::endl;
+
+                    stream << "Tensor C elements (" << elementsC << "):\n";
+                    hiptensorPrintArrayElements<bfloat16_t>(
+                        stream, (bfloat16_t*)resource->hostC().get(), elementsC);
+                    stream << std::endl;
+
+                    stream << "Refenrence elements (" << elementsC << "):\n";
+                    hiptensorPrintArrayElements<bfloat16_t>(
+                        stream, (bfloat16_t*)resource->hostReference().get(), elementsC);
+                    stream << std::endl;
+                }
+                else if(dataType == HIP_R_16F)
+                {
+                    stream << "Tensor A elements (" << elementsA << "):\n";
+                    hiptensorPrintArrayElements<float16_t>(
+                        stream, (float16_t*)resource->hostA().get(), elementsA);
+                    stream << std::endl;
+
+                    stream << "Tensor C elements (" << elementsC << "):\n";
+                    hiptensorPrintArrayElements<float16_t>(
+                        stream, (float16_t*)resource->hostC().get(), elementsC);
+                    stream << std::endl;
+
+                    stream << "Refenrence elements (" << elementsC << "):\n";
+                    hiptensorPrintArrayElements<float16_t>(
+                        stream, (float16_t*)resource->hostReference().get(), elementsC);
+                    stream << std::endl;
+                }
+                else if(dataType == HIP_R_32F)
+                {
+                    stream << "Tensor A elements (" << elementsA << "):\n";
+                    hiptensorPrintArrayElements<float32_t>(
+                        stream, (float32_t*)resource->hostA().get(), elementsA);
+                    stream << std::endl;
+
+                    stream << "Tensor C elements (" << elementsC << "):\n";
+                    hiptensorPrintArrayElements<float32_t>(
+                        stream, (float32_t*)resource->hostC().get(), elementsC);
+                    stream << std::endl;
+
+                    stream << "Refenrence elements (" << elementsC << "):\n";
+                    hiptensorPrintArrayElements<float32_t>(
+                        stream, (float32_t*)resource->hostReference().get(), elementsC);
+                    stream << std::endl;
+                }
+                else if(dataType == HIP_R_64F)
+                {
+                    stream << "Tensor A elements (" << elementsA << "):\n";
+                    hiptensorPrintArrayElements<float64_t>(
+                        stream, (float64_t*)resource->hostA().get(), elementsA);
+                    stream << std::endl;
+
+                    stream << "Tensor C elements (" << elementsC << "):\n";
+                    hiptensorPrintArrayElements<float64_t>(
+                        stream, (float64_t*)resource->hostC().get(), elementsC);
+                    stream << std::endl;
+
+                    stream << "Refenrence elements (" << elementsC << "):\n";
+                    hiptensorPrintArrayElements<float64_t>(
+                        stream, (float64_t*)resource->hostReference().get(), elementsC);
+                    stream << std::endl;
+                }
+            }
+        }
+    }
+
+    void ReductionTest::RunKernel()
+    {
+        auto param      = Base::GetParam();
+        auto testType   = std::get<0>(param);
+        auto logLevel   = std::get<1>(param);
+        auto lengths    = std::get<2>(param);
+        auto outputDims = std::get<3>(param);
+        auto alpha      = std::get<4>(param);
+        auto beta       = std::get<5>(param);
+        auto opReduce   = std::get<6>(param);
+
+        auto acDataType      = testType[0];
+        auto computeDataType = convertToComputeType(testType[1]);
+
+        if(!mRunFlag)
+        {
+            GTEST_SKIP();
+        }
+        auto resource = getResource();
+
+        if(mRunFlag)
+        {
+            std::vector<int> modeA(lengths.size());
+            std::iota(modeA.begin(), modeA.end(), 'a');
+            std::vector<int> modeC;
+            for(auto dim : outputDims)
+            {
+                modeC.push_back(modeA[dim]);
+            }
+
+            int                              nmodeA = modeA.size();
+            int                              nmodeC = modeC.size();
+            std::unordered_map<int, int64_t> extent;
+            for(auto [modeIt, i] = std::tuple{modeA.begin(), 0}; modeIt != modeA.end();
+                ++modeIt, ++i)
+            {
+                extent[*modeIt] = lengths[i];
+            }
+
+            std::vector<int64_t> extentA;
+            for(auto mode : modeA)
+                extentA.push_back(extent[mode]);
+            std::vector<int64_t> extentC;
+            for(auto mode : modeC)
+                extentC.push_back(extent[mode]);
+
+            hiptensorStatus_t  err;
+            hiptensorHandle_t* handle;
+            CHECK_HIPTENSOR_ERROR(hiptensorCreate(&handle));
+
+            hiptensorTensorDescriptor_t descA;
+            CHECK_HIPTENSOR_ERROR(hiptensorInitTensorDescriptor(handle,
+                                                                &descA,
+                                                                nmodeA,
+                                                                extentA.data(),
+                                                                NULL /* stride */,
+                                                                acDataType,
+                                                                HIPTENSOR_OP_IDENTITY));
+
+            hiptensorTensorDescriptor_t descC;
+            CHECK_HIPTENSOR_ERROR(hiptensorInitTensorDescriptor(handle,
+                                                                &descC,
+                                                                nmodeC,
+                                                                extentC.data(),
+                                                                NULL /* stride */,
+                                                                acDataType,
+                                                                HIPTENSOR_OP_IDENTITY));
+
+            uint64_t worksize = 0;
+            CHECK_HIPTENSOR_ERROR(hiptensorReductionGetWorkspaceSize(handle,
+                                                                     resource->deviceA().get(),
+                                                                     &descA,
+                                                                     modeA.data(),
+                                                                     resource->deviceC().get(),
+                                                                     &descC,
+                                                                     modeC.data(),
+                                                                     resource->deviceC().get(),
+                                                                     &descC,
+                                                                     modeC.data(),
+                                                                     opReduce,
+                                                                     computeDataType,
+                                                                     &worksize));
+            void* work = nullptr;
+            if(worksize > 0)
+            {
+                if(hipSuccess != hipMalloc(&work, worksize))
+                {
+                    work     = nullptr;
+                    worksize = 0;
+                }
+            }
+
+            double alphaValue{};
+            double betaValue{};
+            writeVal(&alphaValue, computeDataType, {computeDataType, alpha});
+            writeVal(&betaValue, computeDataType, {computeDataType, beta});
+            CHECK_HIPTENSOR_ERROR(hiptensorReduction(handle,
+                                                     (const void*)&alphaValue,
+                                                     resource->deviceA().get(),
+                                                     &descA,
+                                                     modeA.data(),
+                                                     (const void*)&betaValue,
+                                                     resource->deviceC().get(),
+                                                     &descC,
+                                                     modeC.data(),
+                                                     resource->deviceC().get(),
+                                                     &descC,
+                                                     modeC.data(),
+                                                     opReduce,
+                                                     computeDataType,
+                                                     work,
+                                                     worksize,
+                                                     0 /* stream */));
+
+            resource->copyCToHost();
+
+            CHECK_HIPTENSOR_ERROR(hiptensorReductionReference(&alphaValue,
+                                                              resource->hostA().get(),
+                                                              &descA,
+                                                              modeA.data(),
+                                                              &betaValue,
+                                                              resource->hostReference().get(),
+                                                              &descC,
+                                                              modeC.data(),
+                                                              resource->hostReference().get(),
+                                                              &descC,
+                                                              modeC.data(),
+                                                              opReduce,
+                                                              computeDataType,
+                                                              0 /* stream */));
+            resource->copyReferenceToDevice();
+
+            if(acDataType == HIP_R_16F)
+            {
+                std::tie(mValidationResult, mMaxRelativeError)
+                    = compareEqualLaunchKernel<float16_t>(
+                        (float16_t*)resource->deviceC().get(),
+                        (float16_t*)resource->deviceReference().get(),
+                        resource->getCurrentMatrixCElement(),
+                        computeDataType);
+            }
+            else if(acDataType == HIP_R_16BF)
+            {
+                std::tie(mValidationResult, mMaxRelativeError)
+                    = compareEqualLaunchKernel<bfloat16_t>(
+                        (bfloat16_t*)resource->deviceC().get(),
+                        (bfloat16_t*)resource->deviceReference().get(),
+                        resource->getCurrentMatrixCElement(),
+                        computeDataType);
+            }
+            else if(acDataType == HIP_R_32F)
+            {
+                std::tie(mValidationResult, mMaxRelativeError)
+                    = compareEqualLaunchKernel<float32_t>(
+                        (float32_t*)resource->deviceC().get(),
+                        (float32_t*)resource->deviceReference().get(),
+                        resource->getCurrentMatrixCElement(),
+                        computeDataType);
+            }
+            else if(acDataType == HIP_R_64F)
+            {
+                std::tie(mValidationResult, mMaxRelativeError)
+                    = compareEqualLaunchKernel<float64_t>(
+                        (float64_t*)resource->deviceC().get(),
+                        (float64_t*)resource->deviceReference().get(),
+                        resource->getCurrentMatrixCElement(),
+                        computeDataType);
+            }
+        }
+
+        EXPECT_TRUE(mValidationResult) << "Max relative error: " << mMaxRelativeError;
+
+        using Options        = hiptensor::HiptensorOptions;
+        auto& loggingOptions = Options::instance();
+
+        if(!loggingOptions->omitCout())
+        {
+            reportResults(std::cout,
+                          acDataType,
+                          loggingOptions->omitSkipped(),
+                          loggingOptions->omitFailed(),
+                          loggingOptions->omitPassed());
+        }
+
+        if(loggingOptions->ostream().isOpen())
+        {
+            reportResults(loggingOptions->ostream().fstream(),
+                          acDataType,
+                          loggingOptions->omitSkipped(),
+                          loggingOptions->omitFailed(),
+                          loggingOptions->omitPassed());
+        }
+    }
+
+    void ReductionTest::TearDown() {}
+
+} // namespace hiptensor
diff --git a/test/03_reduction/reduction_test.hpp b/test/03_reduction/reduction_test.hpp
new file mode 100644
index 00000000..51cd30bd
--- /dev/null
+++ b/test/03_reduction/reduction_test.hpp
@@ -0,0 +1,107 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2021-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#ifndef HIPTENSOR_REDUCTION_TEST_HPP
+#define HIPTENSOR_REDUCTION_TEST_HPP
+
+#include <sstream>
+
+#include <hiptensor/hiptensor_types.hpp>
+
+#include "reduction_resource.hpp"
+#include "reduction_test_params.hpp"
+
+#include <gtest/gtest.h>
+
+namespace hiptensor
+{
+    static void logMessage(int32_t logLevel, const char* funcName = "", const char* msg = "");
+
+    using ReductionTestParams_t = std::tuple<typename ReductionTestParams::TestTypesT,
+                                             typename ReductionTestParams::LogLevelT,
+                                             typename ReductionTestParams::LengthsT,
+                                             typename ReductionTestParams::ReducedDimsT,
+                                             typename ReductionTestParams::AlphaT,
+                                             typename ReductionTestParams::BetaT,
+                                             typename ReductionTestParams::OperatorT>;
+    class ReductionTest : public ::testing::TestWithParam<ReductionTestParams_t>
+    {
+    protected: // Types
+        using Base = ::testing::TestWithParam<ReductionTestParams_t>;
+
+        // Shared access to Reduction storage
+        using DataStorage = ReductionResource;
+
+        friend void logMessage(int32_t, const char*, const char*);
+
+    public:
+        ReductionTest();
+        virtual ~ReductionTest() = default;
+
+    protected: // Functions
+        ReductionTest(ReductionTest&&)            = delete;
+        ReductionTest(ReductionTest const&)       = delete;
+        ReductionTest& operator=(ReductionTest&)  = delete;
+        ReductionTest& operator=(ReductionTest&&) = delete;
+
+        bool checkDevice(hipDataType datatype, hiptensorComputeType_t computeDataType) const;
+        bool checkSizes() const;
+        void reset();
+
+        ReductionResource* getResource() const;
+
+        void SetUp() final;
+        void TearDown() final;
+
+        void Warmup() {}
+        void RunKernel();
+
+        void reportResults(std::ostream& stream,
+                           hipDataType   DDataType,
+                           bool          omitSkipped,
+                           bool          omitFailed,
+                           bool          omitPassed) const;
+
+    protected:
+        // Workspace items
+        hiptensorHandle_t* handle = nullptr;
+
+        hiptensorTensorDescriptor_t a_ms_ks, b_ns_ks, c_ms_ns, d_ms_ns;
+
+        // Execution flow control
+        uint32_t mRepeats;
+        bool     mRunFlag          = true;
+        bool     mValidationResult = false;
+        bool     mPrintElements    = false;
+        double   mMaxRelativeError;
+
+        // Output buffer
+        static std::stringstream sAPILogBuff;
+    };
+
+} // namespace hiptensor
+
+#endif // HIPTENSOR_REDUCTION_TEST_HPP
diff --git a/test/03_reduction/reduction_test_helpers.hpp b/test/03_reduction/reduction_test_helpers.hpp
new file mode 100644
index 00000000..dc521504
--- /dev/null
+++ b/test/03_reduction/reduction_test_helpers.hpp
@@ -0,0 +1,78 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2021-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#ifndef HIPTENSOR_REDUCTION_TEST_HELPERS_HPP
+#define HIPTENSOR_REDUCTION_TEST_HELPERS_HPP
+
+#include <gtest/gtest.h>
+
+#include "llvm/hiptensor_options.hpp"
+#include "llvm/yaml_parser.hpp"
+
+#ifdef HIPTENSOR_TEST_YAML_INCLUDE
+#include HIPTENSOR_TEST_YAML_INCLUDE
+#define HIPTENSOR_TEST_YAML_BUNDLE 1
+#else
+#define HIPTENSOR_TEST_YAML_BUNDLE 0
+#endif // HIPTENSOR_TEST_YAML_INCLUDE
+
+auto inline load_config_helper()
+{
+    hiptensor::ReductionTestParams testParams;
+    using Options     = hiptensor::HiptensorOptions;
+    auto& testOptions = Options::instance();
+
+    if(testOptions->usingDefaultConfig() && HIPTENSOR_TEST_YAML_BUNDLE)
+    {
+        auto params = hiptensor::YamlConfigLoader<hiptensor::ReductionTestParams>::loadFromString(
+            HIPTENSOR_TEST_GET_YAML);
+        if(params)
+        {
+            testParams = params.value();
+        }
+    }
+    else
+    {
+        auto params = hiptensor::YamlConfigLoader<hiptensor::ReductionTestParams>::loadFromFile(
+            testOptions->inputFilename());
+        if(params)
+        {
+            testParams = params.value();
+        }
+    }
+
+    // testParams.printParams();
+
+    return ::testing::Combine(::testing::ValuesIn(testParams.dataTypes()),
+                              ::testing::Values(testParams.logLevelMask()),
+                              ::testing::ValuesIn(testParams.problemLengths()),
+                              ::testing::ValuesIn(testParams.reducedDims()),
+                              ::testing::ValuesIn(testParams.alphas()),
+                              ::testing::ValuesIn(testParams.betas()),
+                              ::testing::ValuesIn(testParams.operators()));
+}
+
+#endif // HIPTENSOR_REDUCTION_TEST_HELPERS_HPP
diff --git a/test/03_reduction/reduction_test_params.hpp b/test/03_reduction/reduction_test_params.hpp
new file mode 100644
index 00000000..4f538359
--- /dev/null
+++ b/test/03_reduction/reduction_test_params.hpp
@@ -0,0 +1,111 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (C) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * 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.
+ *
+ *******************************************************************************/
+
+#ifndef HIPTENSOR_REDUCTION_TEST_PARAMS_HPP
+#define HIPTENSOR_REDUCTION_TEST_PARAMS_HPP
+
+#include <tuple>
+#include <vector>
+
+#include <hiptensor/hiptensor.hpp>
+#include <hiptensor/hiptensor_types.hpp>
+
+#include "utils.hpp"
+
+namespace hiptensor
+{
+
+    struct ReductionTestParams
+    {
+        using TestTypesT   = std::vector<hipDataType>;
+        using LogLevelT    = hiptensorLogLevel_t;
+        using LengthsT     = std::vector<std::size_t>;
+        using AlphaT       = double;
+        using BetaT        = double;
+        using ReducedDimsT = std::vector<std::size_t>;
+        using OperatorT    = hiptensorOperator_t;
+
+    public:
+        std::vector<TestTypesT>& dataTypes()
+        {
+            return mDataTypes;
+        }
+
+        LogLevelT& logLevelMask()
+        {
+            return mLogLevelMask;
+        }
+
+        std::vector<LengthsT>& problemLengths()
+        {
+            return mProblemLengths;
+        }
+
+        std::vector<OperatorT>& operators()
+        {
+            return mOperators;
+        }
+
+        std::vector<ReducedDimsT>& reducedDims()
+        {
+            return mReducedDims;
+        }
+
+        std::vector<AlphaT>& alphas()
+        {
+            return mAlphas;
+        }
+
+        std::vector<BetaT>& betas()
+        {
+            return mBetas;
+        }
+
+        void printParams()
+        {
+            std::cout << "DataTypes: " << mDataTypes << "\n"
+                      << "LogLevelMask: " << mLogLevelMask << "\n"
+                      << "ProblemLengths: " << mProblemLengths << "\n"
+                      << "Operators: " << mOperators << "\n"
+                      << "Alphas: " << mAlphas << "\n"
+                      << "Betas: " << mBetas << "\n"
+                      << "ReducedDims: " << mReducedDims << "\n";
+        }
+
+    private:
+        //Data types of input and output tensors
+        std::vector<TestTypesT>   mDataTypes;
+        LogLevelT                 mLogLevelMask;
+        std::vector<LengthsT>     mProblemLengths;
+        std::vector<AlphaT>       mAlphas;
+        std::vector<BetaT>        mBetas;
+        std::vector<OperatorT>    mOperators;
+        std::vector<ReducedDimsT> mReducedDims;
+    };
+
+} // namespace hiptensor
+
+#endif // HIPTENSOR_REDUCTION_TEST_PARAMS_HPP
diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt
index 880baaeb..5a02a151 100644
--- a/test/CMakeLists.txt
+++ b/test/CMakeLists.txt
@@ -186,6 +186,7 @@ endfunction()
 add_subdirectory(00_unit)
 add_subdirectory(01_contraction)
 add_subdirectory(02_permutation)
+add_subdirectory(03_reduction)
 
 rocm_install(
     FILES "${INSTALL_TEST_FILE}"
diff --git a/test/llvm/yaml_parser_config.cpp b/test/llvm/yaml_parser_config.cpp
index 6bb0701f..177af789 100644
--- a/test/llvm/yaml_parser_config.cpp
+++ b/test/llvm/yaml_parser_config.cpp
@@ -32,6 +32,7 @@
 
 #include "01_contraction/contraction_test_params.hpp"
 #include "02_permutation/permutation_test_params.hpp"
+#include "03_reduction/reduction_test_params.hpp"
 #include "yaml_parser_impl.hpp"
 
 // Fwd declare NoneType
@@ -145,6 +146,10 @@ namespace llvm
             {
                 io.enumCase(value, "HIPTENSOR_OP_IDENTITY", HIPTENSOR_OP_IDENTITY);
                 io.enumCase(value, "HIPTENSOR_OP_SQRT", HIPTENSOR_OP_SQRT);
+                io.enumCase(value, "HIPTENSOR_OP_ADD", HIPTENSOR_OP_ADD);
+                io.enumCase(value, "HIPTENSOR_OP_MUL", HIPTENSOR_OP_MUL);
+                io.enumCase(value, "HIPTENSOR_OP_MIN", HIPTENSOR_OP_MIN);
+                io.enumCase(value, "HIPTENSOR_OP_MAX", HIPTENSOR_OP_MAX);
                 io.enumCase(value, "HIPTENSOR_OP_UNKNOWN", HIPTENSOR_OP_UNKNOWN);
             }
         };
@@ -353,6 +358,58 @@ namespace llvm
             }
         };
 
+        ///
+        // Mapping of the test param elements of ReductionTestParams for reading / writing.
+        ///
+        template <>
+        struct MappingTraits<hiptensor::ReductionTestParams>
+        {
+            static void mapping(IO& io, hiptensor::ReductionTestParams& doc)
+            {
+                // Logging bitfield
+                io.mapRequired("Log Level", doc.logLevelMask());
+
+                // Sequences of combinatorial fields
+                io.mapRequired("Tensor Data Types", doc.dataTypes());
+                io.mapRequired("Alphas", (std::vector<AlphaT>&)(doc.alphas()));
+                io.mapRequired("Betas", (std::vector<BetaT>&)(doc.betas()));
+                io.mapRequired("Lengths", doc.problemLengths());
+                io.mapRequired("Reduced Dims", doc.reducedDims());
+                io.mapRequired("Operators", (doc.operators()));
+            }
+
+            // Additional validation for input / output of the config
+            static std::string validate(IO& io, hiptensor::ReductionTestParams& doc)
+            {
+                if(doc.problemLengths().size() == 0)
+                {
+                    return "Error: Empty Lengths";
+                }
+
+                if(doc.alphas().size() == 0)
+                {
+                    return "Error: Empty Alphas";
+                }
+
+                if(doc.betas().size() == 0)
+                {
+                    return "Error: Empty Betas";
+                }
+
+                if(doc.reducedDims().size() == 0)
+                {
+                    return "Error: Empty Reduced Dims";
+                }
+
+                if(doc.operators().size() == 0)
+                {
+                    return "Error: Empty Operators";
+                }
+
+                return std::string{};
+            }
+        };
+
     } // namespace yaml
 
 } // namespace llvm
@@ -362,4 +419,5 @@ namespace hiptensor
 {
     template struct YamlConfigLoader<ContractionTestParams>;
     template struct YamlConfigLoader<PermutationTestParams>;
+    template struct YamlConfigLoader<ReductionTestParams>;
 }