diff --git a/src/targets/gpu/gemm_impl.cpp b/src/targets/gpu/gemm_impl.cpp index 88da849bdc0..bb0c60fa03b 100644 --- a/src/targets/gpu/gemm_impl.cpp +++ b/src/targets/gpu/gemm_impl.cpp @@ -129,7 +129,21 @@ auto rocblas_invoke(F f, Pack p, Ts... xs) }); } -static bool is_transposed(const shape& s) { return s.transposed() and s.strides().back() != 1; } +static bool is_transposed(const shape& s) +{ + if(s.transposed()) + { + return s.strides().back() != 1; + } + + if(not s.broadcasted() and s.strides() != s.as_standard().strides()) + { + auto perm = find_permutation(s); + return not std::is_sorted(perm.begin(), perm.end()); + } + + return false; +} static rocblas_int get_batch_stride(const shape& s) { diff --git a/test/verify/test_gemm_transposeb_detect.cpp b/test/verify/test_gemm_transposeb_detect.cpp new file mode 100644 index 00000000000..1a9098f9b2c --- /dev/null +++ b/test/verify/test_gemm_transposeb_detect.cpp @@ -0,0 +1,49 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * 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. + */ + +#include "verify_program.hpp" +#include +#include +#include + +template +struct test_gemm_transposeb_detect : verify_program> +{ + migraphx::program create_program() const + { + migraphx::program p; + auto* mm = p.get_main_module(); + auto a = mm->add_parameter("a", migraphx::shape{DType, {2, 2, 1}}); + auto b = mm->add_parameter("b", migraphx::shape{DType, {2, 2, 1}}); + auto bt = + mm->add_instruction(migraphx::make_op("transpose", {{"permutation", {0, 2, 1}}}), b); + mm->add_instruction(migraphx::make_op("dot"), a, bt); + return p; + } + std::string section() const { return "gemm"; } +}; + +template struct test_gemm_transposeb_detect; +template struct test_gemm_transposeb_detect; +template struct test_gemm_transposeb_detect;