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 0000000..bf0a5c8 Binary files /dev/null and b/core/test/input_models/ConvWithAutopadSameUpper.onnx differ 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