ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-01 20:27:42 +00:00
Code Issues Packages Projects Releases Wiki Activity

Merge branch 'migx-jit-lib' into migraphx

Browse Source
This commit is contained in:
Umang Yadav
2023-09-28 14:06:00 +00:00
parent 11cab2d533 8f9c0243c7
commit 000c8bcf79
2 changed files with 7 additions and 5 deletions
Download Patch File Download Diff File Expand all files Collapse all files

4
include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_xdl_cshuffle.hpp
View File

@@ -577,9 +577,7 @@ struct DeviceGemmMultipleD_Xdl_CShuffle : public DeviceGemmMultipleD<ALayout,
static bool IsSupportedArgument(const Argument& arg)
{
if(!(ck::get_device_name() == "gfx908" || ck::get_device_name() == "gfx90a" ||
ck::get_device_name() == "gfx940" || ck::get_device_name() == "gfx941" ||
ck::get_device_name() == "gfx942"))
if(!ck::is_xdl_supported())
{
return false;
}

8
library/src/jit_library/src/device_gemm_multiple_d.cpp
View File

@@ -134,13 +134,15 @@ Solution Problem::MakeSolution(std::size_t idx, const std::string& arch) const
const std::size_t k_per_block = std::stoi(k_per_block_str);
const std::size_t grid_size = GetGridSize(M, N, m_per_block, n_per_block);
params[gemm_spec_idx] = GetGemmSpec(M, N, K, m_per_block, n_per_block, k_per_block);
std::string str = std::accumulate(
params.begin() + 1,
params.end(),
std::string{},
[](const std::string& a, const std::string& b) { return a.empty() ? b : a + ", " + b; });
str = params.front() + "< " + str + ">";
if (params.back().find("v2") != std::string::npos and K % k_per_block != 0)
str = "";
return Solution{str, block_size, grid_size};
}
@@ -156,7 +158,9 @@ std::vector<Solution> Problem::GetSolutions(const std::string& arch) const
const std::size_t num_instances = GetInstances(arch).size();
for(std::size_t i = 0; i < num_instances; ++i)
{
solutions.push_back(MakeSolution(i, arch));
auto solution = MakeSolution(i, arch);
if (solution.template_str != "")
solutions.push_back(solution);
}
return solutions;