CodeQL / Analyze (actions) (push) Canceled after 0s
CodeQL / Analyze (python) (push) Canceled after 0s
pre-commit / pre-commit (push) Canceled after 0s
Multi-Arch CI / setup (push) Canceled after 0s
Publish build_manylinux_x86_64 images / publish_build_manylinux_x86_64 (push) Canceled after 0s
Setup (Multi-Arch) / setup (push) Canceled after 0s
Multi-Arch CI / Linux::${{ fromJSON(needs.setup.outputs.linux_build_config || '{}').build_variant_label || 'skip' }} (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Publish no_rocm_image_ubi10 image / publish_no_rocm_image_ubi10 (push) Canceled after 0s
Multi-Arch CI / Windows::${{ fromJSON(needs.setup.outputs.windows_build_config || '{}').build_variant_label || 'skip' }} (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Multi-Arch CI / Manifest Diff (push) Canceled after 0s
TheRock Manifest Diff Report / Generate Manifest Diff Report (push) Canceled after 0s
Multi-Arch CI / CI Summary (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04 (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04_ocl_rt images / publish_no_rocm_image_ubuntu24_04_ocl_rt (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_rocgdb (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04_rocgdb images / publish_no_rocm_image_ubuntu24_04_rocgdb (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_ocl_rt (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Unit Tests / Unit Tests :: ubuntu-24.04 (push) Canceled after 0s
Unit Tests / Unit Tests :: windows-2022 (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_media (push) Canceled after 0s
Unit Tests / Unit Tests Summary (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_openmpi (push) Canceled after 0s
673 lines
25 KiB
C++
673 lines
25 KiB
C++
// Copyright © Advanced Micro Devices, Inc., or its affiliates.
|
|
// SPDX-License-Identifier: MIT
|
|
|
|
#include "MiopenUtils.hpp"
|
|
#include "common/MiopenHandleFixture.hpp"
|
|
|
|
#include <gtest/gtest.h>
|
|
#include <hipdnn_flatbuffers_sdk/data_objects/graph_generated.h>
|
|
#include <hipdnn_flatbuffers_sdk/data_objects/tensor_attributes_generated.h>
|
|
#include <hipdnn_plugin_sdk/PluginApiDataTypes.h>
|
|
|
|
using namespace miopen_plugin;
|
|
using namespace miopen_utils;
|
|
|
|
TEST(TestMiopenUtils, TensorDataTypeToMiopenDataType)
|
|
{
|
|
using namespace hipdnn_flatbuffers_sdk::data_objects;
|
|
|
|
EXPECT_EQ(miopen_utils::tensorDataTypeToMiopenDataType(DataType::FLOAT), miopenFloat);
|
|
EXPECT_EQ(miopen_utils::tensorDataTypeToMiopenDataType(DataType::HALF), miopenHalf);
|
|
EXPECT_EQ(miopen_utils::tensorDataTypeToMiopenDataType(DataType::BFLOAT16), miopenBFloat16);
|
|
}
|
|
|
|
TEST(TestMiopenUtils, TensorDataTypeToMiopenDataTypeThrowsOnUnsupported)
|
|
{
|
|
// Use a value not in the enum
|
|
EXPECT_THROW(miopen_utils::tensorDataTypeToMiopenDataType(
|
|
static_cast<hipdnn_flatbuffers_sdk::data_objects::DataType>(-1)),
|
|
hipdnn_plugin_sdk::HipdnnPluginException);
|
|
}
|
|
|
|
TEST(TestMiopenUtils, FindTensorAttributesReturnsCorrectValue)
|
|
{
|
|
flatbuffers::FlatBufferBuilder builder1;
|
|
auto attrOffset1
|
|
= hipdnn_flatbuffers_sdk::data_objects::CreateTensorAttributesDirect(builder1, 1);
|
|
builder1.Finish(attrOffset1);
|
|
|
|
flatbuffers::FlatBufferBuilder builder2;
|
|
auto attrOffset2
|
|
= hipdnn_flatbuffers_sdk::data_objects::CreateTensorAttributesDirect(builder2, 2);
|
|
builder2.Finish(attrOffset2);
|
|
|
|
auto attrPtr1 = flatbuffers::GetRoot<hipdnn_flatbuffers_sdk::data_objects::TensorAttributes>(
|
|
builder1.GetBufferPointer());
|
|
auto attrPtr2 = flatbuffers::GetRoot<hipdnn_flatbuffers_sdk::data_objects::TensorAttributes>(
|
|
builder2.GetBufferPointer());
|
|
|
|
auto attrMap
|
|
= std::unordered_map<int64_t,
|
|
const hipdnn_flatbuffers_sdk::data_objects::TensorAttributes*>{
|
|
{1, attrPtr1}, {2, attrPtr2}};
|
|
|
|
EXPECT_EQ(miopen_utils::findTensorAttributes(attrMap, 1).uid(), 1);
|
|
EXPECT_EQ(miopen_utils::findTensorAttributes(attrMap, 2).uid(), 2);
|
|
}
|
|
|
|
TEST(TestMiopenUtils, FindTensorAttributesThrowsIfNotFound)
|
|
{
|
|
auto attrMap
|
|
= std::unordered_map<int64_t,
|
|
const hipdnn_flatbuffers_sdk::data_objects::TensorAttributes*>{};
|
|
|
|
EXPECT_THROW(miopen_utils::findTensorAttributes(attrMap, 1),
|
|
hipdnn_plugin_sdk::HipdnnPluginException);
|
|
}
|
|
|
|
TEST(TestMiopenUtils, GetSpatialDimCountReturnsCorrectValue)
|
|
{
|
|
const std::vector<int64_t> dims = {1, 1, 1, 1, 1};
|
|
|
|
flatbuffers::FlatBufferBuilder builder;
|
|
auto attrOffset = hipdnn_flatbuffers_sdk::data_objects::CreateTensorAttributesDirect(
|
|
builder, 1, "", hipdnn_flatbuffers_sdk::data_objects::DataType::UNSET, nullptr, &dims);
|
|
builder.Finish(attrOffset);
|
|
|
|
auto attrPtr1 = flatbuffers::GetRoot<hipdnn_flatbuffers_sdk::data_objects::TensorAttributes>(
|
|
builder.GetBufferPointer());
|
|
|
|
EXPECT_EQ(miopen_utils::getSpatialDimCount(*attrPtr1), 3);
|
|
}
|
|
|
|
TEST(TestMiopenUtils, GetSpatialDimCountThrowsOnInvalidDims)
|
|
{
|
|
const std::vector<int64_t> dims = {1, 1};
|
|
|
|
flatbuffers::FlatBufferBuilder builder;
|
|
auto attrOffset = hipdnn_flatbuffers_sdk::data_objects::CreateTensorAttributesDirect(
|
|
builder, 1, "", hipdnn_flatbuffers_sdk::data_objects::DataType::UNSET, nullptr, &dims);
|
|
builder.Finish(attrOffset);
|
|
|
|
auto attrPtr1 = flatbuffers::GetRoot<hipdnn_flatbuffers_sdk::data_objects::TensorAttributes>(
|
|
builder.GetBufferPointer());
|
|
|
|
EXPECT_THROW(miopen_utils::getSpatialDimCount(*attrPtr1),
|
|
hipdnn_plugin_sdk::HipdnnPluginException);
|
|
}
|
|
|
|
TEST(TestMiopenUtils, CreateBatchnormTensor4dPassthrough)
|
|
{
|
|
// 4D NCHW tensor should pass through unchanged
|
|
const std::vector<int64_t> dims = {2, 3, 14, 14};
|
|
const std::vector<int64_t> strides = {588, 196, 14, 1};
|
|
|
|
flatbuffers::FlatBufferBuilder builder;
|
|
auto attrOffset = hipdnn_flatbuffers_sdk::data_objects::CreateTensorAttributesDirect(
|
|
builder, 42, "", hipdnn_flatbuffers_sdk::data_objects::DataType::FLOAT, &strides, &dims);
|
|
builder.Finish(attrOffset);
|
|
|
|
auto attrPtr = flatbuffers::GetRoot<hipdnn_flatbuffers_sdk::data_objects::TensorAttributes>(
|
|
builder.GetBufferPointer());
|
|
|
|
auto tensorMap
|
|
= std::unordered_map<int64_t,
|
|
const hipdnn_flatbuffers_sdk::data_objects::TensorAttributes*>{
|
|
{42, attrPtr}};
|
|
|
|
auto result = miopen_utils::createBatchnormTensor(tensorMap, 42);
|
|
|
|
EXPECT_EQ(result.uid(), 42);
|
|
|
|
// Verify tensor is unchanged (4D with exact dims/strides)
|
|
int numDims = 0;
|
|
miopenGetTensorDescriptorSize(result.tensorDescriptor(), &numDims);
|
|
ASSERT_EQ(numDims, 4);
|
|
|
|
std::vector<int> resultDims(4);
|
|
std::vector<int> resultStrides(4);
|
|
miopenDataType_t dataType;
|
|
miopenGetTensorDescriptor(
|
|
result.tensorDescriptor(), &dataType, resultDims.data(), resultStrides.data());
|
|
|
|
EXPECT_EQ(resultDims[0], 2);
|
|
EXPECT_EQ(resultDims[1], 3);
|
|
EXPECT_EQ(resultDims[2], 14);
|
|
EXPECT_EQ(resultDims[3], 14);
|
|
|
|
EXPECT_EQ(resultStrides[0], 588);
|
|
EXPECT_EQ(resultStrides[1], 196);
|
|
EXPECT_EQ(resultStrides[2], 14);
|
|
EXPECT_EQ(resultStrides[3], 1);
|
|
}
|
|
|
|
TEST(TestMiopenUtils, CreateBatchnormTensor5dPassthrough)
|
|
{
|
|
// 5D NCDHW tensor should pass through unchanged
|
|
const std::vector<int64_t> dims = {2, 3, 4, 14, 14};
|
|
const std::vector<int64_t> strides = {2352, 784, 196, 14, 1};
|
|
|
|
flatbuffers::FlatBufferBuilder builder;
|
|
auto attrOffset = hipdnn_flatbuffers_sdk::data_objects::CreateTensorAttributesDirect(
|
|
builder, 42, "", hipdnn_flatbuffers_sdk::data_objects::DataType::FLOAT, &strides, &dims);
|
|
builder.Finish(attrOffset);
|
|
|
|
auto attrPtr = flatbuffers::GetRoot<hipdnn_flatbuffers_sdk::data_objects::TensorAttributes>(
|
|
builder.GetBufferPointer());
|
|
|
|
auto tensorMap
|
|
= std::unordered_map<int64_t,
|
|
const hipdnn_flatbuffers_sdk::data_objects::TensorAttributes*>{
|
|
{42, attrPtr}};
|
|
|
|
auto result = miopen_utils::createBatchnormTensor(tensorMap, 42);
|
|
|
|
EXPECT_EQ(result.uid(), 42);
|
|
|
|
// Verify tensor is unchanged (5D with exact dims/strides)
|
|
int numDims = 0;
|
|
miopenGetTensorDescriptorSize(result.tensorDescriptor(), &numDims);
|
|
ASSERT_EQ(numDims, 5);
|
|
|
|
std::vector<int> resultDims(5);
|
|
std::vector<int> resultStrides(5);
|
|
miopenDataType_t dataType;
|
|
miopenGetTensorDescriptor(
|
|
result.tensorDescriptor(), &dataType, resultDims.data(), resultStrides.data());
|
|
|
|
EXPECT_EQ(resultDims[0], 2);
|
|
EXPECT_EQ(resultDims[1], 3);
|
|
EXPECT_EQ(resultDims[2], 4);
|
|
EXPECT_EQ(resultDims[3], 14);
|
|
EXPECT_EQ(resultDims[4], 14);
|
|
|
|
EXPECT_EQ(resultStrides[0], 2352);
|
|
EXPECT_EQ(resultStrides[1], 784);
|
|
EXPECT_EQ(resultStrides[2], 196);
|
|
EXPECT_EQ(resultStrides[3], 14);
|
|
EXPECT_EQ(resultStrides[4], 1);
|
|
}
|
|
|
|
TEST(TestMiopenUtils, CreateBatchnormTensor3dNclPadsToNchw)
|
|
{
|
|
// NCL (channels-first): dims [N, C, L], strides [C*L, L, 1]
|
|
// C stride (14) > L stride (1) = channels-first
|
|
const std::vector<int64_t> dims = {1, 3, 14};
|
|
const std::vector<int64_t> strides = {42, 14, 1};
|
|
|
|
flatbuffers::FlatBufferBuilder builder;
|
|
auto attrOffset = hipdnn_flatbuffers_sdk::data_objects::CreateTensorAttributesDirect(
|
|
builder, 42, "", hipdnn_flatbuffers_sdk::data_objects::DataType::FLOAT, &strides, &dims);
|
|
builder.Finish(attrOffset);
|
|
|
|
auto attrPtr = flatbuffers::GetRoot<hipdnn_flatbuffers_sdk::data_objects::TensorAttributes>(
|
|
builder.GetBufferPointer());
|
|
|
|
auto tensorMap
|
|
= std::unordered_map<int64_t,
|
|
const hipdnn_flatbuffers_sdk::data_objects::TensorAttributes*>{
|
|
{42, attrPtr}};
|
|
|
|
auto result = miopen_utils::createBatchnormTensor(tensorMap, 42);
|
|
|
|
// Should be padded to 4D with W=1 and stride[3]=1 for NCHW
|
|
int numDims = 0;
|
|
miopenGetTensorDescriptorSize(result.tensorDescriptor(), &numDims);
|
|
ASSERT_EQ(numDims, 4);
|
|
|
|
std::vector<int> resultDims(4);
|
|
std::vector<int> resultStrides(4);
|
|
miopenDataType_t dataType;
|
|
miopenGetTensorDescriptor(
|
|
result.tensorDescriptor(), &dataType, resultDims.data(), resultStrides.data());
|
|
|
|
EXPECT_EQ(resultDims[0], 1);
|
|
EXPECT_EQ(resultDims[1], 3);
|
|
EXPECT_EQ(resultDims[2], 14);
|
|
EXPECT_EQ(resultDims[3], 1);
|
|
|
|
EXPECT_EQ(resultStrides[3], 1); // W stride = 1 for channels-first
|
|
}
|
|
|
|
TEST(TestMiopenUtils, CreateBatchnormTensor3dNlcPadsToNhwc)
|
|
{
|
|
// NLC (channels-last): dims [N, C, L], strides [C*L, 1, C]
|
|
// C stride (1) < L stride (3) = channels-last
|
|
const std::vector<int64_t> dims = {1, 3, 14};
|
|
const std::vector<int64_t> strides = {42, 1, 3};
|
|
|
|
flatbuffers::FlatBufferBuilder builder;
|
|
auto attrOffset = hipdnn_flatbuffers_sdk::data_objects::CreateTensorAttributesDirect(
|
|
builder, 42, "", hipdnn_flatbuffers_sdk::data_objects::DataType::FLOAT, &strides, &dims);
|
|
builder.Finish(attrOffset);
|
|
|
|
auto attrPtr = flatbuffers::GetRoot<hipdnn_flatbuffers_sdk::data_objects::TensorAttributes>(
|
|
builder.GetBufferPointer());
|
|
|
|
auto tensorMap
|
|
= std::unordered_map<int64_t,
|
|
const hipdnn_flatbuffers_sdk::data_objects::TensorAttributes*>{
|
|
{42, attrPtr}};
|
|
|
|
auto result = miopen_utils::createBatchnormTensor(tensorMap, 42);
|
|
|
|
// Should be padded to 4D with W=1 and stride[3]=C for NHWC
|
|
int numDims = 0;
|
|
miopenGetTensorDescriptorSize(result.tensorDescriptor(), &numDims);
|
|
ASSERT_EQ(numDims, 4);
|
|
|
|
std::vector<int> resultDims(4);
|
|
std::vector<int> resultStrides(4);
|
|
miopenDataType_t dataType;
|
|
miopenGetTensorDescriptor(
|
|
result.tensorDescriptor(), &dataType, resultDims.data(), resultStrides.data());
|
|
|
|
EXPECT_EQ(resultDims[0], 1);
|
|
EXPECT_EQ(resultDims[1], 3);
|
|
EXPECT_EQ(resultDims[2], 14);
|
|
EXPECT_EQ(resultDims[3], 1);
|
|
|
|
EXPECT_EQ(resultStrides[3], 3); // W stride = C (channel size) for channels-last
|
|
}
|
|
|
|
TEST(TestMiopenUtils, MapPointwiseModeStandardRelu)
|
|
{
|
|
flatbuffers::FlatBufferBuilder builder;
|
|
auto attrOffset = hipdnn_flatbuffers_sdk::data_objects::CreatePointwiseAttributes(
|
|
builder, hipdnn_flatbuffers_sdk::data_objects::PointwiseMode::RELU_FWD);
|
|
builder.Finish(attrOffset);
|
|
const auto* attr
|
|
= flatbuffers::GetRoot<hipdnn_flatbuffers_sdk::data_objects::PointwiseAttributes>(
|
|
builder.GetBufferPointer());
|
|
|
|
ActivationParams result;
|
|
ASSERT_NO_THROW(result = mapPointwiseModeToMiopenActivation(*attr));
|
|
EXPECT_EQ(result.mode, miopenActivationRELU);
|
|
EXPECT_DOUBLE_EQ(result.alpha, 0.0);
|
|
}
|
|
|
|
TEST(TestMiopenUtils, MapPointwiseModeReluNonZeroLowerClipOnly)
|
|
{
|
|
const float lowerClip = 0.1f;
|
|
flatbuffers::FlatBufferBuilder builder;
|
|
auto attrOffset = hipdnn_flatbuffers_sdk::data_objects::CreatePointwiseAttributes(
|
|
builder, hipdnn_flatbuffers_sdk::data_objects::PointwiseMode::RELU_FWD, lowerClip);
|
|
builder.Finish(attrOffset);
|
|
const auto* attr
|
|
= flatbuffers::GetRoot<hipdnn_flatbuffers_sdk::data_objects::PointwiseAttributes>(
|
|
builder.GetBufferPointer());
|
|
|
|
ASSERT_THROW(mapPointwiseModeToMiopenActivation(*attr),
|
|
hipdnn_plugin_sdk::HipdnnPluginException);
|
|
}
|
|
|
|
TEST(TestMiopenUtils, MapPointwiseModeReluZeroLowerClipOnly)
|
|
{
|
|
const float lowerClip = 0.0f;
|
|
flatbuffers::FlatBufferBuilder builder;
|
|
auto attrOffset = hipdnn_flatbuffers_sdk::data_objects::CreatePointwiseAttributes(
|
|
builder, hipdnn_flatbuffers_sdk::data_objects::PointwiseMode::RELU_FWD, lowerClip);
|
|
builder.Finish(attrOffset);
|
|
const auto* attr
|
|
= flatbuffers::GetRoot<hipdnn_flatbuffers_sdk::data_objects::PointwiseAttributes>(
|
|
builder.GetBufferPointer());
|
|
|
|
ActivationParams result;
|
|
ASSERT_NO_THROW(result = mapPointwiseModeToMiopenActivation(*attr));
|
|
|
|
EXPECT_DOUBLE_EQ(result.alpha, static_cast<double>(lowerClip));
|
|
}
|
|
|
|
TEST(TestMiopenUtils, MapPointwiseModeClippedRelu)
|
|
{
|
|
const float upperClip = 6.0f;
|
|
flatbuffers::FlatBufferBuilder builder;
|
|
auto attrOffset = hipdnn_flatbuffers_sdk::data_objects::CreatePointwiseAttributes(
|
|
builder,
|
|
hipdnn_flatbuffers_sdk::data_objects::PointwiseMode::RELU_BWD,
|
|
flatbuffers::nullopt,
|
|
upperClip);
|
|
builder.Finish(attrOffset);
|
|
const auto* attr
|
|
= flatbuffers::GetRoot<hipdnn_flatbuffers_sdk::data_objects::PointwiseAttributes>(
|
|
builder.GetBufferPointer());
|
|
|
|
ActivationParams result;
|
|
ASSERT_NO_THROW(result = miopen_utils::mapPointwiseModeToMiopenActivation(*attr));
|
|
EXPECT_EQ(result.mode, miopenActivationCLIPPEDRELU);
|
|
EXPECT_DOUBLE_EQ(result.alpha, static_cast<double>(upperClip));
|
|
}
|
|
|
|
TEST(TestMiopenUtils, MapPointwiseModeLeakyRelu)
|
|
{
|
|
const float slope = 0.01f;
|
|
flatbuffers::FlatBufferBuilder builder;
|
|
auto attrOffset = hipdnn_flatbuffers_sdk::data_objects::CreatePointwiseAttributes(
|
|
builder,
|
|
hipdnn_flatbuffers_sdk::data_objects::PointwiseMode::RELU_FWD,
|
|
flatbuffers::nullopt,
|
|
flatbuffers::nullopt,
|
|
slope);
|
|
builder.Finish(attrOffset);
|
|
const auto* attr
|
|
= flatbuffers::GetRoot<hipdnn_flatbuffers_sdk::data_objects::PointwiseAttributes>(
|
|
builder.GetBufferPointer());
|
|
|
|
ActivationParams result;
|
|
ASSERT_NO_THROW(result = miopen_utils::mapPointwiseModeToMiopenActivation(*attr));
|
|
EXPECT_EQ(result.mode, miopenActivationLEAKYRELU);
|
|
EXPECT_DOUBLE_EQ(result.alpha, static_cast<double>(slope));
|
|
}
|
|
|
|
TEST(TestMiopenUtils, MapPointwiseModeClamp)
|
|
{
|
|
const float lowerClip = -1.0f;
|
|
const float upperClip = 6.0f;
|
|
flatbuffers::FlatBufferBuilder builder;
|
|
auto attrOffset = hipdnn_flatbuffers_sdk::data_objects::CreatePointwiseAttributes(
|
|
builder,
|
|
hipdnn_flatbuffers_sdk::data_objects::PointwiseMode::RELU_BWD,
|
|
lowerClip,
|
|
upperClip);
|
|
builder.Finish(attrOffset);
|
|
const auto* attr
|
|
= flatbuffers::GetRoot<hipdnn_flatbuffers_sdk::data_objects::PointwiseAttributes>(
|
|
builder.GetBufferPointer());
|
|
|
|
ActivationParams result;
|
|
ASSERT_NO_THROW(result = miopen_utils::mapPointwiseModeToMiopenActivation(*attr));
|
|
EXPECT_EQ(result.mode, miopenActivationCLAMP);
|
|
EXPECT_DOUBLE_EQ(result.alpha, static_cast<double>(lowerClip));
|
|
EXPECT_DOUBLE_EQ(result.beta, static_cast<double>(upperClip));
|
|
}
|
|
|
|
TEST(TestMiopenUtils, MapPointwiseModeSigmoid)
|
|
{
|
|
flatbuffers::FlatBufferBuilder builder;
|
|
auto attrOffset = hipdnn_flatbuffers_sdk::data_objects::CreatePointwiseAttributes(
|
|
builder, hipdnn_flatbuffers_sdk::data_objects::PointwiseMode::SIGMOID_FWD);
|
|
builder.Finish(attrOffset);
|
|
const auto* attr
|
|
= flatbuffers::GetRoot<hipdnn_flatbuffers_sdk::data_objects::PointwiseAttributes>(
|
|
builder.GetBufferPointer());
|
|
|
|
ActivationParams result;
|
|
ASSERT_NO_THROW(result = miopen_utils::mapPointwiseModeToMiopenActivation(*attr));
|
|
EXPECT_EQ(result.mode, miopenActivationLOGISTIC);
|
|
}
|
|
|
|
TEST(TestMiopenUtils, MapPointwiseModeTanh)
|
|
{
|
|
flatbuffers::FlatBufferBuilder builder;
|
|
auto attrOffset = hipdnn_flatbuffers_sdk::data_objects::CreatePointwiseAttributes(
|
|
builder, hipdnn_flatbuffers_sdk::data_objects::PointwiseMode::TANH_BWD);
|
|
builder.Finish(attrOffset);
|
|
const auto* attr
|
|
= flatbuffers::GetRoot<hipdnn_flatbuffers_sdk::data_objects::PointwiseAttributes>(
|
|
builder.GetBufferPointer());
|
|
|
|
ActivationParams result;
|
|
ASSERT_NO_THROW(result = miopen_utils::mapPointwiseModeToMiopenActivation(*attr));
|
|
EXPECT_EQ(result.mode, miopenActivationTANH);
|
|
EXPECT_DOUBLE_EQ(result.alpha, 1.0);
|
|
EXPECT_DOUBLE_EQ(result.beta, 1.0);
|
|
}
|
|
|
|
TEST(TestMiopenUtils, MapPointwiseModeEluWithCustomAlpha)
|
|
{
|
|
const float eluAlpha = 2.0f;
|
|
flatbuffers::FlatBufferBuilder builder;
|
|
auto attrOffset = hipdnn_flatbuffers_sdk::data_objects::CreatePointwiseAttributes(
|
|
builder,
|
|
hipdnn_flatbuffers_sdk::data_objects::PointwiseMode::ELU_FWD,
|
|
flatbuffers::nullopt,
|
|
flatbuffers::nullopt,
|
|
flatbuffers::nullopt,
|
|
flatbuffers::nullopt,
|
|
0,
|
|
flatbuffers::nullopt,
|
|
flatbuffers::nullopt,
|
|
1,
|
|
flatbuffers::nullopt,
|
|
eluAlpha);
|
|
builder.Finish(attrOffset);
|
|
const auto* attr
|
|
= flatbuffers::GetRoot<hipdnn_flatbuffers_sdk::data_objects::PointwiseAttributes>(
|
|
builder.GetBufferPointer());
|
|
|
|
ActivationParams result;
|
|
ASSERT_NO_THROW(result = miopen_utils::mapPointwiseModeToMiopenActivation(*attr));
|
|
EXPECT_EQ(result.mode, miopenActivationELU);
|
|
EXPECT_DOUBLE_EQ(result.alpha, static_cast<double>(eluAlpha));
|
|
}
|
|
|
|
TEST(TestMiopenUtils, MapPointwiseModeEluWithDefaultAlpha)
|
|
{
|
|
flatbuffers::FlatBufferBuilder builder;
|
|
auto attrOffset = hipdnn_flatbuffers_sdk::data_objects::CreatePointwiseAttributes(
|
|
builder, hipdnn_flatbuffers_sdk::data_objects::PointwiseMode::ELU_BWD);
|
|
builder.Finish(attrOffset);
|
|
const auto* attr
|
|
= flatbuffers::GetRoot<hipdnn_flatbuffers_sdk::data_objects::PointwiseAttributes>(
|
|
builder.GetBufferPointer());
|
|
|
|
ActivationParams result;
|
|
ASSERT_NO_THROW(result = miopen_utils::mapPointwiseModeToMiopenActivation(*attr));
|
|
EXPECT_EQ(result.mode, miopenActivationELU);
|
|
EXPECT_DOUBLE_EQ(result.alpha, 1.0);
|
|
}
|
|
|
|
TEST(TestMiopenUtils, MapPointwiseModeSoftplusWithoutBeta)
|
|
{
|
|
flatbuffers::FlatBufferBuilder builder;
|
|
auto attrOffset = hipdnn_flatbuffers_sdk::data_objects::CreatePointwiseAttributes(
|
|
builder, hipdnn_flatbuffers_sdk::data_objects::PointwiseMode::SOFTPLUS_FWD);
|
|
builder.Finish(attrOffset);
|
|
const auto* attr
|
|
= flatbuffers::GetRoot<hipdnn_flatbuffers_sdk::data_objects::PointwiseAttributes>(
|
|
builder.GetBufferPointer());
|
|
|
|
ActivationParams result;
|
|
ASSERT_NO_THROW(result = miopen_utils::mapPointwiseModeToMiopenActivation(*attr));
|
|
EXPECT_EQ(result.mode, miopenActivationSOFTRELU);
|
|
}
|
|
|
|
TEST(TestMiopenUtils, MapPointwiseModeSoftplusWithBetaOne)
|
|
{
|
|
const float beta = 1.0f;
|
|
flatbuffers::FlatBufferBuilder builder;
|
|
auto attrOffset = hipdnn_flatbuffers_sdk::data_objects::CreatePointwiseAttributes(
|
|
builder,
|
|
hipdnn_flatbuffers_sdk::data_objects::PointwiseMode::SOFTPLUS_BWD,
|
|
flatbuffers::nullopt,
|
|
flatbuffers::nullopt,
|
|
flatbuffers::nullopt,
|
|
flatbuffers::nullopt,
|
|
0,
|
|
flatbuffers::nullopt,
|
|
flatbuffers::nullopt,
|
|
1,
|
|
flatbuffers::nullopt,
|
|
flatbuffers::nullopt,
|
|
beta);
|
|
builder.Finish(attrOffset);
|
|
const auto* attr
|
|
= flatbuffers::GetRoot<hipdnn_flatbuffers_sdk::data_objects::PointwiseAttributes>(
|
|
builder.GetBufferPointer());
|
|
|
|
ActivationParams result;
|
|
ASSERT_NO_THROW(result = miopen_utils::mapPointwiseModeToMiopenActivation(*attr));
|
|
EXPECT_EQ(result.mode, miopenActivationSOFTRELU);
|
|
}
|
|
|
|
TEST(TestMiopenUtils, MapPointwiseModeSoftplusWithInvalidBeta)
|
|
{
|
|
const float beta = 2.0f;
|
|
flatbuffers::FlatBufferBuilder builder;
|
|
auto attrOffset = hipdnn_flatbuffers_sdk::data_objects::CreatePointwiseAttributes(
|
|
builder,
|
|
hipdnn_flatbuffers_sdk::data_objects::PointwiseMode::SOFTPLUS_FWD,
|
|
flatbuffers::nullopt,
|
|
flatbuffers::nullopt,
|
|
flatbuffers::nullopt,
|
|
flatbuffers::nullopt,
|
|
0,
|
|
flatbuffers::nullopt,
|
|
flatbuffers::nullopt,
|
|
1,
|
|
flatbuffers::nullopt,
|
|
flatbuffers::nullopt,
|
|
beta);
|
|
builder.Finish(attrOffset);
|
|
const auto* attr
|
|
= flatbuffers::GetRoot<hipdnn_flatbuffers_sdk::data_objects::PointwiseAttributes>(
|
|
builder.GetBufferPointer());
|
|
|
|
ASSERT_THROW(miopen_utils::mapPointwiseModeToMiopenActivation(*attr),
|
|
hipdnn_plugin_sdk::HipdnnPluginException);
|
|
}
|
|
|
|
TEST(TestMiopenUtils, MapPointwiseModeAbs)
|
|
{
|
|
flatbuffers::FlatBufferBuilder builder;
|
|
auto attrOffset = hipdnn_flatbuffers_sdk::data_objects::CreatePointwiseAttributes(
|
|
builder, hipdnn_flatbuffers_sdk::data_objects::PointwiseMode::ABS);
|
|
builder.Finish(attrOffset);
|
|
const auto* attr
|
|
= flatbuffers::GetRoot<hipdnn_flatbuffers_sdk::data_objects::PointwiseAttributes>(
|
|
builder.GetBufferPointer());
|
|
|
|
ActivationParams result;
|
|
ASSERT_NO_THROW(result = miopen_utils::mapPointwiseModeToMiopenActivation(*attr));
|
|
EXPECT_EQ(result.mode, miopenActivationABS);
|
|
}
|
|
|
|
TEST(TestMiopenUtils, MapPointwiseModeIdentity)
|
|
{
|
|
flatbuffers::FlatBufferBuilder builder;
|
|
auto attrOffset = hipdnn_flatbuffers_sdk::data_objects::CreatePointwiseAttributes(
|
|
builder, hipdnn_flatbuffers_sdk::data_objects::PointwiseMode::IDENTITY);
|
|
builder.Finish(attrOffset);
|
|
const auto* attr
|
|
= flatbuffers::GetRoot<hipdnn_flatbuffers_sdk::data_objects::PointwiseAttributes>(
|
|
builder.GetBufferPointer());
|
|
|
|
ActivationParams result;
|
|
ASSERT_NO_THROW(result = miopen_utils::mapPointwiseModeToMiopenActivation(*attr));
|
|
EXPECT_EQ(result.mode, miopenActivationPASTHRU);
|
|
}
|
|
|
|
TEST(TestMiopenUtils, MapPointwiseModeUnsupported)
|
|
{
|
|
flatbuffers::FlatBufferBuilder builder;
|
|
auto attrOffset = hipdnn_flatbuffers_sdk::data_objects::CreatePointwiseAttributes(
|
|
builder, hipdnn_flatbuffers_sdk::data_objects::PointwiseMode::ADD);
|
|
builder.Finish(attrOffset);
|
|
const auto* attr
|
|
= flatbuffers::GetRoot<hipdnn_flatbuffers_sdk::data_objects::PointwiseAttributes>(
|
|
builder.GetBufferPointer());
|
|
|
|
ASSERT_THROW(miopen_utils::mapPointwiseModeToMiopenActivation(*attr),
|
|
hipdnn_plugin_sdk::HipdnnPluginException);
|
|
}
|
|
|
|
// =============================================================================
|
|
// GPU-required tests for ScopedTuningPolicy
|
|
// =============================================================================
|
|
|
|
class TestGpuScopedTuningPolicy : public test_common::MiopenHandleFixture
|
|
{
|
|
};
|
|
|
|
TEST_F(TestGpuScopedTuningPolicy, SetsSearchPolicyWhenBenchmarkingEnabled)
|
|
{
|
|
{
|
|
const ScopedTuningPolicy guard(_miopenHandle, true);
|
|
|
|
miopenTuningPolicy_t currentPolicy;
|
|
auto status = miopenGetTuningPolicy(_miopenHandle, ¤tPolicy);
|
|
ASSERT_EQ(status, miopenStatusSuccess);
|
|
EXPECT_EQ(currentPolicy, miopenTuningPolicySearch);
|
|
}
|
|
}
|
|
|
|
TEST_F(TestGpuScopedTuningPolicy, SetsNonePolicyWhenBenchmarkingDisabled)
|
|
{
|
|
{
|
|
const ScopedTuningPolicy guard(_miopenHandle, false);
|
|
|
|
miopenTuningPolicy_t currentPolicy;
|
|
auto status = miopenGetTuningPolicy(_miopenHandle, ¤tPolicy);
|
|
ASSERT_EQ(status, miopenStatusSuccess);
|
|
EXPECT_EQ(currentPolicy, miopenTuningPolicyNone);
|
|
}
|
|
}
|
|
|
|
TEST_F(TestGpuScopedTuningPolicy, RestoresOriginalPolicyOnDestruction)
|
|
{
|
|
// Set a non-default policy first
|
|
auto preSetStatus = miopenSetTuningPolicy(_miopenHandle, miopenTuningPolicyDbUpdate);
|
|
ASSERT_EQ(preSetStatus, miopenStatusSuccess);
|
|
|
|
{
|
|
const ScopedTuningPolicy guard(_miopenHandle, true);
|
|
|
|
// Verify it was changed during scope
|
|
miopenTuningPolicy_t duringPolicy;
|
|
miopenGetTuningPolicy(_miopenHandle, &duringPolicy);
|
|
EXPECT_EQ(duringPolicy, miopenTuningPolicySearch);
|
|
}
|
|
|
|
// Verify it was restored to original
|
|
miopenTuningPolicy_t afterPolicy;
|
|
auto status = miopenGetTuningPolicy(_miopenHandle, &afterPolicy);
|
|
ASSERT_EQ(status, miopenStatusSuccess);
|
|
EXPECT_EQ(afterPolicy, miopenTuningPolicyDbUpdate);
|
|
}
|
|
|
|
TEST_F(TestGpuScopedTuningPolicy, RestoresToNoneWhenOriginalWasNone)
|
|
{
|
|
// Ensure policy starts as None
|
|
miopenSetTuningPolicy(_miopenHandle, miopenTuningPolicyNone);
|
|
|
|
{
|
|
const ScopedTuningPolicy guard(_miopenHandle, true);
|
|
}
|
|
|
|
miopenTuningPolicy_t afterPolicy;
|
|
auto status = miopenGetTuningPolicy(_miopenHandle, &afterPolicy);
|
|
ASSERT_EQ(status, miopenStatusSuccess);
|
|
EXPECT_EQ(afterPolicy, miopenTuningPolicyNone);
|
|
}
|
|
|
|
TEST_F(TestGpuScopedTuningPolicy, NestedScopesRestoreCorrectly)
|
|
{
|
|
miopenSetTuningPolicy(_miopenHandle, miopenTuningPolicyNone);
|
|
|
|
{
|
|
const ScopedTuningPolicy outerGuard(_miopenHandle, true);
|
|
|
|
miopenTuningPolicy_t afterOuterSet;
|
|
miopenGetTuningPolicy(_miopenHandle, &afterOuterSet);
|
|
EXPECT_EQ(afterOuterSet, miopenTuningPolicySearch);
|
|
|
|
{
|
|
const ScopedTuningPolicy innerGuard(_miopenHandle, false);
|
|
|
|
miopenTuningPolicy_t afterInnerSet;
|
|
miopenGetTuningPolicy(_miopenHandle, &afterInnerSet);
|
|
EXPECT_EQ(afterInnerSet, miopenTuningPolicyNone);
|
|
}
|
|
|
|
// After inner scope, should restore to Search (what outer set)
|
|
miopenTuningPolicy_t afterInnerDestroy;
|
|
miopenGetTuningPolicy(_miopenHandle, &afterInnerDestroy);
|
|
EXPECT_EQ(afterInnerDestroy, miopenTuningPolicySearch);
|
|
}
|
|
|
|
// After outer scope, should restore to None (original)
|
|
miopenTuningPolicy_t finalPolicy;
|
|
miopenGetTuningPolicy(_miopenHandle, &finalPolicy);
|
|
EXPECT_EQ(finalPolicy, miopenTuningPolicyNone);
|
|
}
|