Skip to content

fix(ck): build gfx1250-only conv instances only for gfx1250 #12883

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Closed
afagaj wants to merge 13 commits into from
+1,085 −368
Closed

fix(ck): build gfx1250-only conv instances only for gfx1250 #12883

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
13 commits
Select commit Hold shift + click to select a range
b0ba97d
fix(ck): build gfx1250-only conv instances only for gfx1250 (AICK-2468)
afagaj Oct 1, 2026
f22a696
fix(ck): address review — cover XDL large_k, drop bad instance restor…
afagaj Oct 1, 2026
a1e361d
fix(ck): subtract unwanted offload archs so per-source filtering actu…
afagaj Oct 2, 2026
01ccac1
fix(ck): gate gfx1250 instance registrations on an exported macro
afagaj Oct 2, 2026
3ab1d0d
ci(ck): drop gfx1250 from the build-instances-only stage
afagaj Oct 2, 2026
13e3ad4
fix(ck): restrict gfx1250 arch handling to _gfx1250 sources
afagaj Oct 2, 2026
5271c5d
refactor(ck): drop redundant gfx1250 guards in conv instance CMakeLists
afagaj Oct 2, 2026
5a197c4
refactor(ck): check for gfx1250 inside the gfx1250 add_ functions
afagaj Oct 2, 2026
e89e0b8
fix(ck): gate gemm v3 prefetch registrations on CK_HAS_GFX1250_INSTANCES
afagaj Oct 2, 2026
ee3c94f
docs(ck): correct the KPerBlock=48 comment; document the gfx1250 macr…
afagaj Oct 2, 2026
0fe5c6d
fix(miopen): keep CK objects the per-arch unbundler cannot thin
afagaj Oct 2, 2026
408ed62
style(ck): clang-format-18 the files this PR changes
afagaj Oct 2, 2026
7d4e581
feat(ck): fail safe when a WMMA conv tile exceeds the device's LDS
afagaj Oct 2, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
10 changes: 10 additions & 0 deletions projects/composablekernel/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -510,6 +510,16 @@ if (SUPPORTED_GPU_TARGETS MATCHES "gfx950")
endif()
if (SUPPORTED_GPU_TARGETS MATCHES "gfx1250")
add_definitions(-DCK_USE_GFX1250)
# Distinct from CK_USE_GFX1250, and exported via config.h so that consumers compiling
# against the installed headers see it. It means only "this CK build contains gfx1250
# instance translation units" -- which is what the instance-factory guards need to know.
# CK_USE_GFX1250 itself cannot be exported: ck/ck.hpp:205 keys the host half of
# CK_ENABLE_CLUSTER_LAUNCH on it and the device half on __gfx125__, so exporting it
# desynchronises the two on every non-gfx1250 offload pass. CK_GFX1250_SUPPORT is not
# reused either: dispatcher/bindings/ctypes/CMakeLists.txt undefines and redefines it per
# target architecture, so it does not say what this library build contains.
add_definitions(-DCK_HAS_GFX1250_INSTANCES)
set(CK_HAS_GFX1250_INSTANCES "ON")
add_definitions(-DCK_USE_NATIVE_MX_SUPPORT)
set(CK_USE_NATIVE_MX_SUPPORT "ON")
add_definitions(-DCK_GFX1250_SUPPORT)
Expand Down
9 changes: 9 additions & 0 deletions projects/composablekernel/include/ck/config.h.in
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -139,6 +139,15 @@
#cmakedefine CK_USE_NATIVE_MX_SUPPORT @CK_USE_NATIVE_MX_SUPPORT@
#endif

// Set when this CK build contains gfx1250-exclusive instance translation units. The
// instance-factory headers are installed, so the guards around their gfx1250 registrations
// must be visible to downstream consumers; CK_USE_GFX1250 is build-local only and would
// silently evaluate false for them. Deliberately not CK_USE_GFX1250 or CK_GFX1250_SUPPORT --
// see the comment where it is set in CMakeLists.txt.
#ifndef CK_HAS_GFX1250_INSTANCES
#cmakedefine CK_HAS_GFX1250_INSTANCES @CK_HAS_GFX1250_INSTANCES@
#endif

// clang-format on

#endif // CK_CONFIG_H_IN
2 changes: 1 addition & 1 deletion ...l/include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_wmma_cshuffle_v3.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -12,11 +12,11 @@
#include "ck/tensor_operation/gpu/device/tensor_layout.hpp"
#include "ck/tensor_operation/gpu/device/device_gemm_multiple_d.hpp"
#include "ck/tensor_operation/gpu/device/gemm_specialization.hpp"
#include "ck/tensor_operation/gpu/device/impl/device_gemm_wmma_cshuffle_v3_common.hpp"
#include "ck/tensor_operation/gpu/grid/gridwise_gemm_wmma_cshuffle_v3.hpp"
#include "ck/host_utility/device_prop.hpp"
#include "ck/host_utility/kernel_launch.hpp"
#include "ck/host_utility/flush_cache.hpp"
#include "ck/tensor_operation/gpu/device/impl/device_gemm_wmma_cshuffle_v3_common.hpp"

