ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-04-20 06:49:15 +00:00
Code Issues Packages Projects Releases Wiki Activity

[CK][Examples] Fixing stride issues in ck examples by workaround - Bypassing hostTensor validation.

Browse Source 
Signed-off-by: Michal Kulikowski <Michal.Kulikowski@amd.com>
This commit is contained in:
Michal Kulikowski
2025-10-16 13:01:24 +02:00
committed by Michał Kulikowski
parent 0d3860dfdb
commit b9789a0742
33 changed files with 195 additions and 199 deletions
Download Patch File Download Diff File Expand all files Collapse all files

7
example/60_gemm_multi_ABD/gemm_multi_ABD_wmma_bias_fastgelu_bf16_i8.cpp
View File

@@ -27,7 +27,8 @@ using BF16 = ck::bhalf_t;
using I8 = int8_t;
using F32 = float;
using Row = ck::tensor_layout::gemm::RowMajor;
using Row = ck::tensor_layout::gemm::RowMajor;
using Bypass = ck::tensor_layout::BypassLayoutVerification;
using A0DataType = BF16;
using AsDataType = ck::Tuple<A0DataType>;
@@ -161,11 +162,11 @@ int main(int argc, char* argv[])
if(std::is_same<decltype(layout), ck::tensor_layout::gemm::RowMajor>::value)
{
return HostTensorDescriptor({row, col}, {stride, 1_uz});
return HostTensorDescriptor({row, col}, {stride, 1_uz}, Bypass{});
}
else
{
return HostTensorDescriptor({row, col}, {1_uz, stride});
return HostTensorDescriptor({row, col}, {1_uz, stride}, Bypass{});
}
};

7
example/60_gemm_multi_ABD/gemm_multi_ABD_wmma_fastgelu_bf16_i8.cpp
View File

@@ -27,7 +27,8 @@ using BF16 = ck::bhalf_t;
using I8 = int8_t;
using F32 = float;
using Row = ck::tensor_layout::gemm::RowMajor;
using Row = ck::tensor_layout::gemm::RowMajor;
using Bypass = ck::tensor_layout::BypassLayoutVerification;
using A0DataType = BF16;
using AsDataType = ck::Tuple<A0DataType>;
@@ -157,11 +158,11 @@ int main(int argc, char* argv[])
if(std::is_same<decltype(layout), ck::tensor_layout::gemm::RowMajor>::value)
{
return HostTensorDescriptor({row, col}, {stride, 1_uz});
return HostTensorDescriptor({row, col}, {stride, 1_uz}, Bypass{});
}
else
{
return HostTensorDescriptor({row, col}, {1_uz, stride});
return HostTensorDescriptor({row, col}, {1_uz, stride}, Bypass{});
}
};

7
example/60_gemm_multi_ABD/gemm_multi_ABD_wmma_fp16.cpp
View File

@@ -24,7 +24,8 @@ using S = ck::Sequence<Is...>;
using F16 = ck::half_t;
using F32 = float;
using Row = ck::tensor_layout::gemm::RowMajor;
using Row = ck::tensor_layout::gemm::RowMajor;
using Bypass = ck::tensor_layout::BypassLayoutVerification;
using PassThrough = ck::tensor_operation::element_wise::PassThrough;
@@ -220,11 +221,11 @@ int main(int argc, char* argv[])
if(std::is_same<decltype(layout), ck::tensor_layout::gemm::RowMajor>::value)
{
return HostTensorDescriptor({row, col}, {stride, 1_uz});
return HostTensorDescriptor({row, col}, {stride, 1_uz}, Bypass{});
}
else
{
return HostTensorDescriptor({row, col}, {1_uz, stride});
return HostTensorDescriptor({row, col}, {1_uz, stride}, Bypass{});
}
};

7
example/60_gemm_multi_ABD/gemm_multi_ABD_wmma_multiply_bias_fastgelu_bf16_i8.cpp
View File

@@ -27,7 +27,8 @@ using BF16 = ck::bhalf_t;
using I8 = int8_t;
using F32 = float;
using Row = ck::tensor_layout::gemm::RowMajor;
using Row = ck::tensor_layout::gemm::RowMajor;
using Bypass = ck::tensor_layout::BypassLayoutVerification;
using A0DataType = BF16;
using AsDataType = ck::Tuple<A0DataType>;
@@ -160,11 +161,11 @@ int main(int argc, char* argv[])
if(std::is_same<decltype(layout), ck::tensor_layout::gemm::RowMajor>::value)
{
return HostTensorDescriptor({row, col}, {stride, 1_uz});
return HostTensorDescriptor({row, col}, {stride, 1_uz}, Bypass{});
}
else
{
return HostTensorDescriptor({row, col}, {1_uz, stride});
return HostTensorDescriptor({row, col}, {1_uz, stride}, Bypass{});
}
};

37
example/60_gemm_multi_ABD/gemm_multi_ABD_xdl_bias_fastgelu_bf16_i8.cpp
View File

