/******************************************************************************* * * MIT License * * Copyright (c) 2019 Advanced Micro Devices, Inc. * * 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 GUARD_MIOPEN_TEST_POOLING_COMMON_HPP #define GUARD_MIOPEN_TEST_POOLING_COMMON_HPP #include "test.hpp" #include #include #include #include #include #include #include #include #include #include #include #include #include #include // #include "network_data.hpp" #include "driver.hpp" #include "get_handle.hpp" #include "tensor_holder.hpp" #include "verify.hpp" #include "cpu_conv.hpp" #include "workspace.hpp" #define TEST_PADDING_MODE 0 // NOLINTNEXTLINE (cppcoreguidelines-avoid-non-const-global-variables) static int num_uint16_case = 0; // NOLINTNEXTLINE (cppcoreguidelines-avoid-non-const-global-variables) static int num_uint32_case = 0; // NOLINTNEXTLINE (cppcoreguidelines-avoid-non-const-global-variables) static int num_uint32_case_imgidx = 0; // NOLINTNEXTLINE (cppcoreguidelines-avoid-non-const-global-variables) static int num_uint64_case = 0; // NOLINTNEXTLINE (cppcoreguidelines-avoid-non-const-global-variables) static int num_uint64_case_imgidx = 0; static inline void print(const miopen::PoolingDescriptor& filter) { std::cout << "Pooling: "; if(filter.GetMode() == miopenPoolingAverage) std::cout << "Average"; else if(filter.GetMode() == miopenPoolingAverageInclusive) std::cout << "AverageInclusive"; else std::cout << "Max"; std::cout << std::endl; std::cout << "Lengths: "; miopen::LogRange(std::cout, filter.GetLengths(), ", ") << std::endl; std::cout << "Pads: "; miopen::LogRange(std::cout, filter.GetPads(), ", ") << std::endl; std::cout << "Strides: "; miopen::LogRange(std::cout, filter.GetStrides(), ", ") << std::endl; } template tensor get_output_tensor(const miopen::PoolingDescriptor& filter, const tensor& input) { return tensor{filter.GetForwardOutputTensor(input.desc)}; } template struct pooling_operators { miopen::PoolingDescriptor filter; pooling_operators(miopen::PoolingDescriptor f) : filter(f) {} double start() const { if(filter.GetMode() == miopenPoolingMax) return std::numeric_limits::lowest(); else return 0.0; } double operator()(double x, double y) const { if(filter.GetMode() == miopenPoolingMax) { double m = std::max(x, y); return (m); } else { return x + y; } } double final(double x, double y) { if(filter.GetMode() == miopenPoolingMax) return (x); else return x / y; } }; template struct verify_forward_pooling { template tensor cpu(const tensor& input, const miopen::PoolingDescriptor& filter, std::vector&) const { auto out = get_output_tensor(filter, input); const auto in_strides = input.desc.GetStrides(); const std::size_t in_n_stride = in_strides[0]; const std::size_t in_c_stride = in_strides[1]; std::array in_spatial_strides{}; std::copy_n(in_strides.begin() + 2, SptDim, in_spatial_strides.begin()); const auto out_strides = out.desc.GetStrides(); const std::size_t out_n_stride = out_strides[0]; const std::size_t out_c_stride = out_strides[1]; std::array out_spatial_strides{}; std::copy_n(out_strides.begin() + 2, SptDim, out_spatial_strides.begin()); const T* input_ptr = input.data.data(); T* out_ptr = out.data.data(); std::array in_dim{}; std::copy_n(input.desc.GetLengths().begin() + 2, SptDim, in_dim.begin()); std::array strides{}; std::copy_n(filter.GetStrides().begin(), SptDim, strides.begin()); std::array pads{}; std::copy_n(filter.GetPads().begin(), SptDim, pads.begin()); std::array kers{}; std::copy_n(filter.GetLengths().begin(), SptDim, kers.begin()); auto op = pooling_operators{filter}; int b_n = out.desc.GetLengths()[0]; int k_n = out.desc.GetLengths()[1]; std::array out_spatial_len{}; std::copy_n(out.desc.GetLengths().begin() + 2, SptDim, out_spatial_len.begin()); auto par_ford_out = miopen::unpacker(miopen::prepender(miopen::par_ford, b_n, k_n))(out_spatial_len); par_ford_out([&](int o, int w, auto... out_spatial_id_pack) { auto out_spatial_id = make_array(out_spatial_id_pack...); std::array start_idx{}; std::array win_sz{}; for(int i = 0; i < SptDim; ++i) { start_idx[i] = out_spatial_id[i] * strides[i] - pads[i]; int end_idx = start_idx[i] + kers[i]; end_idx = std::min(end_idx, in_dim[i]); start_idx[i] = std::max(start_idx[i], 0); win_sz[i] = std::max(end_idx - start_idx[i], 0); } int pool_size = filter.GetMode() == miopenPoolingAverageInclusive ? std::accumulate(kers.begin(), kers.end(), 1, std::multiplies()) : std::accumulate(win_sz.begin(), win_sz.end(), 1, std::multiplies()); pool_size = std::max(pool_size, 1); // Avoid division by zero when window is in padding double acc = op.start(); miopen::unpacker(miopen::ford)(win_sz)([&](auto... in_spatial_id_pack) { auto in_spatial_id = make_array(in_spatial_id_pack...); std::array idx{}; idx[0] = o; idx[1] = w; bool in_cmp_idx = true; for(int i = 0; i < SptDim; ++i) { idx[i + 2] = start_idx[i] + in_spatial_id[i]; in_cmp_idx &= (in_dim[i] > idx[i + 2]); } if(in_cmp_idx) { // Compute input linear index std::size_t in_linear_idx = o * in_n_stride + w * in_c_stride; for(int i = 0; i < SptDim; ++i) { in_linear_idx += idx[i + 2] * in_spatial_strides[i]; } acc = op(acc, input_ptr[in_linear_idx]); } }); // Compute output linear index std::size_t out_linear_idx = o * out_n_stride + w * out_c_stride; for(int i = 0; i < SptDim; ++i) { out_linear_idx += out_spatial_id[i] * out_spatial_strides[i]; } out_ptr[out_linear_idx] = T(op.final(acc, pool_size)); }); return out; } template tensor gpu(const tensor& input, const miopen::PoolingDescriptor& filter, std::vector& indices) const { auto&& handle = get_handle(); auto out = get_output_tensor(filter, input); indices.resize(out.data.size(), 0); auto in_dev = handle.Write(input.data); auto out_dev = handle.Create(out.data.size()); Workspace wspace{}; wspace.Write(indices); float alpha = 1, beta = 0; filter.Forward(handle, &alpha, input.desc, in_dev.get(), &beta, out.desc, out_dev.get(), true, wspace.ptr(), wspace.size()); indices = wspace.Read>(); out.data = handle.Read(out_dev, out.data.size()); return out; } template void fail(float, const tensor& input, const miopen::PoolingDescriptor& filter, const std::vector&) const { std::cout << "Forward "; print(filter); std::cout << "Input tensor: " << input.desc.ToString() << std::endl; std::cout << "Output tensor: " << filter.GetForwardOutputTensor(input.desc).ToString() << std::endl; } }; template struct verify_backward_pooling { template tensor cpu(const tensor& input, const tensor& dout, const tensor& out, const miopen::PoolingDescriptor& filter, const std::vector& indices, bool use_global_index, bool verify_index) const { auto dinput = input; std::vector din_vec(input.desc.GetElementSpace(), 0.0); CHECK(dout.desc == out.desc); const auto in_strides = input.desc.GetStrides(); const std::size_t in_n_stride = in_strides[0]; const std::size_t in_c_stride = in_strides[1]; std::array in_spatial_strides{}; std::copy_n(in_strides.begin() + 2, SptDim, in_spatial_strides.begin()); const auto out_strides = dout.desc.GetStrides(); const std::size_t out_n_stride = out_strides[0]; const std::size_t out_c_stride = out_strides[1]; std::array out_spatial_strides{}; std::copy_n(out_strides.begin() + 2, SptDim, out_spatial_strides.begin()); const T* input_ptr = input.data.data(); const T* dout_ptr = dout.data.data(); const T* out_ptr = out.data.data(); std::array in_dim{}; std::copy_n(input.desc.GetLengths().begin(), SptDim + 2, in_dim.begin()); std::array in_str{}; std::copy_n(input.desc.GetStrides().begin(), SptDim + 2, in_str.begin()); std::array strides{}; std::copy_n(filter.GetStrides().begin(), SptDim, strides.begin()); std::array pads{}; std::copy_n(filter.GetPads().begin(), SptDim, pads.begin()); std::array kers{}; std::copy_n(filter.GetLengths().begin(), SptDim, kers.begin()); auto ford_ker = miopen::unpacker(miopen::ford)(kers); int out_n = out.desc.GetLengths()[0]; int out_c = out.desc.GetLengths()[1]; std::array out_spatial_len{}; std::copy_n(out.desc.GetLengths().begin() + 2, SptDim, out_spatial_len.begin()); auto ford_out = miopen::unpacker(miopen::ford)(out_spatial_len); miopen::par_ford(out_n, out_c)([&](int o, int w) { if(filter.GetMode() == miopenPoolingMax) { ford_out([&](auto... out_spatial_id_pack) { auto mx_idx = indices.at(dout.desc.GetIndex(o, w, out_spatial_id_pack...)); std::array idx{}; bool in_cmp_idx = true; if(use_global_index) { auto total_spatial = std::accumulate( in_dim.begin() + 2, in_dim.end(), 1ULL, std::multiplies()); in_cmp_idx = (mx_idx < total_spatial); auto out_spatial_id = make_array(out_spatial_id_pack...); for(int i = 0; i < SptDim && in_cmp_idx; i++) { int win_start = out_spatial_id[i] * strides[i] - pads[i]; int win_end = win_start + kers[i] - 1; if(win_end < 0 || win_start >= in_dim[i + 2]) { in_cmp_idx = false; } } for(int i = 0; i < SptDim; i++) { std::size_t mx_idx_dim = mx_idx; mx_idx_dim /= std::accumulate(in_dim.begin() + i + 3, in_dim.end(), 1ULL, std::multiplies()); mx_idx_dim %= in_dim[i + 2]; idx[i + 2] = mx_idx_dim; } } else { auto out_spatial_id = make_array(out_spatial_id_pack...); for(int i = 0; i < SptDim; i++) { int mx_idx_dim = mx_idx; mx_idx_dim /= std::accumulate( kers.begin() + i + 1, kers.end(), 1, std::multiplies()); mx_idx_dim %= kers[i]; mx_idx_dim += (out_spatial_id[i] * strides[i] - pads[i]); in_cmp_idx &= (in_dim[i + 2] > mx_idx_dim && mx_idx_dim >= 0); idx[i + 2] = std::size_t(mx_idx_dim); } } if(in_cmp_idx) { idx[0] = o; idx[1] = w; if(verify_index) { // Compute input and output linear indices for verification std::size_t in_verify_idx = o * in_n_stride + w * in_c_stride; auto out_spatial_id_verify = make_array(out_spatial_id_pack...); std::size_t out_verify_idx = o * out_n_stride + w * out_c_stride; for(int i = 0; i < SptDim; ++i) { in_verify_idx += idx[i + 2] * in_spatial_strides[i]; out_verify_idx += out_spatial_id_verify[i] * out_spatial_strides[i]; } CHECK(miopen::float_equal(input_ptr[in_verify_idx], out_ptr[out_verify_idx])); } std::size_t din_idx = 0; for(int i = 0; i < SptDim + 2; i++) { din_idx += idx[i] * in_str[i]; } // Compute dout linear index auto out_spatial_id = make_array(out_spatial_id_pack...); std::size_t dout_linear_idx = o * out_n_stride + w * out_c_stride; for(int i = 0; i < SptDim; ++i) { dout_linear_idx += out_spatial_id[i] * out_spatial_strides[i]; } din_vec.at(din_idx) += dout_ptr[dout_linear_idx]; } }); } else { ford_out([&](auto... out_spatial_id_pack) { auto out_spatial_id = make_array(out_spatial_id_pack...); std::array start_idx{}; std::array win_sz{}; for(int i = 0; i < SptDim; ++i) { start_idx[i] = out_spatial_id[i] * strides[i] - pads[i]; int end_idx = start_idx[i] + kers[i]; end_idx = std::min(end_idx, in_dim[i + 2]); win_sz[i] = end_idx - std::max(start_idx[i], 0); win_sz[i] = std::max(win_sz[i], 0); } int pool_size = filter.GetMode() == miopenPoolingAverageInclusive ? std::accumulate(kers.begin(), kers.end(), 1, std::multiplies()) : std::accumulate( win_sz.begin(), win_sz.end(), 1, std::multiplies()); pool_size = std::max(pool_size, 1); // Avoid division by zero ford_ker([&](auto... ker_id_pack) { auto ker_id = make_array(ker_id_pack...); bool in_cmp_idx = true; std::array in_idx{}; in_idx[0] = o; in_idx[1] = w; for(int i = 0; i < SptDim; ++i) { in_idx[i + 2] = start_idx[i] + ker_id[i]; in_cmp_idx &= (in_dim[i + 2] > in_idx[i + 2] && in_idx[i + 2] >= 0); } if(in_cmp_idx) { std::size_t din_idx = 0; for(int i = 0; i < SptDim + 2; i++) { din_idx += in_idx[i] * in_str[i]; } // Compute dout linear index std::size_t dout_linear_idx = o * out_n_stride + w * out_c_stride; for(int i = 0; i < SptDim; ++i) { dout_linear_idx += out_spatial_id[i] * out_spatial_strides[i]; } din_vec.at(din_idx) += static_cast(dout_ptr[dout_linear_idx]) / pool_size; } }); }); } }); miopen::unpacker(miopen::ford)(in_dim)([&](auto... in_id_pack) { auto in_id = make_array(in_id_pack...); std::size_t din_idx = 0; for(int i = 0; i < SptDim + 2; i++) { din_idx += in_id[i] * in_str[i]; } dinput(in_id_pack...) = din_vec.at(din_idx); }); return dinput; } template tensor gpu(const tensor& input, const tensor& dout, const tensor& out, const miopen::PoolingDescriptor& filter, const std::vector& indices, bool, bool) const { auto&& handle = get_handle(); auto dinput = input; auto in_dev = handle.Write(input.data); auto dout_dev = handle.Write(dout.data); auto out_dev = handle.Write(out.data); auto din_dev = handle.Create(dinput.data.size()); Workspace wspace{}; wspace.Write(indices); float alpha = 1, beta = 0; filter.Backward(handle, &alpha, // y out.desc, out_dev.get(), // dy dout.desc, dout_dev.get(), // x input.desc, in_dev.get(), &beta, // dx dinput.desc, din_dev.get(), wspace.ptr()); dinput.data = handle.Read(din_dev, dinput.data.size()); return dinput; } template void fail(float, const tensor& input, const tensor&, const tensor& out, const miopen::PoolingDescriptor& filter, const std::vector&, bool, bool) const { std::cout << "Backward "; print(filter); std::cout << "Input tensor: " << input.desc.ToString() << std::endl; std::cout << "Output tensor: " << out.desc.ToString() << std::endl; } }; template struct pooling_driver : test_driver { miopen::PoolingDescriptor filter; std::vector in_shape; std::vector lens; std::vector pads; std::vector strides; std::string index_type; std::string mode; #if TEST_PADDING_MODE == 1 std::string pmode; #endif int verify_indices{}; int wsidx{}; std::string in_layout; std::unordered_map index_type_lookup = { {miopen::ToUpper("miopenIndexUint8"), miopenIndexUint8}, {miopen::ToUpper("miopenIndexUint16"), miopenIndexUint16}, {miopen::ToUpper("miopenIndexUint32"), miopenIndexUint32}, {miopen::ToUpper("miopenIndexUint64"), miopenIndexUint64}, }; std::unordered_map mode_lookup = { {"MAX", miopenPoolingMax}, {"MIOPENPOOLINGMAX", miopenPoolingMax}, {"AVERAGE", miopenPoolingAverage}, {"MIOPENPOOLINGAVERAGE", miopenPoolingAverage}, {"AVERAGEINCLUSIVE", miopenPoolingAverageInclusive}, {"MIOPENPOOLINGAVERAGEINCLUSIVE", miopenPoolingAverageInclusive}, }; #if TEST_PADDING_MODE == 1 std::unordered_map pmode_lookup = { {"DEFAULT", miopenPaddingDefault}, {"SAME", miopenPaddingSame}, {"VALID", miopenPaddingValid}, }; #endif pooling_driver() { add(index_type, "index_type", generate_multi_data( // {{"miopenIndexUint8", "miopenIndexUint16", "miopenIndexUint32", "miopenIndexUint64"}, // {"miopenIndexUint8", "miopenIndexUint32"}, // {"miopenIndexUint32"}} // )); add(mode, "mode", generate_data( {"miopenPoolingMax", "miopenPoolingAverage", "miopenPoolingAverageInclusive"})); #if TEST_PADDING_MODE == 1 add(pmode, "pmode", generate_data({"default", "same", "valid"})); #endif add(verify_indices, "verify_indices", generate_data({1})); add(in_layout, "in_layout", generate_data({"NCHW"})); } template void run_impl() { std::vector indices{}; auto input = tensor{in_shape}; // Set up tensor descriptor with correct layout BEFORE generating data if(in_layout != "NCHW" && in_layout != "NCDHW") { const std::vector dim_lens = input.desc.GetLengths(); std::vector dim_strides; miopen::tensor_layout_to_strides( dim_lens, miopen::tensor_layout_get_default(input.desc.GetNumDims()), in_layout, dim_strides); input.desc = miopen::TensorDescriptor(miopen_type{}, dim_lens, dim_strides); } input.generate(tensor_elem_gen_integer{ (miopen_type{} == miopenHalf || miopen_type{} == miopenBFloat16) ? 5 : 17}); auto out = verify(verify_forward_pooling{}, input, filter, indices); auto dout = out.first; dout.generate(tensor_elem_gen_integer{2503}); verify(verify_backward_pooling{}, input, dout, out.first, filter, indices, wsidx != 0, static_cast(this->verify_indices)); } void run() { auto idx_typ = index_type_lookup.at(miopen::ToUpper(index_type)); auto idx_sz = sizeof(uint8_t); int spt_dim = in_shape.size() - 2; const bool skip_many_configs_with_non_int8_index = (dataset_id == 0) && full_set; // Otherwise the default dataset takes too much time. const bool wide_dataset = (dataset_id == 2) && full_set; filter = miopen::PoolingDescriptor{mode_lookup.at(miopen::ToUpper(mode)), #if TEST_PADDING_MODE == 1 pmode_lookup.at(miopen::ToUpper(pmode)), #else miopenPaddingDefault, #endif lens, strides, pads}; filter.SetIndexType(idx_typ); filter.SetWorkspaceIndexMode(miopenPoolingWorkspaceIndexMode_t(wsidx)); if(wsidx == 0 && spt_dim == 3 && filter.GetMode() == miopenPoolingMax && full_set) { show_command(); std::cout << "Warning: Config skipped. Workspace index mask mode is not implemented " "yet in 3D max pooling solvers." << std::endl; return; } if(wsidx == 0 && spt_dim == 2 && filter.GetMode() == miopenPoolingMax && wide_dataset) { show_command(); std::cout << "Warning: Config skipped. Workspace index mask mode is not implemented " "yet in 2D max backward solvers that support wide pooling window." << std::endl; return; } if(wsidx == 0 && (filter.GetMode() == miopenPoolingAverage || filter.GetMode() == miopenPoolingAverageInclusive) && full_set) { show_command(); std::cout << "Warning: Config skipped. Workspace index modes are irrelevant for " "Average pooling. " "In order to optimize performance of full tests, we " "skip average pooling configs when (wsidx == 0). " "Please make sure that dataset includes counterparts with (wsidx == 1)." << std::endl; return; } switch(idx_typ) { /// The "index is too small" limitation is an approximation /// of the real limitation, and therefore applied only when /// the "full test" is ran. See: /// \ref max_pooling_index_max_restriction case miopenIndexUint8: { if((spt_dim == 3 || (spt_dim == 2 && wsidx == 1)) && full_set && filter.GetMode() == miopenPoolingMax) { show_command(); std::cout << "Warning: Config skipped: uint8 index is too small " "(spt_dim == 3 || (spt_dim == 2 && wsidx == 1)) " "&& filter.GetMode() == miopenPoolingMax" << std::endl; return; } break; } case miopenIndexUint16: { if((spt_dim == 3 || (spt_dim == 2 && wsidx == 1)) && full_set && filter.GetMode() == miopenPoolingMax) { show_command(); std::cout << "Warning: Config skipped: uint16 index is too small " "(spt_dim == 3 || (spt_dim == 2 && wsidx == 1)) " "&& filter.GetMode() == miopenPoolingMax" << std::endl; return; } if(skip_many_configs_with_non_int8_index) { // test_pooling_test --all only test 5 uint16 cases if(num_uint16_case > 5) { show_command(); std::cout << "Warning: Config skipped for the default dataset to speed " "up testing (num_uint16_case > 5)" << std::endl; return; } ++num_uint16_case; } idx_sz = sizeof(uint16_t); break; } case miopenIndexUint32: { if(skip_many_configs_with_non_int8_index) { // test_pooling_test --all only test 5 uint32 cases if(wsidx == 0) { if(num_uint32_case > 5) { show_command(); std::cout << "Warning: Config skipped for the default dataset to speed up " "testing (wsidx == 0 && num_uint32_case > 5)" << std::endl; return; } ++num_uint32_case; } else { if(num_uint32_case_imgidx > 5) { show_command(); std::cout << "Warning: Config skipped for the default dataset to speed up " "testing (wsidx != 0 && num_uint32_case_imgidx > 5)" << std::endl; return; } ++num_uint32_case_imgidx; } } idx_sz = sizeof(uint32_t); break; } case miopenIndexUint64: { if(skip_many_configs_with_non_int8_index) { if(wsidx == 0) { if(num_uint64_case > 5) { show_command(); std::cout << "Warning: Config skipped for the default dataset to speed up " "testing (wsidx == 0) && (num_uint64_case > 5)" << std::endl; return; } ++num_uint64_case; } else { if(num_uint64_case_imgidx > 5 && spt_dim == 2) { show_command(); std::cout << "Warning: Config skipped to speed up testing of the " "default dataset (wsidx != 0) && (num_uint64_case_imgidx > 5 " "&& spt_dim == 2)" << std::endl; return; } ++num_uint64_case_imgidx; } } idx_sz = sizeof(uint64_t); break; } } auto input_desc = miopen::TensorDescriptor(this->type, in_shape); if(spt_dim != 2 && spt_dim != 3) { show_command(); std::cout << "Warning: Config skipped becuse it is not supported " // "(spt_dim != 2 && spt_dim != 3)" << std::endl; return; } for(int i = 0; i < spt_dim; i++) { if(lens[i] > (input_desc.GetLengths()[i + 2] + static_cast(2) * pads[i])) { show_command(); std::cout << "Warning: Config skipped becuse it is invalid " "(lens[i] > (input_desc.GetLengths()[i + 2] + 2 * pads[i]))" << std::endl; return; } } if(full_set) { auto output_desc = filter.GetForwardOutputTensor(input_desc); size_t total_mem = 3 * input_desc.GetNumBytes() + output_desc.GetNumBytes() + idx_sz * output_desc.GetElementSize(); // estimate based on backward pass size_t device_mem = get_handle().GetGlobalMemorySize(); if(total_mem >= device_mem) { show_command(); std::cout << "Config skipped because it requires " << total_mem << " Bytes to write all necessary tensors to GPU. GPU has " << device_mem << " Bytes of memory." << std::endl; return; } } std::vector in_dim(input_desc.GetLengths().begin() + 2, input_desc.GetLengths().end()); std::vector out_dim(spt_dim); std::vector ker_dim(filter.GetLengths().begin(), filter.GetLengths().end()); #if TEST_PADDING_MODE == 1 if(filter.pmode == miopenPaddingSame) { if(std::any_of(filter.GetStrides().begin(), filter.GetStrides().end(), [](int i) { return i == 0; })) return; for(int i = 0; i < spt_dim; i++) { filter.pads[i] = ((in_dim[i] % filter.GetStrides()[i] == 0) ? (std::max((ker_dim[i] - filter.GetStrides()[i]), 0)) : (std::max((ker_dim[i] - (in_dim[i] % filter.GetStrides()[i])), 0))) / 2; out_dim[i] = std::ceil(static_cast(in_dim[i]) / filter.strides[i]); } if(std::any_of(out_dim.begin(), out_dim.end(), [](int i) { return i <= 0; })) return; } else if(filter.pmode == miopenPaddingValid) { if(std::any_of(filter.GetStrides().begin(), filter.GetStrides().end(), [](int i) { return i == 0; })) return; for(int i = 0; i < spt_dim; i++) { filter.pads[i] = 0; out_dim[i] = std::ceil(static_cast(in_dim[i] - filter.lens[i] + 1) / filter.strides[i]); } if(std::any_of(out_dim.begin(), out_dim.end(), [](int i) { return i <= 0; })) return; } #endif switch(filter.GetIndexType()) { case miopenIndexUint8: { if(spt_dim == 3) { run_impl(); } else { run_impl(); } break; } case miopenIndexUint16: { if(spt_dim == 3) { run_impl(); } else { run_impl(); } break; } case miopenIndexUint32: { if(spt_dim == 3) { run_impl(); } else { run_impl(); } break; } case miopenIndexUint64: { if(spt_dim == 3) { run_impl(); } else { run_impl(); } break; } } } }; #endif