From e432fd993dc7e1d1ae9bd55e7e922289f75845d1 Mon Sep 17 00:00:00 2001 From: Harsh Chauhan Date: Tue, 9 Jun 2026 22:54:47 +0530 Subject: [PATCH] add gtest for conv gpu SAME_UPPER autopad --- .../TestCustomModelsFromONNXForAlpakaCuda.cxx | 40 ++++++++++++++++++ .../ConvWithAutopadSameUpper.onnx | Bin 0 -> 277 bytes .../ConvWithAutopadSameUpper.ref.hxx | 3 ++ 3 files changed, 43 insertions(+) create mode 100644 core/test/input_models/ConvWithAutopadSameUpper.onnx create mode 100644 core/test/input_models/references/ConvWithAutopadSameUpper.ref.hxx diff --git a/core/test/TestCustomModelsFromONNXForAlpakaCuda.cxx b/core/test/TestCustomModelsFromONNXForAlpakaCuda.cxx index 5ad9383..df49e49 100644 --- a/core/test/TestCustomModelsFromONNXForAlpakaCuda.cxx +++ b/core/test/TestCustomModelsFromONNXForAlpakaCuda.cxx @@ -155,6 +155,9 @@ #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" @@ -2328,6 +2331,43 @@ TEST_F(SofieAlpakaTest, ConvWithAutopadSameLower) } +// SAME_UPPER with odd total padding (x[1,1,4,4], k3, s2): the extra pad goes at the end, +// so the output differs from SAME_LOWER. Guards the SAME_UPPER pad split (issue #33). +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/core/test/input_models/ConvWithAutopadSameUpper.onnx b/core/test/input_models/ConvWithAutopadSameUpper.onnx new file mode 100644 index 0000000000000000000000000000000000000000..bf0a5c8334e04d217f58387b99b2b00ca8e54d2f GIT binary patch literal 277 zcmajZ!3u&v5C&jpS8y3BWT_~+MJKP_x@1L{Adn=4E@HKyRJ!OEdFxSn&n`{SF^J)3 zU_SU6;mD@0O?!3 z*f^;&SLe7&6Rmvi^Cvi-i%w rXf?8|$Mj!z1Nu>Ru`NKl$OL?2)3!et{XkC0_k{HiWQ1&Pqvf{Vi0U+( literal 0 HcmV?d00001 diff --git a/core/test/input_models/references/ConvWithAutopadSameUpper.ref.hxx b/core/test/input_models/references/ConvWithAutopadSameUpper.ref.hxx new file mode 100644 index 0000000..60844c7 --- /dev/null +++ b/core/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