diff --git a/test/alpaka/TestAlpakaConv.cxx b/test/alpaka/TestAlpakaConv.cxx index 3d11f94b..d0577f5a 100644 --- a/test/alpaka/TestAlpakaConv.cxx +++ b/test/alpaka/TestAlpakaConv.cxx @@ -6,6 +6,8 @@ #include "input_models/references/ConvWithoutPadding.ref.hxx" #include "ConvWithAutopadSameLower_FromONNX_GPU_ALPAKA.hxx" #include "input_models/references/ConvWithAutopadSameLower.ref.hxx" +#include "ConvWithAutopadSameUpper_FromONNX_GPU_ALPAKA.hxx" +#include "input_models/references/ConvWithAutopadSameUpper.ref.hxx" #include "ConvWithStridesPadding_FromONNX_GPU_ALPAKA.hxx" #include "input_models/references/ConvWithStridesPadding.ref.hxx" #include "ConvWithStridesNoPadding_FromONNX_GPU_ALPAKA.hxx" @@ -150,6 +152,41 @@ TEST_F(SofieAlpakaTest, ConvWithAutopadSameLower) } +TEST_F(SofieAlpakaTest, ConvWithAutopadSameUpper) +{ + constexpr float TOLERANCE = DEFAULT_TOLERANCE; + + std::vector input(16); + std::iota(input.begin(), input.end(), 0.0f); + auto input_h = alpaka::allocBuf(host, Ext1D::all(Idx{input.size()})); + float* input_ptr = reinterpret_cast(alpaka::getPtrNative(input_h)); + for (Idx i = 0; i < input.size(); ++i) input_ptr[i] = input[i]; + + auto input_d = alpaka::allocBuf(device, Ext1D::all(Idx{input.size()})); + alpaka::memcpy(queue, input_d, input_h); + alpaka::wait(queue); + + auto result_h = alpaka::allocBuf(host, Ext1D::all(Idx{sizeof(ConvWithAutopadSameUpper_ExpectedOutput::all_ones) / sizeof(float)})); + + { + SOFIE_ConvWithAutopadSameUpper::Session session("ConvWithAutopadSameUpper_FromONNX_GPU_ALPAKA.dat"); + auto result = session.infer(input_d); + alpaka::wait(queue); + cudaDeviceSynchronize(); + alpaka::memcpy(queue, result_h, result); + alpaka::wait(queue); + } + + float* res_ptr = reinterpret_cast(alpaka::getPtrNative(result_h)); + float *correct = ConvWithAutopadSameUpper_ExpectedOutput::all_ones; + constexpr size_t nOut_sameUpper = sizeof(ConvWithAutopadSameUpper_ExpectedOutput::all_ones) / sizeof(float); + + for (size_t i = 0; i < nOut_sameUpper; ++i) { + EXPECT_LE(std::abs(res_ptr[i] - correct[i]), TOLERANCE) << "i=" << i; + } +} + + TEST_F(SofieAlpakaTest, ConvWithStridesPadding) { constexpr float TOLERANCE = DEFAULT_TOLERANCE; diff --git a/test/input_models/ConvWithAutopadSameUpper.onnx b/test/input_models/ConvWithAutopadSameUpper.onnx new file mode 100644 index 00000000..bf0a5c83 Binary files /dev/null and b/test/input_models/ConvWithAutopadSameUpper.onnx differ diff --git a/test/input_models/references/ConvWithAutopadSameUpper.ref.hxx b/test/input_models/references/ConvWithAutopadSameUpper.ref.hxx new file mode 100644 index 00000000..60844c7a --- /dev/null +++ b/test/input_models/references/ConvWithAutopadSameUpper.ref.hxx @@ -0,0 +1,3 @@ +namespace ConvWithAutopadSameUpper_ExpectedOutput { +float all_ones[] = {45.0, 39.0, 66.0, 50.0}; +} // namespace ConvWithAutopadSameUpper_ExpectedOutput