ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-17 09:08:35 +00:00
Code Issues Packages Projects Releases Wiki Activity

[rocm-libraries] ROCm/rocm-libraries#7734 (commit 03ffb9d)

Browse Source 
[CK] Grouped Convolution Global Load/Store instances

## Motivation

Support global load and store in grouped convolutions using instance
factory.

## Technical Details

- add new instances for each direction
- add new tests for large cases

## Test Plan

New test for large cases

## Test Result

pending

## Submission Checklist

- [x] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
AICK-1255
This commit is contained in:
Bartłomiej Kocot
2026-06-06 22:52:59 +00:00
committed by assistant-librarian[bot]
parent 2c363870d9
commit 28f2966762
91 changed files with 3469 additions and 1638 deletions
Download Patch File Download Diff File Expand all files Collapse all files

16
test/gpu_reference/gpu_reference_utils.hpp
View File

@@ -404,22 +404,22 @@ bool test_conv_fwd_with_d_tensor_impl(const ck::utils::conv::ConvParam& params,
using WeiElementOp = tensor_operation::element_wise::PassThrough;
// Create D tensor lengths and strides for GPU reference
std::vector<index_t> d_lengths_vec(NDimSpatial + 3);
std::vector<long_index_t> d_lengths_vec(NDimSpatial + 3);
d_lengths_vec[0] = params.G_;
d_lengths_vec[1] = params.N_;
d_lengths_vec[2] = params.K_;
for(index_t i = 0; i < NDimSpatial; ++i)
{
d_lengths_vec[3 + i] = static_cast<index_t>(params.output_spatial_lengths_[i]);
d_lengths_vec[3 + i] = static_cast<long_index_t>(params.output_spatial_lengths_[i]);
}
std::vector<index_t> d_strides_vec =
std::vector<long_index_t> d_strides_vec =
ref::compute_conv_tensor_strides<OutLayout>(d_lengths_vec, params.num_dim_spatial_);
std::array<const OutDataType*, 1> d_ptrs = {
reinterpret_cast<const OutDataType*>(d_dev.GetDeviceBuffer())};
std::array<std::vector<index_t>, 1> d_lengths = {d_lengths_vec};
std::array<std::vector<index_t>, 1> d_strides = {d_strides_vec};
std::array<std::vector<long_index_t>, 1> d_lengths = {d_lengths_vec};
std::array<std::vector<long_index_t>, 1> d_strides = {d_strides_vec};
// Call GPU reference with D tensor
std::array<const InDataType*, 1> in_ptrs = {
@@ -536,9 +536,9 @@ bool test_conv_fwd_with_multi_ab_impl(const ck::utils::conv::ConvParam& params,
std::array<const WeiDataType*, 2> wei_ptrs = {
reinterpret_cast<const WeiDataType*>(weight_dev.GetDeviceBuffer()),
reinterpret_cast<const WeiDataType*>(b_extra_dev.GetDeviceBuffer())};
std::array<const OutDataType*, 0> d_ptrs = {};
std::array<std::vector<index_t>, 0> d_lengths = {};
std::array<std::vector<index_t>, 0> d_strides = {};
std::array<const OutDataType*, 0> d_ptrs = {};
std::array<std::vector<long_index_t>, 0> d_lengths = {};
std::array<std::vector<long_index_t>, 0> d_strides = {};
ref::naive_conv_fwd_multi_abd<1, 1, 0, InLayout, WeiLayout, OutLayout>(
in_ptrs,