ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-17 19:40:04 +00:00
Code Issues Packages Projects Releases Wiki Activity

Conv3d new (#94)

Browse Source 
* conv3d compiles but has memory error

* conv3d works

* fix performance issue by using __builtin_amdgc_readfirstlane

* change MakeBlock2CTileMap to MakeDefaultBlock2CTileMap; change c_blockid_to* to cblockid_to*

* clang-format

* remove CK_EXPERIMENTAL_PASS_TENSOR_DECRIPTOR_BY_*; moved wrapper into DeviceConv3d

* format

* remove useless marc

* add comment

Co-authored-by: Chao Liu <chao.liu2@amd.com>

[ROCm/composable_kernel commit: 6dfb92bbef]
This commit is contained in:
Jianfeng Yan
2022-02-22 22:45:28 -06:00
committed by GitHub
parent 4dd24cf943
commit c2e3fa5c91
61 changed files with 2255 additions and 1725 deletions
Download Patch File Download Diff File Expand all files Collapse all files

57
example/10_conv3d_fwd_xdl/README.md Normal file
View File

@@ -0,0 +1,57 @@
# Instructions for ```conv3d_fwd_xdl``` Example
## Docker script
```bash
docker run \
-it \
--rm \
--privileged \
--group-add sudo \
-w /root/workspace \
-v ${PATH_TO_LOCAL_WORKSPACE}:/root/workspace \
rocm/tensorflow:rocm4.3.1-tf2.6-dev \
/bin/bash
```
## Build ```conv3d_fwd_xdl```
```bash
mkdir build && cd build
```
```bash
# Need to specify target ID, example below is gfx908
cmake \
-D BUILD_DEV=OFF \
-D CMAKE_BUILD_TYPE=Release \
-D CMAKE_CXX_FLAGS="-DCK_AMD_GPU_GFX908 --amdgpu-target=gfx908 -O3 " \
-D CMAKE_CXX_COMPILER=/opt/rocm/bin/hipcc \
-D CMAKE_PREFIX_PATH=/opt/rocm \
..
```
```bash
make -j conv3d_fwd_xdl
```
## Run ```conv3d_fwd_xdl```
```bash
#arg1: verification (0=no, 1=yes)
#arg2: initialization (0=no init, 1=integer value, 2=decimal value)
#arg3: run kernel # of times (>1)
#arg4 to 24: N, K, C, Z, Y, X, Di, Hi, Wi, Sz, Sy, Sx, Dz, Dy, Dx, leftPz, LeftPy, LeftPx, RightPz, RightPy, RightPx
./example/conv3d_fwd_xdl 0 1 5
```
Result (MI100 dynamic frequency)
```
in: dim 5, lengths {4, 71, 71, 71, 192}, strides {68718912, 967872, 13632, 192, 1}
wei: dim 5, lengths {256, 3, 3, 3, 192}, strides {5184, 1728, 576, 192, 1}
out: dim 5, lengths {4, 36, 36, 36, 256}, strides {11943936, 331776, 9216, 256, 1}
a_grid_desc_b_k0_m_k1{1, 648, 186624, 8}
b_grid_desc_b_k0_n_k1{1, 648, 256, 8}
launch_and_time_kernel: grid_dim {1458, 1, 1}, block_dim {256, 1, 1}
Warm up
Start running 5 times...
Perf: 4.49466 ms, 110.206 TFlops, 144.161 GB/s
```

281
example/10_conv3d_fwd_xdl/conv3d_fwd_xdl.cpp Normal file
View File

@@ -0,0 +1,281 @@
#include <iostream>
#include <numeric>
#include <initializer_list>
#include <cstdlib>
#include <stdlib.h>
#include <half.hpp>
#include "config.hpp"
#include "print.hpp"
#include "device.hpp"
#include "host_tensor.hpp"
#include "host_tensor_generator.hpp"
#include "host_gemm.hpp"
#include "device_tensor.hpp"
#include "device_base.hpp"
#include "device_conv3d_fwd_xdl_ndhwc_kzyxc_ndhwk.hpp"
#include "device_conv3d_fwd_naive_ndhwc_kzyxc_ndhwk.hpp"
#include "convolution_utility.hpp"
// convolution data type
using InDataType = ck::half_t;
using WeiDataType = ck::half_t;
using OutDataType = ck::half_t;
using AccDataType = float;
using InElementOp = ck::tensor_operation::element_wise::PassThrough;
using WeiElementOp = ck::tensor_operation::element_wise::PassThrough;
using OutElementOp = ck::tensor_operation::element_wise::PassThrough;
using F16 = ck::half_t;
using F32 = float;
template <ck::index_t... Is>
using S = ck::Sequence<Is...>;
using InLayout = ck::tensor_layout::convolution::NDHWC;
using WeiLayout = ck::tensor_layout::convolution::KZYXC;
using OutLayout = ck::tensor_layout::convolution::NDHWK;
static constexpr auto ConvFwdDefault =
ck::tensor_operation::device::ConvolutionForwardSpecialization_t::Default;
using DeviceConv3dFwdInstance = ck::tensor_operation::device::
DeviceConv3dFwdXdl_Input_N_Di_Hi_Wi_C_Weight_K_Z_Y_X_C_Output_N_Do_Ho_Wo_K<
InDataType, // InData
WeiDataType, // WeiData
OutDataType, // OutData
AccDataType, // AccData
InElementOp, // InElementwise Operation
WeiElementOp, // WeiElementwise Operation
OutElementOp, // OutElementwise Operation
ConvFwdDefault, // ConvForwardSpecialization
256, // BlockSize
128, // MPerBlock
256, // NPerBlock
4, // K0PerBlock
8, // K1. K0PerBlock * K1 = KPerBlock
32, // MPerXDL
32, // NPerXDL. Each XDL computes a matrix of size (MPerXDL, NPerBlock)
2, // MXdlPerWave
4, // NXdlPerWave
S<4, 64, 1>, // ABlockTransferThreadClusterLengths_K0_M_K1
S<1, 0, 2>, // ABlockTransferThreadClusterArrangeOrder
S<1, 0, 2>, // ABlockTransferSrcAccessOrder
2, // ABlockTransferSrcVectorDim
8, // ABlockTransferSrcScalarPerVector
8, // ABlockTransferDstScalarPerVector_K1
true, // ABlockLdsAddExtraM
S<4, 64, 1>, // BBlockTransferThreadClusterLengths_K0_N_K1
S<1, 0, 2>, // BBlockTransferThreadClusterArrangeOrder
S<1, 0, 2>, // BBlockTransferSrcAccessOrder
2, // BBlockTransferSrcVectorDim
8, // BBlockTransferSrcScalarPerVector
8, // BBlockTransferDstScalarPerVector_K1
true, // BBlockLdsAddExtraN
7, // CThreadTransferSrcDstVectorDim
1>; // CThreadTransferDstScalarPerVector
int main(int argc, char* argv[])
{
bool do_verification = false;
int init_method = 0;
int nrepeat = 5;
// convolution shape
ck::index_t N = 4;
ck::index_t K = 256;
ck::index_t C = 192;
std::vector<ck::index_t> in_spatial_lengths = {71, 71, 71};
std::vector<ck::index_t> filter_spatial_lengths = {3, 3, 3};
std::vector<ck::index_t> conv_filter_strides = {2, 2, 2};
std::vector<ck::index_t> conv_filter_dilations = {1, 1, 1};
std::vector<ck::index_t> in_left_pads = {1, 1, 1};
std::vector<ck::index_t> in_right_pads = {1, 1, 1};
if(argc == 4)
{
do_verification = std::stoi(argv[1]);
init_method = std::stoi(argv[2]);
nrepeat = std::stoi(argv[3]);
}
else if(argc == 25)
{
do_verification = std::stoi(argv[1]);
init_method = std::stoi(argv[2]);
nrepeat = std::stoi(argv[3]);
N = std::stoi(argv[4]);
K = std::stoi(argv[5]);
C = std::stoi(argv[6]);
filter_spatial_lengths[0] = std::stoi(argv[7]);
filter_spatial_lengths[1] = std::stoi(argv[8]);
filter_spatial_lengths[2] = std::stoi(argv[9]);
in_spatial_lengths[0] = std::stoi(argv[10]);
in_spatial_lengths[1] = std::stoi(argv[11]);
in_spatial_lengths[2] = std::stoi(argv[12]);
conv_filter_strides[0] = std::stoi(argv[13]);
conv_filter_strides[1] = std::stoi(argv[14]);
conv_filter_strides[2] = std::stoi(argv[15]);
conv_filter_dilations[0] = std::stoi(argv[16]);
conv_filter_dilations[1] = std::stoi(argv[17]);
conv_filter_dilations[2] = std::stoi(argv[18]);
in_left_pads[0] = std::stoi(argv[19]);
in_left_pads[1] = std::stoi(argv[20]);
in_left_pads[2] = std::stoi(argv[21]);
in_right_pads[0] = std::stoi(argv[22]);
in_right_pads[1] = std::stoi(argv[23]);
in_right_pads[2] = std::stoi(argv[24]);
}
else
{
printf("Usage: 3 or 24 input arguments\n");
printf(" arg1: verification (0=no, 1=yes)\n");
printf(" arg2: initialization (0=no init, 1=integer value, 2=decimal value)\n");
printf(" arg3: run kernel # of times (>1)\n");
printf(" arg4 to 24: N, K, C, Z, Y, X, Di, Hi, Wi, Sz, Sy, Sz, Dz, Dy, Dx, LeftPz, LeftPy, "
"LeftPz, RightPz, RightPy, RightPx\n");
exit(0);
}
auto conv3d = DeviceConv3dFwdInstance{};
const auto out_spatial_lengths =
ck::tensor_operation::ConvolutionUtility::ComputeOutputSpatialLengths(
in_spatial_lengths,
filter_spatial_lengths,
conv_filter_strides,
conv_filter_dilations,
in_left_pads,
in_right_pads);
Tensor<InDataType> in(
{N, in_spatial_lengths[0], in_spatial_lengths[1], in_spatial_lengths[2], C});
Tensor<WeiDataType> wei(
{K, filter_spatial_lengths[0], filter_spatial_lengths[1], filter_spatial_lengths[2], C});
Tensor<OutDataType> out(
{N, out_spatial_lengths[0], out_spatial_lengths[1], out_spatial_lengths[2], K});
std::cout << "in: " << in.mDesc << std::endl;
std::cout << "wei: " << wei.mDesc << std::endl;
std::cout << "out: " << out.mDesc << std::endl;
switch(init_method)
{
case 0: break;
case 1:
in.GenerateTensorValue(GeneratorTensor_2<InDataType>{-5, 5});
wei.GenerateTensorValue(GeneratorTensor_2<WeiDataType>{-5, 5});
break;
default:
in.GenerateTensorValue(GeneratorTensor_3<InDataType>{0.0, 1.0});
wei.GenerateTensorValue(GeneratorTensor_3<WeiDataType>{-0.5, 0.5});
}
DeviceMem in_device_buf(sizeof(InDataType) * in.mDesc.GetElementSpace());
DeviceMem wei_device_buf(sizeof(WeiDataType) * wei.mDesc.GetElementSpace());
DeviceMem out_device_buf(sizeof(OutDataType) * out.mDesc.GetElementSpace());
in_device_buf.ToDevice(in.mData.data());
wei_device_buf.ToDevice(wei.mData.data());
// do Convolution
auto invoker = conv3d.MakeInvoker();
auto argument = conv3d.MakeArgument(static_cast<InDataType*>(in_device_buf.GetDeviceBuffer()),
static_cast<WeiDataType*>(wei_device_buf.GetDeviceBuffer()),
static_cast<OutDataType*>(out_device_buf.GetDeviceBuffer()),
N,
K,
C,
in_spatial_lengths,
filter_spatial_lengths,
out_spatial_lengths,
conv_filter_strides,
conv_filter_dilations,
in_left_pads,
in_right_pads,
InElementOp{},
WeiElementOp{},
OutElementOp{});
if(!conv3d.IsSupportedArgument(argument))
{
throw std::runtime_error(
"wrong! device_conv3d with the specified compilation parameters does "
"not support this GEMM problem");
}
float ave_time = invoker.Run(argument, nrepeat);
const auto Di = in_spatial_lengths[0];
const auto Hi = in_spatial_lengths[1];
const auto Wi = in_spatial_lengths[2];
const auto Do = out_spatial_lengths[0];
const auto Ho = out_spatial_lengths[1];
const auto Wo = out_spatial_lengths[2];
const auto Z = filter_spatial_lengths[0];
const auto Y = filter_spatial_lengths[1];
const auto X = filter_spatial_lengths[2];
std::size_t flop = std::size_t(2) * N * K * Do * Ho * Wo * C * Z * Y * X;
std::size_t num_btype = sizeof(InDataType) * N * Di * Hi * Wi * C +
sizeof(WeiDataType) * K * Z * Y * X * C +
sizeof(OutDataType) * N * Do * Ho * Wo * K;
float tflops = static_cast<float>(flop) / 1.E9 / ave_time;
float gb_per_sec = num_btype / 1.E6 / ave_time;
std::cout << "Perf: " << ave_time << " ms, " << tflops << " TFlops, " << gb_per_sec << " GB/s"
<< std::endl;
out_device_buf.FromDevice(out.mData.data());
if(do_verification)
{
DeviceMem out_ref_device_buf(sizeof(OutDataType) * N * Do * Ho * Wo * K);
using DeviceConv3dFwdNaive = ck::tensor_operation::device::
DeviceConv3dFwdNaive_Input_N_Di_Hi_Wi_C_Weight_K_Z_Y_X_C_Output_N_Do_Ho_Wo_K<
InDataType,
WeiDataType,
OutDataType,
AccDataType,
InElementOp,
WeiElementOp,
OutElementOp>;
auto conv3d_naive = DeviceConv3dFwdNaive{};
auto invoker_naive = conv3d_naive.MakeInvoker();
auto argument_naive = conv3d_naive.MakeArgument(
static_cast<InDataType*>(in_device_buf.GetDeviceBuffer()),
static_cast<WeiDataType*>(wei_device_buf.GetDeviceBuffer()),
static_cast<OutDataType*>(out_ref_device_buf.GetDeviceBuffer()),
N,
K,
C,
in_spatial_lengths,
filter_spatial_lengths,
out_spatial_lengths,
conv_filter_strides,
conv_filter_dilations,
in_left_pads,
in_right_pads,
InElementOp{},
WeiElementOp{},
OutElementOp{});
if(!conv3d_naive.IsSupportedArgument(argument_naive))
{
throw std::runtime_error(
"wrong! device_conv3d_naive does NOT support the specified compilation parameters");
}
invoker_naive.Run(argument_naive);
Tensor<OutDataType> out_ref(
{N, out_spatial_lengths[0], out_spatial_lengths[1], out_spatial_lengths[2], K});
out_ref_device_buf.FromDevice(out_ref.mData.data());
check_error(out_ref, out);
}
return 0;
}

3
example/1_gemm_xdl/gemm_xdl.cpp
View File

@@ -160,7 +160,6 @@ int main(int argc, char* argv[])
a_m_k_device_buf.ToDevice(a_m_k.mData.data());
b_k_n_device_buf.ToDevice(b_k_n.mData.data());
c_m_n_device_buf.ToDevice(c_m_n_device_result.mData.data());
auto a_element_op = AElementOp{};
auto b_element_op = BElementOp{};
@@ -216,4 +215,6 @@ int main(int argc, char* argv[])
check_error(c_m_n_host_result, c_m_n_device_result);
}
return 0;
}