#if __clang_major__ >= 23
#pragma clang diagnostic push
Expand Down
162 changes: 84 additions & 78 deletions ...or_operation/gpu/device/impl/device_grouped_conv_bwd_data_multiple_d_wmma_cshuffle_v3.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -78,63 +78,35 @@ __launch_bounds__(GridwiseGemm::MaxBlockSize, MinimumOccupancy)

constexpr index_t LDS_size =
GridwiseGemm::template GetSharedMemoryNumberOfByte<SelectedEpilogue>();
__shared__ char p_shared[LDS_size];
auto epilogue_args = SelectedEpilogue{};

const index_t block_args_id = __builtin_amdgcn_readfirstlane(blockIdx.x);
index_t left = 0;
index_t right = gemms_count;
index_t group_id = index_t((left + right) / 2);
while((!(block_args_id >= gemm_kernel_args[group_id].BlockStart_ &&
block_args_id < gemm_kernel_args[group_id].BlockEnd_)) &&
left <= right)

if constexpr(LDS_size <= get_lds_size(get_device_arch()))
{
if(block_args_id < gemm_kernel_args[group_id].BlockStart_)
{
right = group_id;
}
else
__shared__ char p_shared[LDS_size];
auto epilogue_args = SelectedEpilogue{};

const index_t block_args_id = __builtin_amdgcn_readfirstlane(blockIdx.x);
index_t left = 0;
index_t right = gemms_count;
index_t group_id = index_t((left + right) / 2);
while((!(block_args_id >= gemm_kernel_args[group_id].BlockStart_ &&
block_args_id < gemm_kernel_args[group_id].BlockEnd_)) &&
left <= right)
{
left = group_id;
if(block_args_id < gemm_kernel_args[group_id].BlockStart_)
{
right = group_id;
}
else
{
left = group_id;
}
group_id = index_t((left + right) / 2);
}
group_id = index_t((left + right) / 2);
}

const auto num_k_per_block =
gemm_kernel_args[group_id].a_grid_desc_ak0_m_ak1_.GetLength(Number<0>{}) / KBatch;