@@ -28,8 +28,9 @@ using BF16 = ck::bhalf_t;
using I8 = int8_t;
using F32 = float;
using Row = ck::tensor_layout::gemm::RowMajor;
using Col = ck::tensor_layout::gemm::ColumnMajor;
using Row = ck::tensor_layout::gemm::RowMajor;
using Col = ck::tensor_layout::gemm::ColumnMajor;
using Bypass = ck::tensor_layout::BypassLayoutVerification;
using A0DataType = BF16;
using AsDataType = ck::Tuple<A0DataType>;
@@ -121,27 +122,19 @@ int main(int argc, char* argv[])
exit(0);
}
auto f_host_tensor_descriptor = [](std::size_t row,
std::size_t col,
ck::index_t& stride,
auto layout) {
using namespace ck::literals;
auto f_host_tensor_descriptor =
[](std::size_t row, std::size_t col, std::size_t stride, auto layout) {
using namespace ck::literals;
if(std::is_same<decltype(layout), ck::tensor_layout::gemm::RowMajor>::value)
{
auto desc = HostTensorDescriptor({row, col}, {static_cast<std::size_t>(stride), 1_uz});
if(stride <= 0)
stride = desc.GetStrides()[0];
return desc;
}
else
{
auto desc = HostTensorDescriptor({row, col}, {1_uz, static_cast<std::size_t>(stride)});
if(stride <= 0)
stride = desc.GetStrides()[1];
return desc;
}
};
if(std::is_same<decltype(layout), ck::tensor_layout::gemm::RowMajor>::value)
{
return HostTensorDescriptor({row, col}, {stride, 1_uz}, Bypass{});
}
else
{
return HostTensorDescriptor({row, col}, {1_uz, stride}, Bypass{});
}
};
Tensor<A0DataType> a0_m_k(f_host_tensor_descriptor(M, K, StrideA, A0Layout{}));
Tensor<B0DataType> b0_k_n(f_host_tensor_descriptor(K, N, StrideB, B0Layout{}));

37
example/60_gemm_multi_ABD/gemm_multi_ABD_xdl_fastgelu_bf16_i8.cpp
View File

@@ -28,8 +28,9 @@ using BF16 = ck::bhalf_t;
using I8 = int8_t;
using F32 = float;
using Row = ck::tensor_layout::gemm::RowMajor;
using Col = ck::tensor_layout::gemm::ColumnMajor;
using Row = ck::tensor_layout::gemm::RowMajor;
using Col = ck::tensor_layout::gemm::ColumnMajor;
using Bypass = ck::tensor_layout::BypassLayoutVerification;
using A0DataType = BF16;
using AsDataType = ck::Tuple<A0DataType>;
@@ -121,27 +122,19 @@ int main(int argc, char* argv[])
exit(0);
}
auto f_host_tensor_descriptor = [](std::size_t row,
std::size_t col,
ck::index_t& stride,
auto layout) {
using namespace ck::literals;
auto f_host_tensor_descriptor =
[](std::size_t row, std::size_t col, std::size_t stride, auto layout) {
using namespace ck::literals;
if(std::is_same<decltype(layout), ck::tensor_layout::gemm::RowMajor>::value)
{
auto desc = HostTensorDescriptor({row, col}, {static_cast<std::size_t>(stride), 1_uz});
if(stride <= 0)
stride = desc.GetStrides()[0];
return desc;
}
else
{
auto desc = HostTensorDescriptor({row, col}, {1_uz, static_cast<std::size_t>(stride)});
if(stride <= 0)
stride = desc.GetStrides()[1];
return desc;
}
};
if(std::is_same<decltype(layout), ck::tensor_layout::gemm::RowMajor>::value)
{
return HostTensorDescriptor({row, col}, {stride, 1_uz}, Bypass{});
}
else
{
return HostTensorDescriptor({row, col}, {1_uz, stride}, Bypass{});
}
};
Tensor<A0DataType> a0_m_k(f_host_tensor_descriptor(M, K, StrideA, A0Layout{}));
Tensor<B0DataType> b0_k_n(f_host_tensor_descriptor(K, N, StrideB, B0Layout{}));

37
example/60_gemm_multi_ABD/gemm_multi_ABD_xdl_multiply_bias_fastgelu_bf16_i8.cpp
View File

@@ -28,8 +28,9 @@ using BF16 = ck::bhalf_t;
using I8 = int8_t;
using F32 = float;
using Row = ck::tensor_layout::gemm::RowMajor;
using Col = ck::tensor_layout::gemm::ColumnMajor;
using Row = ck::tensor_layout::gemm::RowMajor;
using Col = ck::tensor_layout::gemm::ColumnMajor;
using Bypass = ck::tensor_layout::BypassLayoutVerification;
using A0DataType = BF16;
using AsDataType = ck::Tuple<A0DataType>;
@@ -120,27 +121,19 @@ int main(int argc, char* argv[])
exit(0);
}
auto f_host_tensor_descriptor = [](std::size_t row,
std::size_t col,
ck::index_t& stride,
auto layout) {
using namespace ck::literals;
auto f_host_tensor_descriptor =
[](std::size_t row, std::size_t col, std::size_t stride, auto layout) {
using namespace ck::literals;
if(std::is_same<decltype(layout), ck::tensor_layout::gemm::RowMajor>::value)
{
auto desc = HostTensorDescriptor({row, col}, {static_cast<std::size_t>(stride), 1_uz});
if(stride <= 0)
stride = desc.GetStrides()[0];
return desc;
}
else
{
auto desc = HostTensorDescriptor({row, col}, {1_uz, static_cast<std::size_t>(stride)});
if(stride <= 0)
stride = desc.GetStrides()[1];
return desc;
}
};
if(std::is_same<decltype(layout), ck::tensor_layout::gemm::RowMajor>::value)
{
return HostTensorDescriptor({row, col}, {stride, 1_uz}, Bypass{});
}
else
{
return HostTensorDescriptor({row, col}, {1_uz, stride}, Bypass{});
}
};
Tensor<A0DataType> a0_m_k(f_host_tensor_descriptor(M, K, StrideA, A0Layout{}));
Tensor<B0DataType> b0_k_n(f_host_tensor_descriptor(K, N, StrideB, B0Layout{}));