mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-03-16 05:07:37 +00:00
22de6a19d9
[CK_BUILDER] ck builder conv transfer fix ## Motivation This PR fixes how CK Builder is validating transfer vector size and adds proper validation for LDS transfer vector size as well. ## Changes: * [__source vector dim__] -- Before this PR the data transfer validation logic didn't allow to set the source vectorized dimension to 1. However there are CK instances that are doing this when the group merging is used. This is used only for `DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle` kernel. * [__valid vector size__] -- Before this PR the validation logic concerned only single instruction maximum vector size. However our buffer loading logic has implemented support for loading more values through multiple buffer instructions. This again was discovered to be used in some of the convolution instances. Thus this behavior was reflected in validation logic. * [__valid LDS vector size__] -- Before this PR the LDS vector size validation was done in the same way as VMEM. This PR adds proper LDS vector size validation based on the available LDS instruction sizes. ## Test Plan Run CK BUILDER conv fwd factories tests ## Test Result All CK BUILDER conv fwd factories work (except DL one & ck tile since they're not yet added now) ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
17 lines
435 B
Python
Executable File
17 lines
435 B
Python
Executable File
# Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
|
|
# SPDX-License-Identifier: MIT
|
|
|
|
import os
|
|
|
|
root_dir = os.getcwd()
|
|
ck_tile_include = root_dir + "/projects/composablekernel/include/ck_tile"
|
|
ck_tile_example = root_dir + "/projects/composablekernel/example/ck_tile"
|
|
|
|
# Run for include
|
|
os.chdir(ck_tile_include)
|
|
_ = os.system("python remod.py")
|
|
|
|
# Run for example
|
|
os.chdir(ck_tile_example)
|
|
_ = os.system("python remod.py")
|