if constexpr(HasMainKBlockLoopInAllGemm || NoMainKBlockLoopInAllGemm)
{
const auto num_k_per_block =
gemm_kernel_args[group_id].a_grid_desc_ak0_m_ak1_.GetLength(Number<0>{}) / KBatch;

GridwiseGemm::template Run<GridwiseGemm::ConvRegime::BWD_DATA,
AGridDesc_AK0_M_AK1,
BGridDesc_BK0_N_BK1,
DsGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock,
EGridDesc_MBlock_MPerBlock_NBlock_NPerBlock,
decltype(gemm_kernel_args[group_id].block_2_ctile_map_),
ComputePtrOffsetOfBatch,
ComputePtrOffsetOfN,
0,
HasMainKBlockLoopInAllGemm,
EGlobalMemoryDataOperation,
CTranspose,
TailNum,
decltype(epilogue_args)>(
p_shared,
gemm_kernel_args[group_id].a_grid_desc_ak0_m_ak1_,
gemm_kernel_args[group_id].b_grid_desc_bk0_n_bk1_,
gemm_kernel_args[group_id].ds_grid_desc_mblock_mperblock_nblock_nperblock_,
gemm_kernel_args[group_id].e_grid_desc_mblock_mperblock_nblock_nperblock_,
gemm_kernel_args[group_id].block_2_ctile_map_,
compute_ptr_offset_of_batch,
compute_ptr_offset_of_n,
num_k_per_block,
karg,
epilogue_args);
}
else
{
if(gemm_kernel_args[group_id].HasMainKBlockLoop_)
if constexpr(HasMainKBlockLoopInAllGemm || NoMainKBlockLoopInAllGemm)
{

GridwiseGemm::template Run<GridwiseGemm::ConvRegime::BWD_DATA,
Expand All @@ -146,7 +118,7 @@ __launch_bounds__(GridwiseGemm::MaxBlockSize, MinimumOccupancy)
ComputePtrOffsetOfBatch,
ComputePtrOffsetOfN,
0,
true,
HasMainKBlockLoopInAllGemm,
EGlobalMemoryDataOperation,
CTranspose,
TailNum,
Expand All @@ -165,32 +137,66 @@ __launch_bounds__(GridwiseGemm::MaxBlockSize, MinimumOccupancy)
}
else
{
if(gemm_kernel_args[group_id].HasMainKBlockLoop_)
{

GridwiseGemm::template Run<GridwiseGemm::ConvRegime::BWD_DATA,
AGridDesc_AK0_M_AK1,
BGridDesc_BK0_N_BK1,
DsGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock,
EGridDesc_MBlock_MPerBlock_NBlock_NPerBlock,
decltype(gemm_kernel_args[group_id].block_2_ctile_map_),
ComputePtrOffsetOfBatch,
ComputePtrOffsetOfN,
0,
false,
EGlobalMemoryDataOperation,
CTranspose,
TailNum,
decltype(epilogue_args)>(
p_shared,
gemm_kernel_args[group_id].a_grid_desc_ak0_m_ak1_,
gemm_kernel_args[group_id].b_grid_desc_bk0_n_bk1_,
gemm_kernel_args[group_id].ds_grid_desc_mblock_mperblock_nblock_nperblock_,
gemm_kernel_args[group_id].e_grid_desc_mblock_mperblock_nblock_nperblock_,
gemm_kernel_args[group_id].block_2_ctile_map_,
compute_ptr_offset_of_batch,
compute_ptr_offset_of_n,
num_k_per_block,
karg,
epilogue_args);
GridwiseGemm::template Run<
GridwiseGemm::ConvRegime::BWD_DATA,
AGridDesc_AK0_M_AK1,
BGridDesc_BK0_N_BK1,
DsGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock,
EGridDesc_MBlock_MPerBlock_NBlock_NPerBlock,
decltype(gemm_kernel_args[group_id].block_2_ctile_map_),
ComputePtrOffsetOfBatch,
ComputePtrOffsetOfN,
0,
true,
EGlobalMemoryDataOperation,
CTranspose,
TailNum,
decltype(epilogue_args)>(
p_shared,
gemm_kernel_args[group_id].a_grid_desc_ak0_m_ak1_,
gemm_kernel_args[group_id].b_grid_desc_bk0_n_bk1_,
gemm_kernel_args[group_id].ds_grid_desc_mblock_mperblock_nblock_nperblock_,
gemm_kernel_args[group_id].e_grid_desc_mblock_mperblock_nblock_nperblock_,
gemm_kernel_args[group_id].block_2_ctile_map_,
compute_ptr_offset_of_batch,
compute_ptr_offset_of_n,
num_k_per_block,
karg,
epilogue_args);
}
else
{

GridwiseGemm::template Run<
GridwiseGemm::ConvRegime::BWD_DATA,
AGridDesc_AK0_M_AK1,
BGridDesc_BK0_N_BK1,
DsGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock,
EGridDesc_MBlock_MPerBlock_NBlock_NPerBlock,
decltype(gemm_kernel_args[group_id].block_2_ctile_map_),
ComputePtrOffsetOfBatch,
ComputePtrOffsetOfN,
0,
false,
EGlobalMemoryDataOperation,
CTranspose,
TailNum,
decltype(epilogue_args)>(
p_shared,
gemm_kernel_args[group_id].a_grid_desc_ak0_m_ak1_,
gemm_kernel_args[group_id].b_grid_desc_bk0_n_bk1_,
gemm_kernel_args[group_id].ds_grid_desc_mblock_mperblock_nblock_nperblock_,
gemm_kernel_args[group_id].e_grid_desc_mblock_mperblock_nblock_nperblock_,
gemm_kernel_args[group_id].block_2_ctile_map_,
compute_ptr_offset_of_batch,
compute_ptr_offset_of_n,
num_k_per_block,
karg,
epilogue_args);
}
}
}

Expand Down
58 changes: 30 additions & 28 deletions ...e/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_wmma_cshuffle_v3.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -77,34 +77,36 @@ __launch_bounds__(GridwiseGemm::MaxBlockSize, MinimumOccupancy)

constexpr index_t LDS_size =
GridwiseGemm::template GetSharedMemoryNumberOfByte<SelectedEpilogue>();
__shared__ char p_shared[LDS_size];

GridwiseGemm::template Run<GridwiseGemm::ConvRegime::BWD_WEIGHT,
AGridDesc_AK0_M_K1,
BGridDesc_BK0_N_K1,
ck::Tuple<>, // Empty tuple
CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock,
decltype(block_2_ctile_map_),
ComputePtrOffsetOfBatch,
ComputePtrOffsetOfBatch, // placeholder
1,
HasMainKBlockLoop,
CGlobalMemoryDataOperation,
false,
TailNum,
decltype(epilogue_args)>(
p_shared,
a_grid_desc_ak0_m_ak1,
b_grid_desc_bk0_n_bk1,
ck::Tuple<>(), // placeholder
c_grid_desc_mblock_mperblock_nblock_nperblock,
block_2_ctile_map_,
compute_ptr_offset_of_batch,
ComputePtrOffsetOfBatch{}, // placeholder
num_k_per_block,
karg,
epilogue_args);

