Skip to content

Commit

Permalink
Merge pull request #241 from CongMa13/reduce_test
Browse files Browse the repository at this point in the history
Reduce test
  • Loading branch information
CongMa13 authored Jul 9, 2024
2 parents 36179ef + c4192f7 commit c949394
Show file tree
Hide file tree
Showing 77 changed files with 5,358 additions and 188 deletions.
49 changes: 45 additions & 4 deletions library/src/reduction/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -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})
Expand Down
3 changes: 1 addition & 2 deletions library/src/reduction/hiptensor_reduction.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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),
Expand Down
100 changes: 100 additions & 0 deletions library/src/reduction/reduction_cpu_reference.cpp
Original file line number Diff line number Diff line change
@@ -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;
}
49 changes: 49 additions & 0 deletions library/src/reduction/reduction_cpu_reference.hpp
Original file line number Diff line number Diff line change
@@ -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
139 changes: 139 additions & 0 deletions library/src/reduction/reduction_cpu_reference_impl.hpp
Original file line number Diff line number Diff line change
@@ -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
Loading

0 comments on commit c949394

Please sign in to comment.