29
example/4_conv2d_fwd_xdl/conv2d_fwd_xdl.cpp
View File

@@ -14,6 +14,7 @@
#include "element_wise_operation.hpp"
#include "device_conv2d_fwd_xdl_c_shuffle_nhwc_kyxc_nhwk.hpp"
#include "reference_conv_fwd.hpp"
#include "convolution_utility.hpp"
using InDataType = ck::half_t;
using WeiDataType = ck::half_t;
@@ -138,16 +139,20 @@ int main(int argc, char* argv[])
exit(0);
}
const ck::index_t YEff = (Y - 1) * conv_dilation_h + 1;
const ck::index_t XEff = (X - 1) * conv_dilation_w + 1;
const std::vector<ck::index_t> conv_filter_strides{conv_stride_h, conv_stride_w};
const std::vector<ck::index_t> conv_filter_dilations{conv_dilation_h, conv_dilation_w};
const std::vector<ck::index_t> input_left_pads{in_left_pad_h, in_left_pad_w};
const std::vector<ck::index_t> input_right_pads{in_right_pad_h, in_right_pad_w};
const auto output_spatial_lengths =
ck::tensor_operation::ConvolutionUtility::ComputeOutputSpatialLengths({Hi, Wi},
{Y, X},
conv_filter_strides,
conv_filter_dilations,
input_left_pads,
input_right_pads);
const ck::index_t Ho = (Hi + in_left_pad_h + in_right_pad_h - YEff) / conv_stride_h + 1;
const ck::index_t Wo = (Wi + in_left_pad_w + in_right_pad_w - XEff) / conv_stride_w + 1;
const std::vector<ck::index_t> conv_filter_strides{{conv_stride_h, conv_stride_w}};
const std::vector<ck::index_t> conv_filter_dilations{{conv_dilation_h, conv_dilation_w}};
const std::vector<ck::index_t> input_left_pads{{in_left_pad_h, in_left_pad_w}};
const std::vector<ck::index_t> input_right_pads{{in_right_pad_h, in_right_pad_w}};
const ck::index_t Ho = output_spatial_lengths[0];
const ck::index_t Wo = output_spatial_lengths[1];
// tensor layout
auto f_host_tensor_descriptor = [](std::size_t N_,
@@ -214,9 +219,9 @@ int main(int argc, char* argv[])
N,
K,
C,
std::vector<ck::index_t>{{Hi, Wi}},
std::vector<ck::index_t>{{Y, X}},
std::vector<ck::index_t>{{Ho, Wo}},
std::vector<ck::index_t>{Hi, Wi},
std::vector<ck::index_t>{Y, X},
std::vector<ck::index_t>{Ho, Wo},
conv_filter_strides,
conv_filter_dilations,
input_left_pads,

29
example/5_conv2d_fwd_xdl_bias_relu/conv2d_fwd_xdl_bias_relu.cpp
View File

@@ -14,6 +14,7 @@
#include "element_wise_operation.hpp"
#include "device_conv2d_fwd_xdl_c_shuffle_bias_activation_nhwc_kyxc_nhwk.hpp"
#include "reference_conv_fwd_bias_activation.hpp"
#include "convolution_utility.hpp"
using InDataType = ck::half_t;
using WeiDataType = ck::half_t;
@@ -146,16 +147,20 @@ int main(int argc, char* argv[])
exit(0);
}
const ck::index_t YEff = (Y - 1) * conv_dilation_h + 1;
const ck::index_t XEff = (X - 1) * conv_dilation_w + 1;
const std::vector<ck::index_t> conv_filter_strides{conv_stride_h, conv_stride_w};
const std::vector<ck::index_t> conv_filter_dilations{conv_dilation_h, conv_dilation_w};
const std::vector<ck::index_t> input_left_pads{in_left_pad_h, in_left_pad_w};
const std::vector<ck::index_t> input_right_pads{in_right_pad_h, in_right_pad_w};
const auto output_spatial_lengths =
ck::tensor_operation::ConvolutionUtility::ComputeOutputSpatialLengths({Hi, Wi},
{Y, X},
conv_filter_strides,
conv_filter_dilations,
input_left_pads,
input_right_pads);
const ck::index_t Ho = (Hi + in_left_pad_h + in_right_pad_h - YEff) / conv_stride_h + 1;
const ck::index_t Wo = (Wi + in_left_pad_w + in_right_pad_w - XEff) / conv_stride_w + 1;
const std::vector<ck::index_t> conv_filter_strides{{conv_stride_h, conv_stride_w}};
const std::vector<ck::index_t> conv_filter_dilations{{conv_dilation_h, conv_dilation_w}};
const std::vector<ck::index_t> input_left_pads{{in_left_pad_h, in_left_pad_w}};
const std::vector<ck::index_t> input_right_pads{{in_right_pad_h, in_right_pad_w}};
const ck::index_t Ho = output_spatial_lengths[0];
const ck::index_t Wo = output_spatial_lengths[1];
// tensor layout
auto f_host_tensor_descriptor = [](std::size_t N_,
@@ -232,9 +237,9 @@ int main(int argc, char* argv[])
N,
K,
C,
std::vector<ck::index_t>{{Hi, Wi}},
std::vector<ck::index_t>{{Y, X}},
std::vector<ck::index_t>{{Ho, Wo}},
std::vector<ck::index_t>{Hi, Wi},
std::vector<ck::index_t>{Y, X},
std::vector<ck::index_t>{Ho, Wo},
conv_filter_strides,
conv_filter_dilations,
input_left_pads,

29
example/6_conv2d_fwd_xdl_bias_relu_add/conv2d_fwd_xdl_bias_relu_add.cpp
View File

@@ -14,6 +14,7 @@
#include "element_wise_operation.hpp"
#include "device_conv2d_fwd_xdl_c_shuffle_bias_activation_add_nhwc_kyxc_nhwk.hpp"
#include "reference_conv_fwd_bias_activation_add.hpp"
#include "convolution_utility.hpp"
using InDataType = ck::half_t;
using WeiDataType = ck::half_t;
@@ -143,16 +144,20 @@ int main(int argc, char* argv[])
exit(0);
}
const ck::index_t YEff = (Y - 1) * conv_dilation_h + 1;
const ck::index_t XEff = (X - 1) * conv_dilation_w + 1;
const std::vector<ck::index_t> conv_filter_strides{conv_stride_h, conv_stride_w};
const std::vector<ck::index_t> conv_filter_dilations{conv_dilation_h, conv_dilation_w};
const std::vector<ck::index_t> input_left_pads{in_left_pad_h, in_left_pad_w};
const std::vector<ck::index_t> input_right_pads{in_right_pad_h, in_right_pad_w};
const auto output_spatial_lengths =
ck::tensor_operation::ConvolutionUtility::ComputeOutputSpatialLengths({Hi, Wi},
{Y, X},
conv_filter_strides,
conv_filter_dilations,
input_left_pads,
input_right_pads);
const ck::index_t Ho = (Hi + in_left_pad_h + in_right_pad_h - YEff) / conv_stride_h + 1;
const ck::index_t Wo = (Wi + in_left_pad_w + in_right_pad_w - XEff) / conv_stride_w + 1;
const std::vector<ck::index_t> conv_filter_strides{{conv_stride_h, conv_stride_w}};
const std::vector<ck::index_t> conv_filter_dilations{{conv_dilation_h, conv_dilation_w}};
const std::vector<ck::index_t> input_left_pads{{in_left_pad_h, in_left_pad_w}};
const std::vector<ck::index_t> input_right_pads{{in_right_pad_h, in_right_pad_w}};
const ck::index_t Ho = output_spatial_lengths[0];
const ck::index_t Wo = output_spatial_lengths[1];
// tensor layout
auto f_host_tensor_descriptor = [](std::size_t N_,
@@ -242,9 +247,9 @@ int main(int argc, char* argv[])
N,
K,
C,
std::vector<ck::index_t>{{Hi, Wi}},
std::vector<ck::index_t>{{Y, X}},
std::vector<ck::index_t>{{Ho, Wo}},
std::vector<ck::index_t>{Hi, Wi},
std::vector<ck::index_t>{Y, X},
std::vector<ck::index_t>{Ho, Wo},
conv_filter_strides,
conv_filter_dilations,
input_left_pads,

29
example/7_conv2d_fwd_xdl_bias_relu_atomic_add/conv2d_fwd_xdl_bias_relu_atomic_add.cpp
View File

@@ -13,6 +13,7 @@
#include "tensor_layout.hpp"
#include "device_conv2d_fwd_xdl_c_shuffle_bias_activation_nhwc_kyxc_nhwk.hpp"
#include "element_wise_operation.hpp"
#include "convolution_utility.hpp"
using InDataType = ck::half_t;
using WeiDataType = ck::half_t;
@@ -166,16 +167,20 @@ int main(int argc, char* argv[])
exit(0);
}
const ck::index_t YEff = (Y - 1) * conv_dilation_h + 1;
const ck::index_t XEff = (X - 1) * conv_dilation_w + 1;
const std::vector<ck::index_t> conv_filter_strides{conv_stride_h, conv_stride_w};
const std::vector<ck::index_t> conv_filter_dilations{conv_dilation_h, conv_dilation_w};
const std::vector<ck::index_t> input_left_pads{in_left_pad_h, in_left_pad_w};
const std::vector<ck::index_t> input_right_pads{in_right_pad_h, in_right_pad_w};
const auto output_spatial_lengths =
ck::tensor_operation::ConvolutionUtility::ComputeOutputSpatialLengths({Hi, Wi},
{Y, X},
conv_filter_strides,
conv_filter_dilations,
input_left_pads,
input_right_pads);
const ck::index_t Ho = (Hi + in_left_pad_h + in_right_pad_h - YEff) / conv_stride_h + 1;
const ck::index_t Wo = (Wi + in_left_pad_w + in_right_pad_w - XEff) / conv_stride_w + 1;
const std::vector<ck::index_t> conv_filter_strides{{conv_stride_h, conv_stride_w}};
const std::vector<ck::index_t> conv_filter_dilations{{conv_dilation_h, conv_dilation_w}};
const std::vector<ck::index_t> input_left_pads{{in_left_pad_h, in_left_pad_w}};
const std::vector<ck::index_t> input_right_pads{{in_right_pad_h, in_right_pad_w}};
const ck::index_t Ho = output_spatial_lengths[0];
const ck::index_t Wo = output_spatial_lengths[1];
// tensor layout
auto f_host_tensor_descriptor = [](std::size_t N_,
@@ -255,9 +260,9 @@ int main(int argc, char* argv[])
N,
K,
C,
std::vector<ck::index_t>{{Hi, Wi}},
std::vector<ck::index_t>{{Y, X}},
std::vector<ck::index_t>{{Ho, Wo}},
std::vector<ck::index_t>{Hi, Wi},
std::vector<ck::index_t>{Y, X},
std::vector<ck::index_t>{Ho, Wo},
conv_filter_strides,
conv_filter_dilations,
input_left_pads,

29
example/9_conv2d_fwd_xdl_int8/conv2d_fwd_xdl_int8.cpp
View File

@@ -14,6 +14,7 @@
#include "device_conv2d_fwd_xdl_nhwc_kyxc_nhwk.hpp"
#include "element_wise_operation.hpp"
#include "reference_conv_fwd.hpp"
#include "convolution_utility.hpp"
using InDataType = int8_t;
using WeiDataType = int8_t;
@@ -136,16 +137,20 @@ int main(int argc, char* argv[])
exit(0);
}
const ck::index_t YEff = (Y - 1) * conv_dilation_h + 1;
const ck::index_t XEff = (X - 1) * conv_dilation_w + 1;
const std::vector<ck::index_t> conv_filter_strides{conv_stride_h, conv_stride_w};
const std::vector<ck::index_t> conv_filter_dilations{conv_dilation_h, conv_dilation_w};
const std::vector<ck::index_t> input_left_pads{in_left_pad_h, in_left_pad_w};
const std::vector<ck::index_t> input_right_pads{in_right_pad_h, in_right_pad_w};
const auto output_spatial_lengths =
ck::tensor_operation::ConvolutionUtility::ComputeOutputSpatialLengths({Hi, Wi},
{Y, X},
conv_filter_strides,
conv_filter_dilations,
input_left_pads,
input_right_pads);
const ck::index_t Ho = (Hi + in_left_pad_h + in_right_pad_h - YEff) / conv_stride_h + 1;
const ck::index_t Wo = (Wi + in_left_pad_w + in_right_pad_w - XEff) / conv_stride_w + 1;
const std::vector<ck::index_t> conv_filter_strides{{conv_stride_h, conv_stride_w}};
const std::vector<ck::index_t> conv_filter_dilations{{conv_dilation_h, conv_dilation_w}};
const std::vector<ck::index_t> input_left_pads{{in_left_pad_h, in_left_pad_w}};
const std::vector<ck::index_t> input_right_pads{{in_right_pad_h, in_right_pad_w}};
const ck::index_t Ho = output_spatial_lengths[0];
const ck::index_t Wo = output_spatial_lengths[1];
// tensor layout
auto f_host_tensor_descriptor = [](std::size_t N_,
@@ -212,9 +217,9 @@ int main(int argc, char* argv[])
N,
K,
C,
std::vector<ck::index_t>{{Hi, Wi}},
std::vector<ck::index_t>{{Y, X}},
std::vector<ck::index_t>{{Ho, Wo}},
std::vector<ck::index_t>{Hi, Wi},
std::vector<ck::index_t>{Y, X},
std::vector<ck::index_t>{Ho, Wo},
conv_filter_strides,
conv_filter_dilations,
input_left_pads,

5
example/CMakeLists.txt
View File

@@ -10,6 +10,7 @@ include_directories(BEFORE
${PROJECT_SOURCE_DIR}/composable_kernel/include/tensor_operation
${PROJECT_SOURCE_DIR}/composable_kernel/include/problem_transform
${PROJECT_SOURCE_DIR}/external/rocm/include
${PROJECT_SOURCE_DIR}/device_operation_reference/include
)
set(GEMM_XDL_SOURCE 1_gemm_xdl/gemm_xdl.cpp)
@@ -21,6 +22,7 @@ set(CONV2D_FWD_XDL_BIAS_RELU_ADD_SOURCE 6_conv2d_fwd_xdl_bias_relu_add/conv2d_fw
set(CONV2D_FWD_XDL_BIAS_RELU_ATOMIC_ADD_SOURCE 7_conv2d_fwd_xdl_bias_relu_atomic_add/conv2d_fwd_xdl_bias_relu_atomic_add.cpp)
set(GEMM_XDL_ALPHA_BETA_SOURCE 8_gemm_xdl_alpha_beta/gemm_xdl_alpha_beta.cpp)
set(CONV2D_FWD_XDL_INT8_SOURCE 9_conv2d_fwd_xdl_int8/conv2d_fwd_xdl_int8.cpp)
set(CONV3D_FWD_XDL_SOURCE 10_conv3d_fwd_xdl/conv3d_fwd_xdl.cpp)
add_executable(gemm_xdl ${GEMM_XDL_SOURCE})
add_executable(gemm_xdl_bias_relu ${GEMM_XDL_BIAS_RELU_SOURCE})
@@ -31,6 +33,7 @@ add_executable(conv2d_fwd_xdl_bias_relu_add ${CONV2D_FWD_XDL_BIAS_RELU_ADD_SOURC
add_executable(conv2d_fwd_xdl_bias_relu_atomic_add ${CONV2D_FWD_XDL_BIAS_RELU_ATOMIC_ADD_SOURCE})
add_executable(gemm_xdl_alpha_beta ${GEMM_XDL_ALPHA_BETA_SOURCE})
add_executable(conv2d_fwd_xdl_int8 ${CONV2D_FWD_XDL_INT8_SOURCE})
add_executable(conv3d_fwd_xdl ${CONV3D_FWD_XDL_SOURCE})
target_link_libraries(gemm_xdl PRIVATE host_tensor)
target_link_libraries(gemm_xdl_bias_relu PRIVATE host_tensor)
@@ -41,3 +44,5 @@ target_link_libraries(conv2d_fwd_xdl_bias_relu_add PRIVATE host_tensor)
target_link_libraries(conv2d_fwd_xdl_bias_relu_atomic_add PRIVATE host_tensor)
target_link_libraries(gemm_xdl_alpha_beta PRIVATE host_tensor)
target_link_libraries(conv2d_fwd_xdl_int8 PRIVATE host_tensor)
target_link_libraries(conv3d_fwd_xdl PRIVATE host_tensor)