if constexpr(LDS_size <= get_lds_size(get_device_arch()))
{
__shared__ char p_shared[LDS_size];

GridwiseGemm::template Run<GridwiseGemm::ConvRegime::BWD_WEIGHT,
AGridDesc_AK0_M_K1,
BGridDesc_BK0_N_K1,
ck::Tuple<>, // Empty tuple
CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock,
decltype(block_2_ctile_map_),
ComputePtrOffsetOfBatch,
ComputePtrOffsetOfBatch, // placeholder
1,
HasMainKBlockLoop,
CGlobalMemoryDataOperation,
false,
TailNum,
decltype(epilogue_args)>(
p_shared,
a_grid_desc_ak0_m_ak1,
b_grid_desc_bk0_n_bk1,
ck::Tuple<>(), // placeholder
c_grid_desc_mblock_mperblock_nblock_nperblock,
block_2_ctile_map_,
compute_ptr_offset_of_batch,
ComputePtrOffsetOfBatch{}, // placeholder
num_k_per_block,
karg,
epilogue_args);
}
#if defined(__gfx11__)
}
#endif
Expand Down
78 changes: 41 additions & 37 deletions ...ensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_abd_wmma_cshuffle_v3.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -101,43 +101,47 @@ __launch_bounds__(GridwiseGemm::MaxBlockSize, MinimumOccupancy)

constexpr index_t LDS_size =
GridwiseGemm::template GetSharedMemoryNumberOfByte<SelectedEpilogue>();
__shared__ char p_shared[LDS_size];

auto epilogue_args = SelectedEpilogue{};

const auto a_grid_desc_ak0_m_ak1 =
GridwiseGemm::MakeAGridDescriptor_AK0_M_AK1(a_grid_desc_m_k);

const auto b_grid_desc_bk0_n_bk1 =
GridwiseGemm::MakeBGridDescriptor_BK0_N_BK1(b_grid_desc_n_k);

const auto block_2_ctile_map_ = typename GridwiseGemm::Block2CTileMap{karg.M, karg.N, 4};

GridwiseGemm::template Run<GridwiseGemm::ConvRegime::FORWARD,
decltype(a_grid_desc_ak0_m_ak1),
decltype(b_grid_desc_bk0_n_bk1),
DsGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock,
EGridDesc_MBlock_MPerBlock_NBlock_NPerBlock,
decltype(block_2_ctile_map_),
ComputePtrOffset,
ComputePtrOffset,
0,
HasMainKBlockLoop,
EGlobalMemoryDataOperation,
false,
TailNum,
decltype(epilogue_args)>(
p_shared,
a_grid_desc_ak0_m_ak1,
b_grid_desc_bk0_n_bk1,
ds_grid_desc_mblock_mperblock_nblock_nperblock,
e_grid_desc_mblock_mperblock_nblock_nperblock,
block_2_ctile_map_,
compute_ptr_offset_of_batch,
compute_ptr_offset_of_n,
num_k_per_block,
karg,
epilogue_args);
if constexpr(LDS_size <= get_lds_size(get_device_arch()))
{
__shared__ char p_shared[LDS_size];

auto epilogue_args = SelectedEpilogue{};

const auto a_grid_desc_ak0_m_ak1 =
GridwiseGemm::MakeAGridDescriptor_AK0_M_AK1(a_grid_desc_m_k);

const auto b_grid_desc_bk0_n_bk1 =
GridwiseGemm::MakeBGridDescriptor_BK0_N_BK1(b_grid_desc_n_k);

const auto block_2_ctile_map_ =
typename GridwiseGemm::Block2CTileMap{karg.M, karg.N, 4};

GridwiseGemm::template Run<GridwiseGemm::ConvRegime::FORWARD,
decltype(a_grid_desc_ak0_m_ak1),
decltype(b_grid_desc_bk0_n_bk1),
DsGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock,
EGridDesc_MBlock_MPerBlock_NBlock_NPerBlock,
decltype(block_2_ctile_map_),
ComputePtrOffset,
ComputePtrOffset,
0,
HasMainKBlockLoop,
EGlobalMemoryDataOperation,
false,
TailNum,
decltype(epilogue_args)>(
p_shared,
a_grid_desc_ak0_m_ak1,
b_grid_desc_bk0_n_bk1,
ds_grid_desc_mblock_mperblock_nblock_nperblock,
e_grid_desc_mblock_mperblock_nblock_nperblock,
block_2_ctile_map_,
compute_ptr_offset_of_batch,
compute_ptr_offset_of_n,
num_k_per_block,
karg,
epilogue_args);
}

#if defined(__gfx11__)
}
Expand Down
Loading
Loading