From bf1c7eed5ecfeb28f3f85dae62c1ed9d3dea7b69 Mon Sep 17 00:00:00 2001 From: Cong Ma Date: Thu, 2 Nov 2023 14:24:53 +0000 Subject: [PATCH] Implement permutation with CK - implemented `hiptensorPermutation` with CK - added unit tests for `hiptensorPermutation` --- .gitignore | 2 +- .../hiptensor/internal/hiptensor_utility.hpp | 5 + library/src/hiptensor.cpp | 3 +- .../src/permutation/hiptensor_permutation.cpp | 131 +++++++- .../src/permutation/permutation_ck_col.hpp | 48 +++ .../permutation/permutation_ck_col_impl.hpp | 123 +++++++ .../permutation/permutation_cpu_reference.hpp | 3 - .../permutation_cpu_reference_impl.hpp | 9 +- test/01_contraction/contraction_resource.hpp | 4 +- test/01_contraction/contraction_test.cpp | 6 +- .../contraction_test_params.hpp | 2 +- test/02_permutation/CMakeLists.txt | 35 ++ test/02_permutation/configs/test_params.yaml | 41 +++ .../permutation_column_major_test.cpp} | 32 +- test/02_permutation/permutation_resource.cpp | 161 +++++++++ test/02_permutation/permutation_resource.hpp | 104 ++++++ test/02_permutation/permutation_test.cpp | 306 ++++++++++++++++++ test/02_permutation/permutation_test.hpp | 105 ++++++ .../permutation_test_helpers.hpp | 68 ++++ .../permutation_test_params.hpp | 95 ++++++ test/CMakeLists.txt | 4 +- test/device/common.hpp | 2 +- test/llvm/CMakeLists.txt | 4 +- ...tion_config.cpp => yaml_parser_config.cpp} | 47 ++- test/utils.hpp | 45 ++- 25 files changed, 1337 insertions(+), 48 deletions(-) create mode 100644 library/src/permutation/permutation_ck_col.hpp create mode 100644 library/src/permutation/permutation_ck_col_impl.hpp create mode 100644 test/02_permutation/CMakeLists.txt create mode 100644 test/02_permutation/configs/test_params.yaml rename test/{singleton.hpp => 02_permutation/permutation_column_major_test.cpp} (74%) create mode 100644 test/02_permutation/permutation_resource.cpp create mode 100644 test/02_permutation/permutation_resource.hpp create mode 100644 test/02_permutation/permutation_test.cpp create mode 100644 test/02_permutation/permutation_test.hpp create mode 100644 test/02_permutation/permutation_test_helpers.hpp create mode 100644 test/02_permutation/permutation_test_params.hpp rename test/llvm/{yaml_parser_contraction_config.cpp => yaml_parser_config.cpp} (85%) diff --git a/.gitignore b/.gitignore index ae21cc0c..ad44a303 100644 --- a/.gitignore +++ b/.gitignore @@ -9,7 +9,7 @@ hiptensor_version.hpp hiptensor-version.hpp # Generated source file -test/01_contraction/configs/*.hpp +test/*/configs/*.hpp # Precompiled Headers *.gch diff --git a/library/include/hiptensor/internal/hiptensor_utility.hpp b/library/include/hiptensor/internal/hiptensor_utility.hpp index f2df2dd2..5fd9e331 100644 --- a/library/include/hiptensor/internal/hiptensor_utility.hpp +++ b/library/include/hiptensor/internal/hiptensor_utility.hpp @@ -131,6 +131,11 @@ namespace std return os; } + static ostream& operator<<(ostream& os, const _Float16 value) + { + os << static_cast(value); + return os; + } } #endif // HIPTENSOR_UTILITY_INTERNAL_HPP diff --git a/library/src/hiptensor.cpp b/library/src/hiptensor.cpp index 6638f398..3016ef64 100644 --- a/library/src/hiptensor.cpp +++ b/library/src/hiptensor.cpp @@ -151,7 +151,8 @@ hiptensorStatus_t hiptensorInitTensorDescriptor(const hiptensorHandle_t* han return HIPTENSOR_STATUS_NOT_INITIALIZED; } - if((lens == nullptr) || ((dataType != HIP_R_32F) && (dataType != HIP_R_64F)) + if((lens == nullptr) + || ((dataType != HIP_R_16F) && (dataType != HIP_R_32F) && (dataType != HIP_R_64F)) || unaryOp != HIPTENSOR_OP_IDENTITY) { auto errorCode = HIPTENSOR_STATUS_INVALID_VALUE; diff --git a/library/src/permutation/hiptensor_permutation.cpp b/library/src/permutation/hiptensor_permutation.cpp index 493616a0..b3d60379 100644 --- a/library/src/permutation/hiptensor_permutation.cpp +++ b/library/src/permutation/hiptensor_permutation.cpp @@ -25,7 +25,8 @@ *******************************************************************************/ #include -#include "permutation_cpu_reference.hpp" +#include "logger.hpp" +#include "permutation_ck_col.hpp" hiptensorStatus_t hiptensorPermutation(const hiptensorHandle_t* handle, const void* alpha, @@ -38,15 +39,133 @@ hiptensorStatus_t hiptensorPermutation(const hiptensorHandle_t* handle const hipDataType typeScalar, const hipStream_t stream) { - assert(descA->mType == HIP_R_16F || descA->mType == HIP_R_32F); - assert(descA->mType == descB->mType); + using hiptensor::Logger; + auto& logger = Logger::instance(); + + // Log API access + char msg[2048]; + snprintf(msg, + sizeof(msg), + "handle=%p, alpha=%p, A=%p, descA=%p, modeA=%p, B=%p, descB=%p, modeB=%p, " + "typeScalar=0x%02X, stream=%p", + handle, + alpha, + A, + descA, + modeA, + B, + descB, + modeB, + (unsigned int)typeScalar, + stream); + + logger->logAPITrace("hiptensorPermutation", msg); + + if(!handle || !alpha || !A || !descA || !modeA || !B || !descB || !modeB) + { + auto errorCode = HIPTENSOR_STATUS_NOT_INITIALIZED; + auto printErrorMessage = [&logger, errorCode](const std::string& paramName) { + char msg[512]; + snprintf(msg, + sizeof(msg), + "Initialization Error : %s = nullptr (%s)", + paramName.c_str(), + hiptensorGetErrorString(errorCode)); + logger->logError("hiptensorPermutation", msg); + }; + if(!handle) + { + printErrorMessage("handle"); + } + if(!alpha) + { + printErrorMessage("alpha"); + } + if(!A) + { + printErrorMessage("A"); + } + if(!descA) + { + printErrorMessage("descA"); + } + if(!modeA) + { + printErrorMessage("modeA"); + } + if(!B) + { + printErrorMessage("B"); + } + if(!descB) + { + printErrorMessage("descB"); + } + if(!modeB) + { + printErrorMessage("modeB"); + } + return errorCode; + } + + if(descA->mType != HIP_R_16F && descA->mType != HIP_R_32F) + { + auto errorCode = HIPTENSOR_STATUS_NOT_SUPPORTED; + snprintf(msg, + sizeof(msg), + "Unsupported Data Type Error : The supported data types of A and B are HIP_R_16F " + "and HIP_R_32F (%s)", + hiptensorGetErrorString(errorCode)); + logger->logError("hiptensorPermutation", msg); + return errorCode; + } + + if(descA->mType != descB->mType) + { + auto errorCode = HIPTENSOR_STATUS_INVALID_VALUE; + snprintf(msg, + sizeof(msg), + "Mismatched Data Type Error : Data types of A and B are not the same. (%s)", + hiptensorGetErrorString(errorCode)); + logger->logError("hiptensorPermutation", msg); + return errorCode; + } + + if(typeScalar != HIP_R_16F && typeScalar != HIP_R_32F) + { + auto errorCode = HIPTENSOR_STATUS_NOT_SUPPORTED; + snprintf(msg, + sizeof(msg), + "Unsupported Data Type Error : The supported data types of alpha are HIP_R_16F " + "and HIP_R_32F (%s)", + hiptensorGetErrorString(errorCode)); + logger->logError("hiptensorPermutation", msg); + return errorCode; + } + if(descA->mType == HIP_R_16F) { - return hiptensor::detail::permuteByCpu(alpha, static_cast(A), descA, modeA, static_cast<_Float16 *>(B), descB, modeB, typeScalar); + return hiptensor::detail::permuteByCk(alpha, + static_cast(A), + descA, + modeA, + static_cast<_Float16*>(B), + descB, + modeB, + typeScalar, + stream); } else if(descA->mType == HIP_R_32F) { - return hiptensor::detail::permuteByCpu(alpha, static_cast(A), descA, modeA, static_cast(B), descB, modeB, typeScalar); + return hiptensor::detail::permuteByCk(alpha, + static_cast(A), + descA, + modeA, + static_cast(B), + descB, + modeB, + typeScalar, + stream); } - return HIPTENSOR_STATUS_NOT_SUPPORTED; + return HIPTENSOR_STATUS_NOT_SUPPORTED; } diff --git a/library/src/permutation/permutation_ck_col.hpp b/library/src/permutation/permutation_ck_col.hpp new file mode 100644 index 00000000..600658c6 --- /dev/null +++ b/library/src/permutation/permutation_ck_col.hpp @@ -0,0 +1,48 @@ +/******************************************************************************* + * + * 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_PERMUTATION_CK_COL_HPP +#define HIPTENSOR_PERMUTATION_CK_COL_HPP +#include + +namespace hiptensor +{ + namespace detail + { + template + hiptensorStatus_t permuteByCk(const void* alpha, + const DataType* A, + const hiptensorTensorDescriptor_t* descA, + const int32_t modeA[], + DataType* B, + const hiptensorTensorDescriptor_t* descB, + const int32_t modeB[], + const hipDataType typeScalar); + + } +} + +#include "permutation_ck_col_impl.hpp" +#endif // HIPTENSOR_PERMUTATION_CK_COL_HPP diff --git a/library/src/permutation/permutation_ck_col_impl.hpp b/library/src/permutation/permutation_ck_col_impl.hpp new file mode 100644 index 00000000..7f16bb8a --- /dev/null +++ b/library/src/permutation/permutation_ck_col_impl.hpp @@ -0,0 +1,123 @@ +/******************************************************************************* + * + * 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_PERMUTATION_CK_COL_IMPL_HPP +#define HIPTENSOR_PERMUTATION_CK_COL_IMPL_HPP +#include + +#include +#include +#include + +#include "types.hpp" + +namespace hiptensor +{ + namespace detail + { + template + hiptensorStatus_t permuteByCk(const void* alpha, + const DataType* A, + const hiptensorTensorDescriptor_t* descA, + const int32_t modeA[], + DataType* B, + const hiptensorTensorDescriptor_t* descB, + const int32_t modeB[], + const hipDataType typeScalar, + const hipStream_t stream) + { + using PassThrough = ck::tensor_operation::element_wise::PassThrough; + using UnaryOp = ck::tensor_operation::element_wise::PassThrough; + using Scale = ck::tensor_operation::element_wise::Scale; + using DeviceElementwisePermuteInstance + = ck::tensor_operation::device::DeviceElementwiseImpl< + ck::Tuple, // InDataTypeTuple + ck::Tuple, // OutDataTypeTuple + PassThrough, // ElementwiseOp + UnaryOp, // UnaryOp + Scale, // Scalar + 4, // NumDim + 1, // MPerThread + ck::Sequence<1>, // InScalarPerVectorSeq + ck::Sequence<1>>; // OutScalarPerVectorSeq + + const auto modeSize = descA->mLengths.size(); + assert(modeSize == 4); + + std::unordered_map + modeToLength; // for example {'n': 1, 'c': 2, 'w': 3, 'h':0} + + for(int32_t index = 0; index < modeSize; index++) + { + modeToLength[modeA[index]] = descA->mLengths[index]; + } + + std::unordered_map bModeToStrides; + int32_t stride = 1; + bModeToStrides[modeB[0]] = stride; + for(int32_t index = 1; index < modeSize; index++) + { + stride *= modeToLength[modeB[index - 1]]; + bModeToStrides[modeB[index]] = stride; + } + + float alphaValue = readVal(alpha, typeScalar); + std::array input = {A}; + std::array output = {B}; + std::array a_strides + = {1, + modeToLength[modeA[0]], + modeToLength[modeA[0]] * modeToLength[modeA[1]], + modeToLength[modeA[0]] * modeToLength[modeA[1]] * modeToLength[modeA[2]]}; + std::array b_strides = {bModeToStrides[modeA[0]], + bModeToStrides[modeA[1]], + bModeToStrides[modeA[2]], + bModeToStrides[modeA[3]]}; + std::array ab_lengths = {modeToLength[modeA[0]], + modeToLength[modeA[1]], + modeToLength[modeA[2]], + modeToLength[modeA[3]]}; + auto broadcastPermute = DeviceElementwisePermuteInstance{}; + auto argument = broadcastPermute.MakeArgumentPointer(ab_lengths, + {a_strides}, + {b_strides}, + input, + output, + PassThrough{}, + UnaryOp{}, + Scale{alphaValue}); + + if(!broadcastPermute.IsSupportedArgument(argument.get())) + { + return HIPTENSOR_STATUS_NOT_SUPPORTED; + }; + + auto broadcastPermute_invoker_ptr = broadcastPermute.MakeInvokerPointer(); + broadcastPermute_invoker_ptr->Run(argument.get(), StreamConfig{stream, false}); + return HIPTENSOR_STATUS_SUCCESS; + } + } +} +#endif // HIPTENSOR_PERMUTATION_CK_COL_IMPL_HPP diff --git a/library/src/permutation/permutation_cpu_reference.hpp b/library/src/permutation/permutation_cpu_reference.hpp index f7e335f2..0bcff3cf 100644 --- a/library/src/permutation/permutation_cpu_reference.hpp +++ b/library/src/permutation/permutation_cpu_reference.hpp @@ -26,9 +26,6 @@ #ifndef HIPTENSOR_PERMUTATION_CPU_REFERENCE_HPP #define HIPTENSOR_PERMUTATION_CPU_REFERENCE_HPP - -#include - #include namespace hiptensor { diff --git a/library/src/permutation/permutation_cpu_reference_impl.hpp b/library/src/permutation/permutation_cpu_reference_impl.hpp index 6cc34e6b..a04e4176 100644 --- a/library/src/permutation/permutation_cpu_reference_impl.hpp +++ b/library/src/permutation/permutation_cpu_reference_impl.hpp @@ -26,8 +26,8 @@ #ifndef HIPTENSOR_PERMUTATION_CPU_REFERENCE_IMPL_HPP #define HIPTENSOR_PERMUTATION_CPU_REFERENCE_IMPL_HPP #include -#include #include +#include #include "permutation_cpu_reference.hpp" #include "types.hpp" @@ -63,8 +63,9 @@ namespace hiptensor { bStrides[i] = descB->mLengths[i - 1] * bStrides[i - 1]; } - auto bIndices = std::vector(modeSize, 0); - auto elementCount = hiptensor::elementsFromLengths(aLens); + auto bIndices = std::vector(modeSize, 0); + auto elementCount = hiptensor::elementsFromLengths(aLens); + float alphaValue = readVal(alpha, typeScalar); for(int elementIndex = 0; elementIndex < elementCount; elementIndex++) { auto index = elementIndex; @@ -75,7 +76,7 @@ namespace hiptensor } auto bOffset = std::inner_product(bIndices.begin(), bIndices.end(), bStrides.begin(), 0); - B[bOffset] = A[elementIndex]; + B[bOffset] = static_cast(A[elementIndex] * alphaValue); } return HIPTENSOR_STATUS_SUCCESS; diff --git a/test/01_contraction/contraction_resource.hpp b/test/01_contraction/contraction_resource.hpp index 88f13e88..2fd3e62c 100644 --- a/test/01_contraction/contraction_resource.hpp +++ b/test/01_contraction/contraction_resource.hpp @@ -30,8 +30,8 @@ #include #include -#include "../hip_resource.hpp" -#include "../singleton.hpp" +#include "hip_resource.hpp" +#include "singleton.hpp" // ContractionResource class is intended to manage a shared pool of resources for // testing hiptensor contraction kernels on the GPU. diff --git a/test/01_contraction/contraction_test.cpp b/test/01_contraction/contraction_test.cpp index 790676cd..4ad97610 100644 --- a/test/01_contraction/contraction_test.cpp +++ b/test/01_contraction/contraction_test.cpp @@ -25,12 +25,12 @@ *******************************************************************************/ #include -#include "../library/src/include/types.hpp" +#include "types.hpp" #include "llvm/hiptensor_options.hpp" -#include "../utils.hpp" -#include "contraction_cpu_reference.hpp" +#include "contraction/contraction_cpu_reference.hpp" #include "contraction_test.hpp" +#include "utils.hpp" namespace hiptensor { diff --git a/test/01_contraction/contraction_test_params.hpp b/test/01_contraction/contraction_test_params.hpp index 30c99349..29c4aa1b 100644 --- a/test/01_contraction/contraction_test_params.hpp +++ b/test/01_contraction/contraction_test_params.hpp @@ -33,7 +33,7 @@ #include #include -#include "../utils.hpp" +#include "utils.hpp" namespace hiptensor { diff --git a/test/02_permutation/CMakeLists.txt b/test/02_permutation/CMakeLists.txt new file mode 100644 index 00000000..4334901c --- /dev/null +++ b/test/02_permutation/CMakeLists.txt @@ -0,0 +1,35 @@ +############################################################################### + # + # 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(PermutationCommonSources ${HIPTENSOR_COMMON_TEST_SOURCES} + ${CMAKE_CURRENT_SOURCE_DIR}/permutation_resource.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/permutation_test.cpp) + +# tests +set (PermutationTestSources ${PermutationCommonSources} + ${CMAKE_CURRENT_SOURCE_DIR}/permutation_column_major_test.cpp) +set (PermutationTestConfig ${CMAKE_CURRENT_SOURCE_DIR}/configs/test_params.yaml) +add_hiptensor_test(permutation_test ${PermutationTestConfig} ${PermutationTestSources}) + diff --git a/test/02_permutation/configs/test_params.yaml b/test/02_permutation/configs/test_params.yaml new file mode 100644 index 00000000..f6aa92e2 --- /dev/null +++ b/test/02_permutation/configs/test_params.yaml @@ -0,0 +1,41 @@ +--- +Log Level: [ HIPTENSOR_LOG_LEVEL_ERROR, HIPTENSOR_LOG_LEVEL_PERF_TRACE ] +Tensor Data Types: + - [ HIP_R_32F, HIP_R_32F] + - [ HIP_R_16F, HIP_R_16F] + - [ HIP_R_16F, HIP_R_32F] +Alphas: + - 0.0 + - 1.0 + - 2.3 +Lengths: + - [ 1, 1, 1, 1] + - [ 5, 2, 3, 4] + - [ 5, 2, 1, 1] + - [ 15, 12, 23, 11] +Permuted Dims: + - [0, 1, 2, 3] + - [0, 1, 3, 2] + - [0, 2, 1, 3] + - [0, 2, 3, 1] + - [0, 3, 1, 2] + - [0, 3, 2, 1] + - [1, 0, 2, 3] + - [1, 0, 3, 2] + - [1, 2, 0, 3] + - [1, 2, 3, 0] + - [1, 3, 0, 2] + - [1, 3, 2, 0] + - [2, 0, 1, 3] + - [2, 0, 3, 1] + - [2, 1, 0, 3] + - [2, 1, 3, 0] + - [2, 3, 0, 1] + - [2, 3, 1, 0] + - [3, 0, 1, 2] + - [3, 0, 2, 1] + - [3, 1, 0, 2] + - [3, 1, 2, 0] + - [3, 2, 0, 1] + - [3, 2, 1, 0] +... diff --git a/test/singleton.hpp b/test/02_permutation/permutation_column_major_test.cpp similarity index 74% rename from test/singleton.hpp rename to test/02_permutation/permutation_column_major_test.cpp index eedf37c1..9a4ace70 100644 --- a/test/singleton.hpp +++ b/test/02_permutation/permutation_column_major_test.cpp @@ -24,25 +24,25 @@ * *******************************************************************************/ -#ifndef HIPTENSOR_TEST_SINGLETON_HPP -#define HIPTENSOR_TEST_SINGLETON_HPP +#include +#include -#include +#include "permutation_test.hpp" +#include "permutation_test_helpers.hpp" -namespace hiptensor +class PermutationTest : public hiptensor::PermutationTest { +}; - template - class LazySingleton +TEST_P(PermutationTest, RunKernel) +{ + static bool ranWarmup = false; + if(!ranWarmup) { - public: - static inline std::unique_ptr const& instance() - { - static auto sInstance = std::make_unique(); - return sInstance; - } - }; - -} // namespace hiptensor + this->Warmup(); + ranWarmup = true; + } + this->RunKernel(); +} -#endif // HIPTENSOR_TEST_SINGLETON_HPP +INSTANTIATE_TEST_SUITE_P(PermutationTests, PermutationTest, load_config_helper()); diff --git a/test/02_permutation/permutation_resource.cpp b/test/02_permutation/permutation_resource.cpp new file mode 100644 index 00000000..4323e01d --- /dev/null +++ b/test/02_permutation/permutation_resource.cpp @@ -0,0 +1,161 @@ +/******************************************************************************* + * + * 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_PERMUTATION_RESOURCE_IMPL_HPP +#define HIPTENSOR_PERMUTATION_RESOURCE_IMPL_HPP + +#include "permutation_resource.hpp" +#include "types.hpp" +#include "utils.hpp" + +namespace hiptensor +{ + + PermutationResource::PermutationResource() + : HipResource() + , mDeviceA(Base::allocDevice(0)) + , mDeviceB(Base::allocDevice(0)) + , mHostA(Base::allocHost(0)) + , mHostB(Base::allocHost(0)) + , mCurrentMatrixElement(0) + , mCurrentDataType(HIP_R_32F) + , mCurrentAllocByte(0) + { + } + + PermutationResource::PermutationResource(PermutationResource&& rhs) + : HipResource() + , mDeviceA(std::move(rhs.mDeviceA)) + , mDeviceB(std::move(rhs.mDeviceB)) + , mHostA(std::move(rhs.mHostA)) + , mHostB(std::move(rhs.mHostB)) + , mCurrentMatrixElement(rhs.mCurrentMatrixElement) + , mCurrentDataType(rhs.mCurrentDataType) + , mCurrentAllocByte(rhs.mCurrentAllocByte) + { + } + + void PermutationResource::setupStorage(ProblemDims const& dimSizes, hipDataType dataType) + { + auto requiredElementCount = getProduct(dimSizes); + auto requiredMemorySize = requiredElementCount * hipDataTypeSize(dataType); + + bool needFillData = false; + if(requiredMemorySize > mCurrentAllocByte) + { + Base::reallocDeviceHostPair(mDeviceA, mHostA, requiredMemorySize); + Base::reallocDeviceHostPair(mDeviceB, mHostB, requiredMemorySize); + Base::reallocDeviceHostPair(mDeviceReference, mHostReference, requiredMemorySize); + mCurrentAllocByte = requiredMemorySize; + needFillData = true; + } + else if(mCurrentDataType != dataType) + { + needFillData = true; + } + mCurrentMatrixElement = requiredElementCount; + mCurrentDataType = dataType; + if(needFillData) + { + fillRandToA(); + } + } + + void PermutationResource::reset() + { + Base::reallocDeviceHostPair(mDeviceA, mHostA, 0); + Base::reallocDeviceHostPair(mDeviceB, mHostB, 0); + Base::reallocDeviceHostPair(mDeviceReference, mHostReference, 0); + mCurrentMatrixElement = 0; + mCurrentDataType = HIP_R_32F; + mCurrentAllocByte = 0; + } + + void PermutationResource::fillRandToA() + { + if(mCurrentDataType == HIP_R_32F) + { + fillLaunchKernel((float*)deviceA().get(), mCurrentMatrixElement); + } + else + { + fillLaunchKernel<_Float16>((_Float16*)deviceA().get(), mCurrentMatrixElement); + } + Base::copyData(hostA(), deviceA(), getCurrentMatrixMemorySize()); + } + + void PermutationResource::copyBToHost() + { + Base::copyData(hostB(), deviceB(), getCurrentMatrixMemorySize()); + } + + void PermutationResource::copyReferenceToDevice() + { + Base::copyData(deviceReference(), hostReference(), getCurrentMatrixMemorySize()); + } + + size_t PermutationResource::getCurrentMatrixElement() const + { + return mCurrentMatrixElement; + } + + size_t PermutationResource::getCurrentMatrixMemorySize() const + { + return mCurrentMatrixElement * hipDataTypeSize(mCurrentDataType); + } + + auto PermutationResource::hostA() -> HostPtrT& + { + return mHostA; + } + + auto PermutationResource::hostB() -> HostPtrT& + { + return mHostB; + } + + auto PermutationResource::hostReference() -> HostPtrT& + { + return mHostReference; + } + + auto PermutationResource::deviceA() -> DevicePtrT& + { + return mDeviceA; + } + + auto PermutationResource::deviceB() -> DevicePtrT& + { + return mDeviceB; + } + + auto PermutationResource::deviceReference() -> DevicePtrT& + { + return mDeviceReference; + } +} // namespace hiptensor + +#endif // HIPTENSOR_PERMUTATION_RESOURCE_IMPL_HPP diff --git a/test/02_permutation/permutation_resource.hpp b/test/02_permutation/permutation_resource.hpp new file mode 100644 index 00000000..2e0d77ef --- /dev/null +++ b/test/02_permutation/permutation_resource.hpp @@ -0,0 +1,104 @@ +/******************************************************************************* + * + * 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_PERMUTATION_RESOURCE_HPP +#define HIPTENSOR_PERMUTATION_RESOURCE_HPP + +#include +#include + +#include "hip_resource.hpp" +#include "singleton.hpp" + +// PermutationResource class is intended to manage a shared pool of resources for +// testing hiptensor contraction 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 PermutationResource : public HipResource, public LazySingleton + { + // For static initialization + friend std::unique_ptr std::make_unique(); + + using Base = HipResource; + + public: + using DevicePtrT = Base::DevicePtrT; + using HostPtrT = Base::HostPtrT; + + // N, C, W, H + using ProblemDims = std::vector; + + private: // No public instantiation except make_unique. + // No copy + PermutationResource(); + PermutationResource(const PermutationResource&) = delete; + PermutationResource& operator=(const PermutationResource&) = delete; + + public: + PermutationResource(PermutationResource&&); + virtual ~PermutationResource() = default; + + void setupStorage(ProblemDims const& dimSizes, hipDataType dataType); + void fillRandToA(); + void copyBToHost(); + void copyReferenceToDevice(); + + HostPtrT& hostA(); + HostPtrT& hostB(); + HostPtrT& hostReference(); + + DevicePtrT& deviceA(); + DevicePtrT& deviceB(); + DevicePtrT& deviceReference(); + + size_t getCurrentMatrixElement() const; + size_t getCurrentMatrixMemorySize() const; + void reset() final; + + protected: + DevicePtrT mDeviceA, mDeviceB, mDeviceReference; + HostPtrT mHostA, mHostB, mHostReference; + + size_t mCurrentMatrixElement; /**< Element count of A/B */ + hipDataType + mCurrentDataType; /**< Type size of element of A/B, only support HIP_R_16F, HIP_R_32F */ + size_t mCurrentAllocByte; /**< Allocated size of memory */ + }; + +} // namespace hiptensor + +#endif // HIPTENSOR_PERMUTATION_RESOURCE_HPP diff --git a/test/02_permutation/permutation_test.cpp b/test/02_permutation/permutation_test.cpp new file mode 100644 index 00000000..dbb52d6e --- /dev/null +++ b/test/02_permutation/permutation_test.cpp @@ -0,0 +1,306 @@ +/******************************************************************************* + * + * 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 + +#include "logger.hpp" +#include "permutation/permutation_cpu_reference.hpp" +#include "permutation_test.hpp" +#include "types.hpp" +#include "utils.hpp" +#include "llvm/hiptensor_options.hpp" + +namespace hiptensor +{ + /*static*/ std::stringstream PermutationTest::sAPILogBuff = std::stringstream(); + + static void logMessage(int32_t logLevel, const char* funcName /*=""*/, const char* msg /*=""*/) + { + PermutationTest::sAPILogBuff << msg; + } + + PermutationTest::PermutationTest() + : Base() + { + reset(); + + // Handle our own outputs + hiptensorLoggerOpenFile("/dev/null"); + hiptensorLoggerSetCallback(logMessage); + } + + // Kernel run checks. Virtual as different Permutation kernels have different requirements + // True = run test + // False = skip test + bool PermutationTest::checkDevice(hipDataType datatype) const + { + return (isF32Supported() && datatype == HIP_R_32F) + || (isF64Supported() && datatype == HIP_R_64F); + } + + bool PermutationTest::checkSizes() const + { + return true; + } + + void PermutationTest::reset() + { + handle = nullptr; + + mRepeats = 1u; + mRunFlag = true; + mValidationResult = false; + mMaxRelativeError = 0.0; + } + + PermutationResource* PermutationTest::getResource() const + { + return DataStorage::instance().get(); + } + + void PermutationTest::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 permutedDims = std::get<3>(param); + auto alpha = std::get<4>(param); + + // 4D tensors only at the moment. + EXPECT_EQ(lengths.size(), 4); // Format {'n', 'c', 'w', 'h'} + EXPECT_EQ(permutedDims.size(), 4); // permutation of {0, 1, 2, 3} + + EXPECT_EQ(testType.size(), 2); // HIP_R_16F or HIP_R_32F + auto abDataType = testType[0]; + EXPECT_TRUE((abDataType == HIP_R_16F) || (abDataType == HIP_R_32F)); + + getResource()->setupStorage(lengths, abDataType); + + // set mPrintElements to true to print element + mPrintElements = false; + } + + void PermutationTest::reportResults(std::ostream& stream, + hipDataType dataType, + bool omitSkipped, + bool omitFailed, + bool omitPassed) const + { + // Conditionally print outputs + if((mRunFlag || !omitSkipped) && (mValidationResult || !omitFailed) + && (!mValidationResult || !omitPassed)) + { + stream << PermutationTest::sAPILogBuff.str(); + + if(mPrintElements) + { + auto resource = getResource(); + + size_t elementsA = resource->getCurrentMatrixElement(); + size_t elementsB = elementsA; + + if(dataType == HIP_R_32F) + { + stream << "Tensor A elements (" << elementsA << "):\n"; + hiptensorPrintArrayElements( + stream, (float*)resource->hostA().get(), elementsA); + stream << std::endl; + + stream << "Tensor B elements (" << elementsB << "):\n"; + hiptensorPrintArrayElements( + stream, (float*)resource->hostB().get(), elementsB); + stream << std::endl; + } + else + { + stream << "Tensor A elements (" << elementsA << "):\n"; + hiptensorPrintArrayElements<_Float16>( + stream, (_Float16*)resource->hostA().get(), elementsA); + stream << std::endl; + + stream << "Tensor B elements (" << elementsB << "):\n"; + hiptensorPrintArrayElements<_Float16>( + stream, (_Float16*)resource->hostB().get(), elementsB); + stream << std::endl; + } + } + } + } + + void PermutationTest::RunKernel() + { + auto param = Base::GetParam(); + auto testType = std::get<0>(param); + auto logLevel = std::get<1>(param); + auto lengths = std::get<2>(param); + auto permutedDims = std::get<3>(param); + auto alpha = std::get<4>(param); + + auto abDataType = testType[0]; + auto computeDataType = testType[1]; + + if(!mRunFlag) + { + GTEST_SKIP(); + } + auto resource = getResource(); + + if(mRunFlag) + { + /********************** + B_{w, h, c, n} = 1.0 * \textsl{IDENTITY}(A_{c, n, h, w}) + **********************/ + + std::vector modeA{'n', 'c', 'w', 'h'}; + std::vector modeB; + for(auto dim : permutedDims) + { + modeB.push_back(modeA[dim]); + } + + int nmodeA = modeA.size(); + int nmodeB = modeB.size(); + std::unordered_map extent; + for(auto [modeIt, i] = std::tuple{modeA.begin(), 0}; modeIt != modeA.end(); + ++modeIt, ++i) + { + extent[*modeIt] = lengths[i]; + } + + std::vector extentA; + for(auto mode : modeA) + extentA.push_back(extent[mode]); + std::vector extentB; + for(auto mode : modeB) + extentB.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 */, + abDataType, + HIPTENSOR_OP_IDENTITY)); + + hiptensorTensorDescriptor_t descB; + CHECK_HIPTENSOR_ERROR(hiptensorInitTensorDescriptor(handle, + &descB, + nmodeB, + extentB.data(), + NULL /* stride */, + abDataType, + HIPTENSOR_OP_IDENTITY)); + + float alphaValue{}; + if(computeDataType == HIP_R_16F) + { + *(reinterpret_cast<_Float16*>(&alphaValue)) = static_cast<_Float16>(alpha); + } + else + { + *(reinterpret_cast(&alphaValue)) = static_cast(alpha); + } + CHECK_HIPTENSOR_ERROR(hiptensorPermutation(handle, + &alphaValue, + resource->deviceA().get(), + &descA, + modeA.data(), + resource->deviceB().get(), + &descB, + modeB.data(), + computeDataType, + 0 /* stream */)); + resource->copyBToHost(); + + if(abDataType == HIP_R_32F) + { + hiptensor::detail::permuteByCpu(&alphaValue, + (const float*)resource->hostA().get(), + &descA, + modeA.data(), + (float*)resource->hostReference().get(), + &descB, + modeB.data(), + computeDataType); + resource->copyReferenceToDevice(); + std::tie(mValidationResult, mMaxRelativeError) + = compareEqualLaunchKernel((float*)resource->deviceB().get(), + (float*)resource->deviceReference().get(), + resource->getCurrentMatrixElement()); + } + else if(abDataType == HIP_R_16F) + { + hiptensor::detail::permuteByCpu(&alphaValue, + (const _Float16*)resource->hostA().get(), + &descA, + modeA.data(), + (_Float16*)resource->hostReference().get(), + &descB, + modeB.data(), + computeDataType); + resource->copyReferenceToDevice(); + std::tie(mValidationResult, mMaxRelativeError) = compareEqualLaunchKernel<_Float16>( + (_Float16*)resource->deviceB().get(), + (_Float16*)resource->deviceReference().get(), + resource->getCurrentMatrixElement()); + } + } + + EXPECT_TRUE(mValidationResult) << "Max relative error: " << mMaxRelativeError; + + using Options = hiptensor::HiptensorOptions; + auto& loggingOptions = Options::instance(); + + if(!loggingOptions->omitCout()) + { + reportResults(std::cout, + abDataType, + loggingOptions->omitSkipped(), + loggingOptions->omitFailed(), + loggingOptions->omitPassed()); + } + + if(loggingOptions->ostream().isOpen()) + { + reportResults(loggingOptions->ostream().fstream(), + abDataType, + loggingOptions->omitSkipped(), + loggingOptions->omitFailed(), + loggingOptions->omitPassed()); + } + } + + void PermutationTest::TearDown() {} + +} // namespace hiptensor diff --git a/test/02_permutation/permutation_test.hpp b/test/02_permutation/permutation_test.hpp new file mode 100644 index 00000000..9009d8ea --- /dev/null +++ b/test/02_permutation/permutation_test.hpp @@ -0,0 +1,105 @@ +/******************************************************************************* + * + * 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_PERMUTATION_TEST_HPP +#define HIPTENSOR_PERMUTATION_TEST_HPP + +#include + +#include + +#include "permutation_resource.hpp" +#include "permutation_test_params.hpp" + +#include + +namespace hiptensor +{ + static void logMessage(int32_t logLevel, const char* funcName = "", const char* msg = ""); + + using PermutationTestParams_t = std::tuple; + class PermutationTest : public ::testing::TestWithParam + { + protected: // Types + using Base = ::testing::TestWithParam; + + // Shared access to Permutation storage + using DataStorage = PermutationResource; + + friend void logMessage(int32_t, const char*, const char*); + + public: + PermutationTest(); + virtual ~PermutationTest() = default; + + protected: // Functions + PermutationTest(PermutationTest&&) = delete; + PermutationTest(PermutationTest const&) = delete; + PermutationTest& operator=(PermutationTest&) = delete; + PermutationTest& operator=(PermutationTest&&) = delete; + + bool checkDevice(hipDataType datatype) const; + bool checkSizes() const; + void reset(); + + PermutationResource* 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_PERMUTATION_TEST_HPP diff --git a/test/02_permutation/permutation_test_helpers.hpp b/test/02_permutation/permutation_test_helpers.hpp new file mode 100644 index 00000000..ac34a21a --- /dev/null +++ b/test/02_permutation/permutation_test_helpers.hpp @@ -0,0 +1,68 @@ +/******************************************************************************* + * + * 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_PERMUTATION_TEST_HELPERS_HPP +#define HIPTENSOR_PERMUTATION_TEST_HELPERS_HPP + +#include + +#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::PermutationTestParams testParams; + using Options = hiptensor::HiptensorOptions; + auto& testOptions = Options::instance(); + + if(testOptions->usingDefaultConfig() && HIPTENSOR_TEST_YAML_BUNDLE) + { + testParams = hiptensor::YamlConfigLoader::loadFromString( + HIPTENSOR_TEST_GET_YAML); + } + else + { + testParams = hiptensor::YamlConfigLoader::loadFromFile( + testOptions->inputFilename()); + } + + // testParams.printParams(); + + return ::testing::Combine(::testing::ValuesIn(testParams.dataTypes()), + ::testing::Values(testParams.logLevelMask()), + ::testing::ValuesIn(testParams.problemLengths()), + ::testing::ValuesIn(testParams.permutedDims()), + ::testing::ValuesIn(testParams.alphas())); +} + +#endif // HIPTENSOR_PERMUTATION_TEST_HELPERS_HPP diff --git a/test/02_permutation/permutation_test_params.hpp b/test/02_permutation/permutation_test_params.hpp new file mode 100644 index 00000000..f526ad47 --- /dev/null +++ b/test/02_permutation/permutation_test_params.hpp @@ -0,0 +1,95 @@ +/******************************************************************************* + * + * 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_PERMUTATION_TEST_PARAMS_HPP +#define HIPTENSOR_PERMUTATION_TEST_PARAMS_HPP + +#include +#include + +#include +#include + +#include "utils.hpp" + +namespace hiptensor +{ + + struct PermutationTestParams + { + using TestTypesT = std::vector; + using LogLevelT = hiptensorLogLevel_t; + using LengthsT = std::vector; + using AlphaT = double; + using PermutedDimsT = std::vector; + + public: + std::vector& dataTypes() + { + return mDataTypes; + } + + LogLevelT& logLevelMask() + { + return mLogLevelMask; + } + + std::vector& problemLengths() + { + return mProblemLengths; + } + + std::vector& permutedDims() + { + return mPermutedDims; + } + + std::vector& alphas() + { + return mAlphas; + } + + void printParams() + { + std::cout << "DataTypes: " << mDataTypes << "\n" + << "LogLevelMask: " << mLogLevelMask << "\n" + << "ProblemLengths: " << mProblemLengths << "\n" + << "Alphas: " << mAlphas << "\n" + << "PermutedDims: " << mPermutedDims << "\n"; + } + + private: + //Data types of input and output tensors + std::vector mDataTypes; + LogLevelT mLogLevelMask; + std::vector mProblemLengths; + std::vector mAlphas; + std::vector mPermutedDims; + }; + +} // namespace hiptensor + +#endif // HIPTENSOR_PERMUTATION_TEST_PARAMS_HPP diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index e9f791cd..06967a09 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -101,7 +101,8 @@ function(add_hiptensor_test BINARY_NAME YAML_CONFIG_FILE TEST_SOURCES) target_include_directories(${BINARY_NAME} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR} ${PROJECT_SOURCE_DIR}/library/include - ${PROJECT_SOURCE_DIR}/library/src/contraction + ${PROJECT_SOURCE_DIR}/library/src/include + ${PROJECT_SOURCE_DIR}/library/src ${PROJECT_SOURCE_DIR}/test ${hiptensor_test_YAML_DIR}) @@ -179,6 +180,7 @@ endfunction() add_subdirectory(00_unit) add_subdirectory(01_contraction) +add_subdirectory(02_permutation) rocm_install( FILES "${INSTALL_TEST_FILE}" diff --git a/test/device/common.hpp b/test/device/common.hpp index 8ad75dd4..f961abc1 100644 --- a/test/device/common.hpp +++ b/test/device/common.hpp @@ -66,7 +66,7 @@ __global__ static void // fill kernel for 'elementSize' elements template -__global__ void fillKernel(DataType* data, uint32_t elementSize) +__global__ void fillKernel(DataType* data, uint32_t elementSize, uint32_t seed) { uint32_t index = (blockIdx.x * blockDim.x + threadIdx.x); diff --git a/test/llvm/CMakeLists.txt b/test/llvm/CMakeLists.txt index 064746bb..6708ae2e 100644 --- a/test/llvm/CMakeLists.txt +++ b/test/llvm/CMakeLists.txt @@ -56,11 +56,13 @@ set(HIPTENSOR_LLVM_LIBS "-L${LLVM_LIBRARY_DIR}" # Includes set(HIPTENSOR_LLVM_INCLUDES ${LLVM_INCLUDE_DIRS} ${PROJECT_SOURCE_DIR}/library/include + ${PROJECT_SOURCE_DIR}/library/src/include + ${PROJECT_SOURCE_DIR}/library/src ${PROJECT_SOURCE_DIR}/test ) # Sources for this static object -set(HIPTENSOR_LLVM_SOURCES ${CMAKE_CURRENT_SOURCE_DIR}/yaml_parser_contraction_config.cpp +set(HIPTENSOR_LLVM_SOURCES ${CMAKE_CURRENT_SOURCE_DIR}/yaml_parser_config.cpp ${CMAKE_CURRENT_SOURCE_DIR}/hiptensor_options.cpp) # Create hiptensor_llvm object target diff --git a/test/llvm/yaml_parser_contraction_config.cpp b/test/llvm/yaml_parser_config.cpp similarity index 85% rename from test/llvm/yaml_parser_contraction_config.cpp rename to test/llvm/yaml_parser_config.cpp index 6982eeea..46f4c43e 100644 --- a/test/llvm/yaml_parser_contraction_config.cpp +++ b/test/llvm/yaml_parser_config.cpp @@ -31,6 +31,7 @@ #include #include "01_contraction/contraction_test_params.hpp" +#include "02_permutation/permutation_test_params.hpp" #include "yaml_parser_impl.hpp" // Fwd declare NoneType @@ -108,6 +109,7 @@ namespace llvm { static void enumeration(IO& io, hipDataType& value) { + io.enumCase(value, "HIP_R_16F", HIP_R_16F); io.enumCase(value, "HIP_R_32F", HIP_R_32F); io.enumCase(value, "HIP_R_64F", HIP_R_64F); io.enumCase(value, "NONE_TYPE", hiptensor::NONE_TYPE); @@ -209,7 +211,7 @@ namespace llvm }; /// - // Mapping of the test param elements for reading / writing. + // Mapping of the test param elements of ContractionTestParams for reading / writing. /// template <> struct MappingTraits @@ -269,12 +271,53 @@ namespace llvm } }; + /// + // Mapping of the test param elements of PermutationTestParams for reading / writing. + /// + template <> + struct MappingTraits + { + static void mapping(IO& io, hiptensor::PermutationTestParams& 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&)(doc.alphas())); + io.mapRequired("Lengths", doc.problemLengths()); + io.mapRequired("Permuted Dims", doc.permutedDims()); + } + + // Additional validation for input / output of the config + static std::string validate(IO& io, hiptensor::PermutationTestParams& doc) + { + if(doc.problemLengths().size() == 0) + { + return "Error: Empty Lengths"; + } + + if(doc.alphas().size() == 0) + { + return "Error: Empty Alphas"; + } + + if(doc.permutedDims().size() == 0) + { + return "Error: Empty Permuted Dims"; + } + + return std::string{}; + } + }; + } // namespace yaml } // namespace llvm -// Instantiate the yaml loader for the ContractionTestParams +// Instantiate the yaml loader for the ContractionTestParams and PermutationTestParams namespace hiptensor { template struct YamlConfigLoader; + template struct YamlConfigLoader; } diff --git a/test/utils.hpp b/test/utils.hpp index 421c1a04..67418568 100644 --- a/test/utils.hpp +++ b/test/utils.hpp @@ -28,6 +28,7 @@ #define HIPTENSOR_TEST_UTILS_HPP #include +#include #include #include #include @@ -49,10 +50,10 @@ CHECK_HIP_ERROR(hipFree(ptr)); \ } -#define HIPTENSOR_FREE_HOST(ptr) \ - if(ptr != nullptr) \ - { \ - free(ptr); \ +#define HIPTENSOR_FREE_HOST(ptr) \ + if(ptr != nullptr) \ + { \ + CHECK_HIP_ERROR(hipHostFree(ptr)); \ } inline bool isF32Supported() @@ -97,13 +98,30 @@ static constexpr intT1 ceilDiv(const intT1 numerator, const intT2 divisor) return (numerator + divisor - 1) / divisor; } +template +auto getProduct(const Container& container, + typename Container::value_type init = typename Container::value_type{1}) +{ + return std::accumulate(std::begin(container), + std::end(container), + init, + std::multiplies{}); +} + // fill kernel for 'elementSize' elements template __host__ static inline void fillLaunchKernel(DataType* data, uint32_t elementSize) { auto blockDim = dim3(1024, 1, 1); auto gridDim = dim3(ceilDiv(elementSize, blockDim.x), 1, 1); - hipLaunchKernelGGL((fillKernel), gridDim, blockDim, 0, 0, data, elementSize); + hipLaunchKernelGGL((fillKernel), + gridDim, + blockDim, + 0, + 0, + data, + elementSize, + static_cast(std::time(nullptr))); } // fill kernel wrapper for 'elementSize' elements with a specific value @@ -191,6 +209,19 @@ std::pair compareEqual(DDataType const* deviceD, return std::make_pair(retval, max_relative_error); } +template +double getEpsilon() +{ + if(std::is_same_v) + { + return 0.0009765625; // numeric_limits<_Float16>::epsilon() => 0 + } + else + { + return std::numeric_limits::epsilon(); + } +}; + template std::pair compareEqualLaunchKernel(DDataType* deviceD, DDataType* hostD, @@ -257,7 +288,7 @@ std::pair compareEqualLaunchKernel(DDataType* deviceD, auto toDouble = [](DDataType const& val) { return static_cast(static_cast(val)); }; - auto eps = toDouble(std::numeric_limits::epsilon()); + auto eps = getEpsilon(); if(isNaN) { retval = false; @@ -276,6 +307,7 @@ namespace std template ostream& operator<<(ostream& os, const std::vector& vec) { + os << "[ "; for(auto i = 0; i < vec.size(); i++) { if(i < vec.size() - 1) @@ -287,6 +319,7 @@ namespace std os << vec[i]; } } + os << " ]"; return os; }