diff --git a/driver/device_implicit_gemm_convolution_1_chwn_csrk_khwn.hpp b/driver/device_implicit_gemm_convolution_1_chwn_cyxk_khwn.hpp similarity index 91% rename from driver/device_implicit_gemm_convolution_1_chwn_csrk_khwn.hpp rename to driver/device_implicit_gemm_convolution_1_chwn_cyxk_khwn.hpp index fc2b245148..2c27080670 100644 --- a/driver/device_implicit_gemm_convolution_1_chwn_csrk_khwn.hpp +++ b/driver/device_implicit_gemm_convolution_1_chwn_cyxk_khwn.hpp @@ -1,13 +1,13 @@ #pragma once #include #include "device.hpp" -#include "gridwise_implicit_gemm_convolution_1_chwn_csrk_khwn.hip.hpp" +#include "gridwise_implicit_gemm_convolution_1_chwn_cyxk_khwn.hip.hpp" template -void device_implicit_gemm_convolution_1_chwn_csrk_khwn(InDesc, +void device_implicit_gemm_convolution_1_chwn_cyxk_khwn(InDesc, const Tensor& in_nchw, WeiDesc, - const Tensor& wei_kcsr, + const Tensor& wei_kcyx, OutDesc, Tensor& out_nkhw, unsigned nrepeat) @@ -18,7 +18,7 @@ void device_implicit_gemm_convolution_1_chwn_csrk_khwn(InDesc, constexpr auto I3 = Number<3>{}; constexpr auto in_nchw_desc = InDesc{}; - constexpr auto wei_kcsr_desc = WeiDesc{}; + constexpr auto wei_kcyx_desc = WeiDesc{}; constexpr auto out_nkhw_desc = OutDesc{}; constexpr unsigned Hi = in_nchw_desc.GetLength(I2); @@ -28,22 +28,22 @@ void device_implicit_gemm_convolution_1_chwn_csrk_khwn(InDesc, constexpr unsigned Ho = out_nkhw_desc.GetLength(I2); constexpr unsigned Wo = out_nkhw_desc.GetLength(I3); - constexpr unsigned K = wei_kcsr_desc.GetLength(I0); - constexpr unsigned C = wei_kcsr_desc.GetLength(I1); - constexpr unsigned Y = wei_kcsr_desc.GetLength(I2); - constexpr unsigned X = wei_kcsr_desc.GetLength(I3); + constexpr unsigned K = wei_kcyx_desc.GetLength(I0); + constexpr unsigned C = wei_kcyx_desc.GetLength(I1); + constexpr unsigned Y = wei_kcyx_desc.GetLength(I2); + constexpr unsigned X = wei_kcyx_desc.GetLength(I3); // reorder weight - auto wei_csrk_desc = make_ConstantTensorDescriptor(Sequence{}); - ostream_ConstantTensorDescriptor(wei_csrk_desc, std::cout << "wei_csrk_desc: "); + auto wei_cyxk_desc = make_ConstantTensorDescriptor(Sequence{}); + ostream_ConstantTensorDescriptor(wei_cyxk_desc, std::cout << "wei_cyxk_desc: "); - Tensor wei_csrk(make_TensorDescriptor(wei_csrk_desc)); + Tensor wei_cyxk(make_TensorDescriptor(wei_cyxk_desc)); - auto f_reorder_kcsr2csrk = [&](auto k, auto c, auto s, auto r) { - wei_csrk(c, s, r, k) = wei_kcsr(k, c, s, r); + auto f_reorder_kcyx2cyxk = [&](auto k, auto c, auto y, auto x) { + wei_cyxk(c, y, x, k) = wei_kcyx(k, c, y, x); }; - make_ParallelTensorFunctor(f_reorder_kcsr2csrk, K, C, Y, X)( + make_ParallelTensorFunctor(f_reorder_kcyx2cyxk, K, C, Y, X)( std::thread::hardware_concurrency()); // reorder input @@ -67,11 +67,11 @@ void device_implicit_gemm_convolution_1_chwn_csrk_khwn(InDesc, std::size_t data_sz = sizeof(T); DeviceMem in_chwn_device_buf(data_sz * in_chwn.mDesc.GetElementSpace()); - DeviceMem wei_csrk_device_buf(data_sz * wei_csrk.mDesc.GetElementSpace()); + DeviceMem wei_cyxk_device_buf(data_sz * wei_cyxk.mDesc.GetElementSpace()); DeviceMem out_khwn_device_buf(data_sz * out_khwn.mDesc.GetElementSpace()); in_chwn_device_buf.ToDevice(in_chwn.mData.data()); - wei_csrk_device_buf.ToDevice(wei_csrk.mData.data()); + wei_cyxk_device_buf.ToDevice(wei_cyxk.mData.data()); out_khwn_device_buf.ToDevice(out_khwn.mData.data()); #if 1 @@ -257,11 +257,11 @@ void device_implicit_gemm_convolution_1_chwn_csrk_khwn(InDesc, for(unsigned i = 0; i < nrepeat; ++i) { float time = launch_kernel( - gridwise_implicit_gemm_convolution_1_chwn_csrk_khwn(in_chwn_device_buf.GetDeviceBuffer()), - static_cast(wei_csrk_device_buf.GetDeviceBuffer()), + static_cast(wei_cyxk_device_buf.GetDeviceBuffer()), static_cast(out_khwn_device_buf.GetDeviceBuffer())); printf("Elapsed time : %f ms\n", time); diff --git a/driver/device_implicit_gemm_convolution_1_chwn_csrk_khwn_padded.hpp b/driver/device_implicit_gemm_convolution_1_chwn_cyxk_khwn_padded.hpp similarity index 90% rename from driver/device_implicit_gemm_convolution_1_chwn_csrk_khwn_padded.hpp rename to driver/device_implicit_gemm_convolution_1_chwn_cyxk_khwn_padded.hpp index db0cb3aa90..1843061a7a 100644 --- a/driver/device_implicit_gemm_convolution_1_chwn_csrk_khwn_padded.hpp +++ b/driver/device_implicit_gemm_convolution_1_chwn_cyxk_khwn_padded.hpp @@ -1,13 +1,13 @@ #pragma once #include #include "device.hpp" -#include "gridwise_implicit_gemm_convolution_1_chwn_csrk_khwn_padded.hip.hpp" +#include "gridwise_implicit_gemm_convolution_1_chwn_cyxk_khwn_padded.hip.hpp" template -void device_implicit_gemm_convolution_1_chwn_csrk_khwn_padded(InDesc, +void device_implicit_gemm_convolution_1_chwn_cyxk_khwn_padded(InDesc, const Tensor& in_nchw, WeiDesc, - const Tensor& wei_kcsr, + const Tensor& wei_kcyx, OutDesc, Tensor& out_nkhw, LowerPads, @@ -20,7 +20,7 @@ void device_implicit_gemm_convolution_1_chwn_csrk_khwn_padded(InDesc, constexpr auto I3 = Number<3>{}; constexpr auto in_nchw_desc = InDesc{}; - constexpr auto wei_kcsr_desc = WeiDesc{}; + constexpr auto wei_kcyx_desc = WeiDesc{}; constexpr auto out_nkhw_desc = OutDesc{}; constexpr unsigned Hi = in_nchw_desc.GetLength(I2); @@ -30,22 +30,22 @@ void device_implicit_gemm_convolution_1_chwn_csrk_khwn_padded(InDesc, constexpr unsigned Ho = out_nkhw_desc.GetLength(I2); constexpr unsigned Wo = out_nkhw_desc.GetLength(I3); - constexpr unsigned K = wei_kcsr_desc.GetLength(I0); - constexpr unsigned C = wei_kcsr_desc.GetLength(I1); - constexpr unsigned Y = wei_kcsr_desc.GetLength(I2); - constexpr unsigned X = wei_kcsr_desc.GetLength(I3); + constexpr unsigned K = wei_kcyx_desc.GetLength(I0); + constexpr unsigned C = wei_kcyx_desc.GetLength(I1); + constexpr unsigned Y = wei_kcyx_desc.GetLength(I2); + constexpr unsigned X = wei_kcyx_desc.GetLength(I3); // reorder weight - auto wei_csrk_desc = make_ConstantTensorDescriptor(Sequence{}); - ostream_ConstantTensorDescriptor(wei_csrk_desc, std::cout << "wei_csrk_desc: "); + auto wei_cyxk_desc = make_ConstantTensorDescriptor(Sequence{}); + ostream_ConstantTensorDescriptor(wei_cyxk_desc, std::cout << "wei_cyxk_desc: "); - Tensor wei_csrk(make_TensorDescriptor(wei_csrk_desc)); + Tensor wei_cyxk(make_TensorDescriptor(wei_cyxk_desc)); - auto f_reorder_kcsr2csrk = [&](auto k, auto c, auto s, auto r) { - wei_csrk(c, s, r, k) = wei_kcsr(k, c, s, r); + auto f_reorder_kcyx2cyxk = [&](auto k, auto c, auto y, auto x) { + wei_cyxk(c, y, x, k) = wei_kcyx(k, c, y, x); }; - make_ParallelTensorFunctor(f_reorder_kcsr2csrk, K, C, Y, X)( + make_ParallelTensorFunctor(f_reorder_kcyx2cyxk, K, C, Y, X)( std::thread::hardware_concurrency()); // reorder input @@ -69,11 +69,11 @@ void device_implicit_gemm_convolution_1_chwn_csrk_khwn_padded(InDesc, std::size_t data_sz = sizeof(T); DeviceMem in_chwn_device_buf(data_sz * in_chwn.mDesc.GetElementSpace()); - DeviceMem wei_csrk_device_buf(data_sz * wei_csrk.mDesc.GetElementSpace()); + DeviceMem wei_cyxk_device_buf(data_sz * wei_cyxk.mDesc.GetElementSpace()); DeviceMem out_khwn_device_buf(data_sz * out_khwn.mDesc.GetElementSpace()); in_chwn_device_buf.ToDevice(in_chwn.mData.data()); - wei_csrk_device_buf.ToDevice(wei_csrk.mData.data()); + wei_cyxk_device_buf.ToDevice(wei_cyxk.mData.data()); out_khwn_device_buf.ToDevice(out_khwn.mData.data()); #if 0 @@ -250,11 +250,11 @@ void device_implicit_gemm_convolution_1_chwn_csrk_khwn_padded(InDesc, for(unsigned i = 0; i < nrepeat; ++i) { float time = launch_kernel( - gridwise_implicit_gemm_convolution_1_chwn_csrk_khwn_padded(in_chwn_device_buf.GetDeviceBuffer()), - static_cast(wei_csrk_device_buf.GetDeviceBuffer()), + static_cast(wei_cyxk_device_buf.GetDeviceBuffer()), static_cast(out_khwn_device_buf.GetDeviceBuffer())); printf("Elapsed time : %f ms\n", time); diff --git a/driver/device_implicit_gemm_convolution_2_chwn_csrk_khwn.hpp b/driver/device_implicit_gemm_convolution_2_chwn_cyxk_khwn.hpp similarity index 89% rename from driver/device_implicit_gemm_convolution_2_chwn_csrk_khwn.hpp rename to driver/device_implicit_gemm_convolution_2_chwn_cyxk_khwn.hpp index 88fd5a2dea..a657949f35 100644 --- a/driver/device_implicit_gemm_convolution_2_chwn_csrk_khwn.hpp +++ b/driver/device_implicit_gemm_convolution_2_chwn_cyxk_khwn.hpp @@ -1,13 +1,13 @@ #pragma once #include #include "device.hpp" -#include "gridwise_implicit_gemm_convolution_2_chwn_csrk_khwn_lds_double_buffer.hip.hpp" +#include "gridwise_implicit_gemm_convolution_2_chwn_cyxk_khwn_lds_double_buffer.hip.hpp" template -void device_implicit_gemm_convolution_2_chwn_csrk_khwn(InDesc, +void device_implicit_gemm_convolution_2_chwn_cyxk_khwn(InDesc, const Tensor& in_nchw, WeiDesc, - const Tensor& wei_kcsr, + const Tensor& wei_kcyx, OutDesc, Tensor& out_nkhw, unsigned nrepeat) @@ -18,7 +18,7 @@ void device_implicit_gemm_convolution_2_chwn_csrk_khwn(InDesc, constexpr auto I3 = Number<3>{}; constexpr auto in_nchw_desc = InDesc{}; - constexpr auto wei_kcsr_desc = WeiDesc{}; + constexpr auto wei_kcyx_desc = WeiDesc{}; constexpr auto out_nkhw_desc = OutDesc{}; constexpr unsigned N = in_nchw_desc.GetLength(I0); @@ -28,10 +28,10 @@ void device_implicit_gemm_convolution_2_chwn_csrk_khwn(InDesc, constexpr unsigned Ho = out_nkhw_desc.GetLength(I2); constexpr unsigned Wo = out_nkhw_desc.GetLength(I3); - constexpr unsigned K = wei_kcsr_desc.GetLength(I0); - constexpr unsigned C = wei_kcsr_desc.GetLength(I1); - constexpr unsigned Y = wei_kcsr_desc.GetLength(I2); - constexpr unsigned X = wei_kcsr_desc.GetLength(I3); + constexpr unsigned K = wei_kcyx_desc.GetLength(I0); + constexpr unsigned C = wei_kcyx_desc.GetLength(I1); + constexpr unsigned Y = wei_kcyx_desc.GetLength(I2); + constexpr unsigned X = wei_kcyx_desc.GetLength(I3); constexpr unsigned BGhostRead = (Y - 1) * Wi + (X - 1); @@ -48,14 +48,14 @@ void device_implicit_gemm_convolution_2_chwn_csrk_khwn(InDesc, Hi, Wi)(std::thread::hardware_concurrency()); - // convert wei_kcsr to wei_csrk - auto wei_csrk_desc = make_ConstantTensorDescriptor(Sequence{}); - ostream_ConstantTensorDescriptor(wei_csrk_desc, std::cout << "wei_csrk_desc: "); + // convert wei_kcyx to wei_cyxk + auto wei_cyxk_desc = make_ConstantTensorDescriptor(Sequence{}); + ostream_ConstantTensorDescriptor(wei_cyxk_desc, std::cout << "wei_cyxk_desc: "); - Tensor wei_csrk(make_TensorDescriptor(wei_csrk_desc)); + Tensor wei_cyxk(make_TensorDescriptor(wei_cyxk_desc)); make_ParallelTensorFunctor( - [&](auto k, auto c, auto s, auto r) { wei_csrk(c, s, r, k) = wei_kcsr(k, c, s, r); }, + [&](auto k, auto c, auto y, auto x) { wei_cyxk(c, y, x, k) = wei_kcyx(k, c, y, x); }, K, C, Y, @@ -200,22 +200,22 @@ void device_implicit_gemm_convolution_2_chwn_csrk_khwn(InDesc, std::size_t data_sz = sizeof(T); DeviceMem in_chwn_device_buf(data_sz * (in_chwn.mDesc.GetElementSpace() + BGhostRead + BPerBlock)); // reserve extra space for BGhostRead - DeviceMem wei_csrk_device_buf(data_sz * wei_csrk.mDesc.GetElementSpace()); + DeviceMem wei_cyxk_device_buf(data_sz * wei_cyxk.mDesc.GetElementSpace()); DeviceMem out_khwn_device_buf(data_sz * out_khwn.mDesc.GetElementSpace()); in_chwn_device_buf.ToDevice(in_chwn.mData.data()); - wei_csrk_device_buf.ToDevice(wei_csrk.mData.data()); + wei_cyxk_device_buf.ToDevice(wei_cyxk.mData.data()); out_khwn_device_buf.ToDevice(out_khwn.mData.data()); for(unsigned i = 0; i < nrepeat; ++i) { float time = - launch_kernel(gridwise_implicit_gemm_convolution_2_chwn_csrk_khwn_lds_double_buffer< + launch_kernel(gridwise_implicit_gemm_convolution_2_chwn_cyxk_khwn_lds_double_buffer< GridSize, BlockSize, T, decltype(in_chwn_desc), - decltype(wei_csrk_desc), + decltype(wei_cyxk_desc), decltype(out_khwn_desc), BPerBlock, KPerBlock, @@ -240,7 +240,7 @@ void device_implicit_gemm_convolution_2_chwn_csrk_khwn(InDesc, dim3(GridSize), dim3(BlockSize), static_cast(in_chwn_device_buf.GetDeviceBuffer()), - static_cast(wei_csrk_device_buf.GetDeviceBuffer()), + static_cast(wei_cyxk_device_buf.GetDeviceBuffer()), static_cast(out_khwn_device_buf.GetDeviceBuffer())); printf("Elapsed time : %f ms\n", time); diff --git a/driver/driver.hip.cpp b/driver/driver.hip.cpp index b5af7009ce..aca345acfd 100644 --- a/driver/driver.hip.cpp +++ b/driver/driver.hip.cpp @@ -9,9 +9,9 @@ #include "conv_common.hip.hpp" #include "device_direct_convolution_1.hpp" #include "device_direct_convolution_2.hpp" -#include "device_implicit_gemm_convolution_1_chwn_csrk_khwn.hpp" -#include "device_implicit_gemm_convolution_1_chwn_csrk_khwn_padded.hpp" -#include "device_implicit_gemm_convolution_2_chwn_csrk_khwn.hpp" +#include "device_implicit_gemm_convolution_1_chwn_cyxk_khwn.hpp" +#include "device_implicit_gemm_convolution_1_chwn_cyxk_khwn_padded.hpp" +#include "device_implicit_gemm_convolution_2_chwn_cyxk_khwn.hpp" struct GeneratorTensor_1 { @@ -108,7 +108,7 @@ auto make_TensorDescriptor(TConstTensorDesc) template void host_direct_convolution( - const Tensor& in_nchw, const Tensor& wei_kcsr, Tensor& out, LowerPads, UpperPads) + const Tensor& in_nchw, const Tensor& wei_kcyx, Tensor& out, LowerPads, UpperPads) { unsigned h_pad_low = LowerPads{}.Get(Number<0>{}); unsigned w_pad_low = LowerPads{}.Get(Number<1>{}); @@ -118,18 +118,18 @@ void host_direct_convolution( auto f = [&](auto n, auto k, auto ho, auto wo) { double v = 0; - for(int c = 0; c < wei_kcsr.mDesc.GetLengths()[1]; ++c) + for(int c = 0; c < wei_kcyx.mDesc.GetLengths()[1]; ++c) { - for(int y = 0; y < wei_kcsr.mDesc.GetLengths()[2]; ++y) + for(int y = 0; y < wei_kcyx.mDesc.GetLengths()[2]; ++y) { int hi = ho + y - h_pad_low; - for(int x = 0; x < wei_kcsr.mDesc.GetLengths()[3]; ++x) + for(int x = 0; x < wei_kcyx.mDesc.GetLengths()[3]; ++x) { int wi = wo + x - w_pad_low; if(hi >= 0 && hi < in_nchw.mDesc.GetLengths()[2] && wi >= 0 && wi < in_nchw.mDesc.GetLengths()[3]) { - v += in_nchw(n, c, hi, wi) * wei_kcsr(k, c, y, x); + v += in_nchw(n, c, hi, wi) * wei_kcyx(k, c, y, x); } } } @@ -148,7 +148,7 @@ void host_direct_convolution( template void host_winograd_3x3_convolution( - const Tensor& in_nchw, const Tensor& wei_kcsr, Tensor& out, LowerPads, UpperPads) + const Tensor& in_nchw, const Tensor& wei_kcyx, Tensor& out, LowerPads, UpperPads) { constexpr std::size_t HoPerTile = 2; constexpr std::size_t WoPerTile = 2; @@ -158,9 +158,9 @@ void host_winograd_3x3_convolution( std::size_t HI = in_nchw.mDesc.GetLengths()[2]; std::size_t WI = in_nchw.mDesc.GetLengths()[3]; - std::size_t K = wei_kcsr.mDesc.GetLengths()[0]; - std::size_t Y = wei_kcsr.mDesc.GetLengths()[2]; - std::size_t X = wei_kcsr.mDesc.GetLengths()[3]; + std::size_t K = wei_kcyx.mDesc.GetLengths()[0]; + std::size_t Y = wei_kcyx.mDesc.GetLengths()[2]; + std::size_t X = wei_kcyx.mDesc.GetLengths()[3]; std::size_t HO = out.mDesc.GetLengths()[2]; std::size_t WO = out.mDesc.GetLengths()[3]; @@ -259,49 +259,49 @@ void host_winograd_3x3_convolution( }; auto f_wei_transform = [&](auto k, auto c) { - wei_transform(k, c, 0, 0) = wei_kcsr(k, c, 0, 0); + wei_transform(k, c, 0, 0) = wei_kcyx(k, c, 0, 0); wei_transform(k, c, 0, 1) = - 0.5 * wei_kcsr(k, c, 0, 0) + 0.5 * wei_kcsr(k, c, 0, 1) + 0.5 * wei_kcsr(k, c, 0, 2); + 0.5 * wei_kcyx(k, c, 0, 0) + 0.5 * wei_kcyx(k, c, 0, 1) + 0.5 * wei_kcyx(k, c, 0, 2); wei_transform(k, c, 0, 2) = - 0.5 * wei_kcsr(k, c, 0, 0) - 0.5 * wei_kcsr(k, c, 0, 1) + 0.5 * wei_kcsr(k, c, 0, 2); - wei_transform(k, c, 0, 3) = wei_kcsr(k, c, 0, 2); + 0.5 * wei_kcyx(k, c, 0, 0) - 0.5 * wei_kcyx(k, c, 0, 1) + 0.5 * wei_kcyx(k, c, 0, 2); + wei_transform(k, c, 0, 3) = wei_kcyx(k, c, 0, 2); wei_transform(k, c, 1, 0) = - 0.5 * wei_kcsr(k, c, 0, 0) + 0.5 * wei_kcsr(k, c, 1, 0) + 0.5 * wei_kcsr(k, c, 2, 0); - wei_transform(k, c, 1, 1) = 0.25 * wei_kcsr(k, c, 0, 0) + 0.25 * wei_kcsr(k, c, 0, 1) + - 0.25 * wei_kcsr(k, c, 0, 2) + 0.25 * wei_kcsr(k, c, 1, 0) + - 0.25 * wei_kcsr(k, c, 1, 1) + 0.25 * wei_kcsr(k, c, 1, 2) + - 0.25 * wei_kcsr(k, c, 2, 0) + 0.25 * wei_kcsr(k, c, 2, 1) + - 0.25 * wei_kcsr(k, c, 2, 2); - wei_transform(k, c, 1, 2) = 0.25 * wei_kcsr(k, c, 0, 0) - 0.25 * wei_kcsr(k, c, 0, 1) + - 0.25 * wei_kcsr(k, c, 0, 2) + 0.25 * wei_kcsr(k, c, 1, 0) - - 0.25 * wei_kcsr(k, c, 1, 1) + 0.25 * wei_kcsr(k, c, 1, 2) + - 0.25 * wei_kcsr(k, c, 2, 0) - 0.25 * wei_kcsr(k, c, 2, 1) + - 0.25 * wei_kcsr(k, c, 2, 2); + 0.5 * wei_kcyx(k, c, 0, 0) + 0.5 * wei_kcyx(k, c, 1, 0) + 0.5 * wei_kcyx(k, c, 2, 0); + wei_transform(k, c, 1, 1) = 0.25 * wei_kcyx(k, c, 0, 0) + 0.25 * wei_kcyx(k, c, 0, 1) + + 0.25 * wei_kcyx(k, c, 0, 2) + 0.25 * wei_kcyx(k, c, 1, 0) + + 0.25 * wei_kcyx(k, c, 1, 1) + 0.25 * wei_kcyx(k, c, 1, 2) + + 0.25 * wei_kcyx(k, c, 2, 0) + 0.25 * wei_kcyx(k, c, 2, 1) + + 0.25 * wei_kcyx(k, c, 2, 2); + wei_transform(k, c, 1, 2) = 0.25 * wei_kcyx(k, c, 0, 0) - 0.25 * wei_kcyx(k, c, 0, 1) + + 0.25 * wei_kcyx(k, c, 0, 2) + 0.25 * wei_kcyx(k, c, 1, 0) - + 0.25 * wei_kcyx(k, c, 1, 1) + 0.25 * wei_kcyx(k, c, 1, 2) + + 0.25 * wei_kcyx(k, c, 2, 0) - 0.25 * wei_kcyx(k, c, 2, 1) + + 0.25 * wei_kcyx(k, c, 2, 2); wei_transform(k, c, 1, 3) = - 0.5 * wei_kcsr(k, c, 0, 2) + 0.5 * wei_kcsr(k, c, 1, 2) + 0.5 * wei_kcsr(k, c, 2, 2); + 0.5 * wei_kcyx(k, c, 0, 2) + 0.5 * wei_kcyx(k, c, 1, 2) + 0.5 * wei_kcyx(k, c, 2, 2); wei_transform(k, c, 2, 0) = - 0.5 * wei_kcsr(k, c, 0, 0) - 0.5 * wei_kcsr(k, c, 1, 0) + 0.5 * wei_kcsr(k, c, 2, 0); - wei_transform(k, c, 2, 1) = 0.25 * wei_kcsr(k, c, 0, 0) + 0.25 * wei_kcsr(k, c, 0, 1) + - 0.25 * wei_kcsr(k, c, 0, 2) - 0.25 * wei_kcsr(k, c, 1, 0) - - 0.25 * wei_kcsr(k, c, 1, 1) - 0.25 * wei_kcsr(k, c, 1, 2) + - 0.25 * wei_kcsr(k, c, 2, 0) + 0.25 * wei_kcsr(k, c, 2, 1) + - 0.25 * wei_kcsr(k, c, 2, 2); - wei_transform(k, c, 2, 2) = 0.25 * wei_kcsr(k, c, 0, 0) - 0.25 * wei_kcsr(k, c, 0, 1) + - 0.25 * wei_kcsr(k, c, 0, 2) - 0.25 * wei_kcsr(k, c, 1, 0) + - 0.25 * wei_kcsr(k, c, 1, 1) - 0.25 * wei_kcsr(k, c, 1, 2) + - 0.25 * wei_kcsr(k, c, 2, 0) - 0.25 * wei_kcsr(k, c, 2, 1) + - 0.25 * wei_kcsr(k, c, 2, 2); + 0.5 * wei_kcyx(k, c, 0, 0) - 0.5 * wei_kcyx(k, c, 1, 0) + 0.5 * wei_kcyx(k, c, 2, 0); + wei_transform(k, c, 2, 1) = 0.25 * wei_kcyx(k, c, 0, 0) + 0.25 * wei_kcyx(k, c, 0, 1) + + 0.25 * wei_kcyx(k, c, 0, 2) - 0.25 * wei_kcyx(k, c, 1, 0) - + 0.25 * wei_kcyx(k, c, 1, 1) - 0.25 * wei_kcyx(k, c, 1, 2) + + 0.25 * wei_kcyx(k, c, 2, 0) + 0.25 * wei_kcyx(k, c, 2, 1) + + 0.25 * wei_kcyx(k, c, 2, 2); + wei_transform(k, c, 2, 2) = 0.25 * wei_kcyx(k, c, 0, 0) - 0.25 * wei_kcyx(k, c, 0, 1) + + 0.25 * wei_kcyx(k, c, 0, 2) - 0.25 * wei_kcyx(k, c, 1, 0) + + 0.25 * wei_kcyx(k, c, 1, 1) - 0.25 * wei_kcyx(k, c, 1, 2) + + 0.25 * wei_kcyx(k, c, 2, 0) - 0.25 * wei_kcyx(k, c, 2, 1) + + 0.25 * wei_kcyx(k, c, 2, 2); wei_transform(k, c, 2, 3) = - 0.5 * wei_kcsr(k, c, 0, 2) - 0.5 * wei_kcsr(k, c, 1, 2) + 0.5 * wei_kcsr(k, c, 2, 2); + 0.5 * wei_kcyx(k, c, 0, 2) - 0.5 * wei_kcyx(k, c, 1, 2) + 0.5 * wei_kcyx(k, c, 2, 2); - wei_transform(k, c, 3, 0) = wei_kcsr(k, c, 2, 0); + wei_transform(k, c, 3, 0) = wei_kcyx(k, c, 2, 0); wei_transform(k, c, 3, 1) = - 0.5 * wei_kcsr(k, c, 2, 0) + 0.5 * wei_kcsr(k, c, 2, 1) + 0.5 * wei_kcsr(k, c, 2, 2); + 0.5 * wei_kcyx(k, c, 2, 0) + 0.5 * wei_kcyx(k, c, 2, 1) + 0.5 * wei_kcyx(k, c, 2, 2); wei_transform(k, c, 3, 2) = - 0.5 * wei_kcsr(k, c, 2, 0) - 0.5 * wei_kcsr(k, c, 2, 1) + 0.5 * wei_kcsr(k, c, 2, 2); - wei_transform(k, c, 3, 3) = wei_kcsr(k, c, 2, 2); + 0.5 * wei_kcyx(k, c, 2, 0) - 0.5 * wei_kcyx(k, c, 2, 1) + 0.5 * wei_kcyx(k, c, 2, 2); + wei_transform(k, c, 3, 3) = wei_kcyx(k, c, 2, 2); }; auto f_out_transform = [&](auto n, auto k, auto htile, auto wtile) { @@ -569,16 +569,16 @@ int main(int argc, char* argv[]) auto upper_pads = Sequence{}; auto in_nchw_desc = make_ConstantTensorDescriptor(Sequence{}); - auto wei_kcsr_desc = make_ConstantTensorDescriptor(Sequence{}); + auto wei_kcyx_desc = make_ConstantTensorDescriptor(Sequence{}); auto out_nkhw_desc = get_convolution_with_padding_output_default_4d_tensor_descriptor( - in_nchw_desc, wei_kcsr_desc, lower_pads, upper_pads); + in_nchw_desc, wei_kcyx_desc, lower_pads, upper_pads); ostream_ConstantTensorDescriptor(in_nchw_desc, std::cout << "in_nchw_desc: "); - ostream_ConstantTensorDescriptor(wei_kcsr_desc, std::cout << "wei_kcsr_desc: "); + ostream_ConstantTensorDescriptor(wei_kcyx_desc, std::cout << "wei_kcyx_desc: "); ostream_ConstantTensorDescriptor(out_nkhw_desc, std::cout << "out_nkhw_desc: "); Tensor in_nchw(make_TensorDescriptor(in_nchw_desc)); - Tensor wei_kcsr(make_TensorDescriptor(wei_kcsr_desc)); + Tensor wei_kcyx(make_TensorDescriptor(wei_kcyx_desc)); Tensor out_nkhw_host(make_TensorDescriptor(out_nkhw_desc)); Tensor out_nkhw_device(make_TensorDescriptor(out_nkhw_desc)); @@ -597,13 +597,13 @@ int main(int argc, char* argv[]) { #if 0 in_nchw.GenerateTensorValue(GeneratorTensor_1{}, num_thread); - wei_kcsr.GenerateTensorValue(GeneratorTensor_1{}, num_thread); + wei_kcyx.GenerateTensorValue(GeneratorTensor_1{}, num_thread); #elif 1 in_nchw.GenerateTensorValue(GeneratorTensor_2{-5, 5}, num_thread); - wei_kcsr.GenerateTensorValue(GeneratorTensor_2{-5, 5}, num_thread); + wei_kcyx.GenerateTensorValue(GeneratorTensor_2{-5, 5}, num_thread); #elif 1 in_nchw.GenerateTensorValue(GeneratorTensor_2{-2, 2}, num_thread); - wei_kcsr.GenerateTensorValue(GeneratorTensor_1{}, num_thread); + wei_kcyx.GenerateTensorValue(GeneratorTensor_1{}, num_thread); #endif } @@ -613,17 +613,17 @@ int main(int argc, char* argv[]) #elif 0 device_direct_convolution_2 #elif 1 - device_implicit_gemm_convolution_1_chwn_csrk_khwn + device_implicit_gemm_convolution_1_chwn_cyxk_khwn #elif 0 - device_implicit_gemm_convolution_2_chwn_csrk_khwn + device_implicit_gemm_convolution_2_chwn_cyxk_khwn #endif - (in_nchw_desc, in_nchw, wei_kcsr_desc, wei_kcsr, out_nkhw_desc, out_nkhw_device, nrepeat); + (in_nchw_desc, in_nchw, wei_kcyx_desc, wei_kcyx, out_nkhw_desc, out_nkhw_device, nrepeat); #elif 1 - device_implicit_gemm_convolution_1_chwn_csrk_khwn_padded(in_nchw_desc, + device_implicit_gemm_convolution_1_chwn_cyxk_khwn_padded(in_nchw_desc, in_nchw, - wei_kcsr_desc, - wei_kcsr, + wei_kcyx_desc, + wei_kcyx, out_nkhw_desc, out_nkhw_device, lower_pads, @@ -636,18 +636,18 @@ int main(int argc, char* argv[]) #if 0 if(Y == 3 && X == 3) { - host_winograd_3x3_convolution(in_nchw, wei_kcsr, out_nkhw_host, lower_pads, upper_pads); + host_winograd_3x3_convolution(in_nchw, wei_kcyx, out_nkhw_host, lower_pads, upper_pads); } else { - host_direct_convolution(in_nchw, wei_kcsr, out_nkhw_host, lower_pads, upper_pads); + host_direct_convolution(in_nchw, wei_kcyx, out_nkhw_host, lower_pads, upper_pads); } check_error(out_nkhw_host, out_nkhw_device); #endif #if 0 LogRange(std::cout << "in_nchw : ", in_nchw.mData, ",") << std::endl; - LogRange(std::cout << "wei_kcsr: ", wei_kcsr.mData, ",") << std::endl; + LogRange(std::cout << "wei_kcyx: ", wei_kcyx.mData, ",") << std::endl; LogRange(std::cout << "out_nkhw_host : ", out_nkhw_host.mData, ",") << std::endl; LogRange(std::cout << "out_nkhw_device: ", out_nkhw_device.mData, ",") << std::endl; #endif diff --git a/src/include/gridwise_implicit_gemm_convolution_1_chwn_csrk_khwn.hip.hpp b/src/include/gridwise_implicit_gemm_convolution_1_chwn_cyxk_khwn.hip.hpp similarity index 94% rename from src/include/gridwise_implicit_gemm_convolution_1_chwn_csrk_khwn.hip.hpp rename to src/include/gridwise_implicit_gemm_convolution_1_chwn_cyxk_khwn.hip.hpp index b5e810cfd5..1caef669e9 100644 --- a/src/include/gridwise_implicit_gemm_convolution_1_chwn_csrk_khwn.hip.hpp +++ b/src/include/gridwise_implicit_gemm_convolution_1_chwn_cyxk_khwn.hip.hpp @@ -35,7 +35,7 @@ template __global__ void -gridwise_implicit_gemm_convolution_1_chwn_csrk_khwn(const Float* const __restrict__ p_in_global, +gridwise_implicit_gemm_convolution_1_chwn_cyxk_khwn(const Float* const __restrict__ p_in_global, const Float* const __restrict__ p_wei_global, Float* const __restrict__ p_out_global) { @@ -52,7 +52,7 @@ gridwise_implicit_gemm_convolution_1_chwn_csrk_khwn(const Float* const __restric constexpr auto I3 = Number<3>{}; constexpr auto in_chwn_global_desc = InGlobalDesc{}; - constexpr auto wei_csrk_global_desc = WeiGlobalDesc{}; + constexpr auto wei_cyxk_global_desc = WeiGlobalDesc{}; constexpr auto out_khwn_global_desc = OutGlobalDesc{}; constexpr unsigned C = in_chwn_global_desc.GetLength(I0); @@ -62,8 +62,8 @@ gridwise_implicit_gemm_convolution_1_chwn_csrk_khwn(const Float* const __restric constexpr unsigned Wo = out_khwn_global_desc.GetLength(I2); constexpr unsigned N = out_khwn_global_desc.GetLength(I3); - constexpr unsigned Y = wei_csrk_global_desc.GetLength(I1); - constexpr unsigned X = wei_csrk_global_desc.GetLength(I2); + constexpr unsigned Y = wei_cyxk_global_desc.GetLength(I1); + constexpr unsigned X = wei_cyxk_global_desc.GetLength(I2); constexpr unsigned HiPerBlock = HoPerBlock + Y - 1; constexpr unsigned WiPerBlock = WoPerBlock + X - 1; @@ -100,7 +100,7 @@ gridwise_implicit_gemm_convolution_1_chwn_csrk_khwn(const Float* const __restric constexpr auto wei_ek_block_desc = make_ConstantTensorDescriptor_aligned( Sequence{}, Number{}); - constexpr auto wei_csrk_block_desc = make_ConstantTensorDescriptor_aligned( + constexpr auto wei_cyxk_block_desc = make_ConstantTensorDescriptor_aligned( Sequence{}, Number{}); // tensor view of threadwise output in register @@ -133,7 +133,7 @@ gridwise_implicit_gemm_convolution_1_chwn_csrk_khwn(const Float* const __restric // B_matrix[C,Wo*N] is a sub-matrix of in_block[C,Hi,Wi,N] // C_matrix[K,Wo*N] is a sub-matrix of out_block[K,Ho,Wo,N] constexpr auto a_cxk_block_mtx_desc = make_ConstantMatrixDescriptor( - Number{}, Number{}, Number{}); + Number{}, Number{}, Number{}); constexpr auto b_cxwn_block_mtx_desc = make_ConstantMatrixDescriptor(Number{}, @@ -168,7 +168,7 @@ gridwise_implicit_gemm_convolution_1_chwn_csrk_khwn(const Float* const __restric in_chwn_block_desc.GetElementSpace(Number{}); constexpr unsigned wei_block_size = - wei_csrk_block_desc.GetElementSpace(Number{}); + wei_cyxk_block_desc.GetElementSpace(Number{}); constexpr unsigned max_align = InBlockCopyDataPerRead > WeiBlockCopyDataPerRead ? InBlockCopyDataPerRead @@ -188,11 +188,11 @@ gridwise_implicit_gemm_convolution_1_chwn_csrk_khwn(const Float* const __restric 0, hi_block_data_begin, wi_block_data_begin, n_block_data_begin); const Float* p_wei_global_block_begin = - p_wei_global + wei_csrk_global_desc.Get1dIndex(0, 0, 0, k_block_data_begin); + p_wei_global + wei_cyxk_global_desc.Get1dIndex(0, 0, 0, k_block_data_begin); for(unsigned c_block_data_begin = 0; c_block_data_begin < C; c_block_data_begin += CPerBlock, p_in_global_block_begin += CPerBlock * in_chwn_global_desc.GetStride(I0), - p_wei_global_block_begin += CPerBlock * wei_csrk_global_desc.GetStride(I0), + p_wei_global_block_begin += CPerBlock * wei_cyxk_global_desc.GetStride(I0), __syncthreads()) { // input: global mem to LDS @@ -204,12 +204,12 @@ gridwise_implicit_gemm_convolution_1_chwn_csrk_khwn(const Float* const __restric __syncthreads(); // a series of batched GEMM - for(unsigned s = 0; s < Y; ++s) + for(unsigned y = 0; y < Y; ++y) { - for(unsigned r = 0; r < X; ++r) + for(unsigned x = 0; x < X; ++x) { - blockwise_batch_gemm.Run(p_wei_block + wei_csrk_block_desc.Get1dIndex(0, s, r, 0), - p_in_block + in_chwn_block_desc.Get1dIndex(0, s, r, 0), + blockwise_batch_gemm.Run(p_wei_block + wei_cyxk_block_desc.Get1dIndex(0, y, x, 0), + p_in_block + in_chwn_block_desc.Get1dIndex(0, y, x, 0), p_out_thread, [](auto& acc, const auto&& v) { acc += v; }); } diff --git a/src/include/gridwise_implicit_gemm_convolution_1_chwn_csrk_khwn_padded.hip.hpp b/src/include/gridwise_implicit_gemm_convolution_1_chwn_cyxk_khwn_padded.hip.hpp similarity index 93% rename from src/include/gridwise_implicit_gemm_convolution_1_chwn_csrk_khwn_padded.hip.hpp rename to src/include/gridwise_implicit_gemm_convolution_1_chwn_cyxk_khwn_padded.hip.hpp index 5f14fc3373..a4904cdf58 100644 --- a/src/include/gridwise_implicit_gemm_convolution_1_chwn_csrk_khwn_padded.hip.hpp +++ b/src/include/gridwise_implicit_gemm_convolution_1_chwn_cyxk_khwn_padded.hip.hpp @@ -27,7 +27,7 @@ template -__global__ void gridwise_implicit_gemm_convolution_1_chwn_csrk_khwn_padded( +__global__ void gridwise_implicit_gemm_convolution_1_chwn_cyxk_khwn_padded( const Float* const __restrict__ p_in_global, const Float* const __restrict__ p_wei_global, Float* const __restrict__ p_out_global) @@ -45,7 +45,7 @@ __global__ void gridwise_implicit_gemm_convolution_1_chwn_csrk_khwn_padded( constexpr auto I3 = Number<3>{}; constexpr auto in_chwn_global_desc = InGlobalDesc{}; - constexpr auto wei_csrk_global_desc = WeiGlobalDesc{}; + constexpr auto wei_cyxk_global_desc = WeiGlobalDesc{}; constexpr auto out_khwn_global_desc = OutGlobalDesc{}; constexpr unsigned C = in_chwn_global_desc.GetLength(I0); @@ -55,8 +55,8 @@ __global__ void gridwise_implicit_gemm_convolution_1_chwn_csrk_khwn_padded( constexpr unsigned Wo = out_khwn_global_desc.GetLength(I2); constexpr unsigned N = out_khwn_global_desc.GetLength(I3); - constexpr unsigned Y = wei_csrk_global_desc.GetLength(I1); - constexpr unsigned X = wei_csrk_global_desc.GetLength(I2); + constexpr unsigned Y = wei_cyxk_global_desc.GetLength(I1); + constexpr unsigned X = wei_cyxk_global_desc.GetLength(I2); constexpr unsigned HPadLow = LowerPads{}.Get(I0); constexpr unsigned WPadLow = LowerPads{}.Get(I1); @@ -92,7 +92,7 @@ __global__ void gridwise_implicit_gemm_convolution_1_chwn_csrk_khwn_padded( constexpr auto in_chwn_block_desc = make_ConstantTensorDescriptor(Sequence{}); - constexpr auto wei_csrk_block_desc = + constexpr auto wei_cyxk_block_desc = make_ConstantTensorDescriptor(Sequence{}); // flattened (2d) tensor view of wei in LDS @@ -107,7 +107,7 @@ __global__ void gridwise_implicit_gemm_convolution_1_chwn_csrk_khwn_padded( if(get_thread_local_1d_id() == 0 && get_block_1d_id() == 0) { print_ConstantTensorDescriptor(in_chwn_block_desc, "in_chwn_block_desc"); - print_ConstantTensorDescriptor(wei_csrk_block_desc, "wei_csrk_block_desc"); + print_ConstantTensorDescriptor(wei_cyxk_block_desc, "wei_cyxk_block_desc"); print_ConstantTensorDescriptor(out_hkwn_thread_desc, "out_hkwn_thread_desc"); } #endif @@ -148,9 +148,9 @@ __global__ void gridwise_implicit_gemm_convolution_1_chwn_csrk_khwn_padded( constexpr auto blockwise_wei_copy = Blockwise4dTensorCopy1{}; + decltype(wei_cyxk_global_desc), + decltype(wei_cyxk_block_desc), + decltype(wei_cyxk_block_desc.GetLengths())>{}; #elif 0 // weight: format is [C*Y*X,K] constexpr auto blockwise_wei_copy = @@ -177,7 +177,7 @@ __global__ void gridwise_implicit_gemm_convolution_1_chwn_csrk_khwn_padded( // B_matrix[C,Wo*N] is a sub-matrix of in_block[C,Hi,Wi,N] // C_matrix[K,Wo*N] is a sub-matrix of out_block[Ho,K,Wo,N] constexpr auto a_cxk_block_mtx_desc = make_ConstantMatrixDescriptor( - Number{}, Number{}, Number{}); + Number{}, Number{}, Number{}); constexpr auto b_cxwn_block_mtx_desc = make_ConstantMatrixDescriptor(Number{}, @@ -205,7 +205,7 @@ __global__ void gridwise_implicit_gemm_convolution_1_chwn_csrk_khwn_padded( // LDS constexpr unsigned in_block_size = in_chwn_block_desc.GetElementSpace(); - constexpr unsigned wei_block_size = wei_csrk_block_desc.GetElementSpace(); + constexpr unsigned wei_block_size = wei_cyxk_block_desc.GetElementSpace(); __shared__ Float p_in_block[in_block_size]; __shared__ Float p_wei_block[wei_block_size]; @@ -245,14 +245,14 @@ __global__ void gridwise_implicit_gemm_convolution_1_chwn_csrk_khwn_padded( __syncthreads(); // a series of batched GEMM - for(unsigned s = 0; s < Y; ++s) + for(unsigned y = 0; y < Y; ++y) { - for(unsigned r = 0; r < X; ++r) + for(unsigned x = 0; x < X; ++x) { auto f_accum = [](auto& acc, const auto&& v) { acc += v; }; - blockwise_batch_gemm.Run(p_wei_block + wei_csrk_block_desc.Get1dIndex(0, s, r, 0), - p_in_block + in_chwn_block_desc.Get1dIndex(0, s, r, 0), + blockwise_batch_gemm.Run(p_wei_block + wei_cyxk_block_desc.Get1dIndex(0, y, x, 0), + p_in_block + in_chwn_block_desc.Get1dIndex(0, y, x, 0), p_out_thread, f_accum); } diff --git a/src/include/gridwise_implicit_gemm_convolution_2_chwn_csrk_khwn_lds_double_buffer.hip.hpp b/src/include/gridwise_implicit_gemm_convolution_2_chwn_cyxk_khwn_lds_double_buffer.hip.hpp similarity index 92% rename from src/include/gridwise_implicit_gemm_convolution_2_chwn_csrk_khwn_lds_double_buffer.hip.hpp rename to src/include/gridwise_implicit_gemm_convolution_2_chwn_cyxk_khwn_lds_double_buffer.hip.hpp index eba1d09675..7c802266d8 100644 --- a/src/include/gridwise_implicit_gemm_convolution_2_chwn_csrk_khwn_lds_double_buffer.hip.hpp +++ b/src/include/gridwise_implicit_gemm_convolution_2_chwn_cyxk_khwn_lds_double_buffer.hip.hpp @@ -34,7 +34,7 @@ template -__global__ void gridwise_implicit_gemm_convolution_2_chwn_csrk_khwn_lds_double_buffer( +__global__ void gridwise_implicit_gemm_convolution_2_chwn_cyxk_khwn_lds_double_buffer( const Float* const __restrict__ p_in_global, const Float* const __restrict__ p_wei_global, Float* const __restrict__ p_out_global) @@ -45,7 +45,7 @@ __global__ void gridwise_implicit_gemm_convolution_2_chwn_csrk_khwn_lds_double_b constexpr auto I3 = Number<3>{}; constexpr auto in_chwn_global_desc = InGlobalDesc{}; - constexpr auto wei_csrk_global_desc = WeiGlobalDesc{}; + constexpr auto wei_cyxk_global_desc = WeiGlobalDesc{}; constexpr auto out_khwn_global_desc = OutGlobalDesc{}; constexpr unsigned C = in_chwn_global_desc.GetLength(I0); @@ -57,8 +57,8 @@ __global__ void gridwise_implicit_gemm_convolution_2_chwn_csrk_khwn_lds_double_b constexpr unsigned Ho = out_khwn_global_desc.GetLength(I1); constexpr unsigned Wo = out_khwn_global_desc.GetLength(I2); - constexpr unsigned Y = wei_csrk_global_desc.GetLength(I1); - constexpr unsigned X = wei_csrk_global_desc.GetLength(I2); + constexpr unsigned Y = wei_cyxk_global_desc.GetLength(I1); + constexpr unsigned X = wei_cyxk_global_desc.GetLength(I2); constexpr unsigned B = N * Hi * Wi; constexpr unsigned BGhostRead = (Y - 1) * Wi + (X - 1); @@ -85,7 +85,7 @@ __global__ void gridwise_implicit_gemm_convolution_2_chwn_csrk_khwn_lds_double_b constexpr auto wei_ek_block_desc = make_ConstantTensorDescriptor_aligned( Sequence{}, Number{}); - constexpr auto wei_csrk_block_desc = make_ConstantTensorDescriptor_aligned( + constexpr auto wei_cyxk_block_desc = make_ConstantTensorDescriptor_aligned( Sequence{}, Number{}); // tensor view of threadwise output in register @@ -96,14 +96,14 @@ __global__ void gridwise_implicit_gemm_convolution_2_chwn_csrk_khwn_lds_double_b if(get_thread_local_1d_id() == 0 && get_block_1d_id() == 0) { print_ConstantTensorDescriptor(in_chwn_global_desc, "in_chwn_global_desc"); - print_ConstantTensorDescriptor(wei_csrk_global_desc, "wei_csrk_global_desc"); + print_ConstantTensorDescriptor(wei_cyxk_global_desc, "wei_cyxk_global_desc"); print_ConstantTensorDescriptor(out_khwn_global_desc, "out_khwn_global_desc"); print_ConstantTensorDescriptor(in_cb_global_desc, "in_cb_global_desc"); print_ConstantTensorDescriptor(wei_ek_global_desc, "wei_ek_global_desc"); print_ConstantTensorDescriptor(in_cb_block_desc, "in_cb_block_desc"); - print_ConstantTensorDescriptor(wei_csrk_block_desc, "wei_csrk_block_desc"); + print_ConstantTensorDescriptor(wei_cyxk_block_desc, "wei_cyxk_block_desc"); print_ConstantTensorDescriptor(wei_ek_block_desc, "wei_ek_block_desc"); print_ConstantTensorDescriptor(out_kb_thread_desc, "out_kb_thread_desc"); @@ -170,7 +170,7 @@ __global__ void gridwise_implicit_gemm_convolution_2_chwn_csrk_khwn_lds_double_b // b_mtx[C,B] is a subset of in_block[C,B + BGhostRead] // c_mtx[K,B] is out_block[K,B] constexpr auto a_cxk_block_mtx_desc = make_ConstantMatrixDescriptor( - Number{}, Number{}, Number{}); + Number{}, Number{}, Number{}); constexpr auto b_cxb_block_mtx_desc = make_ConstantMatrixDescriptor( Number{}, Number{}, Number{}); @@ -210,7 +210,7 @@ __global__ void gridwise_implicit_gemm_convolution_2_chwn_csrk_khwn_lds_double_b in_cb_block_desc.GetElementSpace(Number{}); constexpr unsigned wei_block_size = - wei_csrk_block_desc.GetElementSpace(Number{}); + wei_cyxk_block_desc.GetElementSpace(Number{}); constexpr unsigned max_align = InBlockCopyDataPerRead > WeiBlockCopyDataPerRead ? InBlockCopyDataPerRead @@ -227,14 +227,14 @@ __global__ void gridwise_implicit_gemm_convolution_2_chwn_csrk_khwn_lds_double_b p_in_global + in_cb_global_desc.Get1dIndex(0, b_block_data_begin); const Float* p_wei_global_block_offset = - p_wei_global + wei_csrk_global_desc.Get1dIndex(0, 0, 0, k_block_data_begin); + p_wei_global + wei_cyxk_global_desc.Get1dIndex(0, 0, 0, k_block_data_begin); // preload data into LDS blockwise_in_copy.Run(p_in_global_block_offset, p_in_block_0); blockwise_wei_copy.Run(p_wei_global_block_offset, p_wei_block_0); p_in_global_block_offset += CPerBlock * in_cb_global_desc.GetStride(I0); - p_wei_global_block_offset += CPerBlock * wei_csrk_global_desc.GetStride(I0); + p_wei_global_block_offset += CPerBlock * wei_cyxk_global_desc.GetStride(I0); // register Float p_out_thread[out_kb_thread_desc.GetElementSpace()]; @@ -247,7 +247,7 @@ __global__ void gridwise_implicit_gemm_convolution_2_chwn_csrk_khwn_lds_double_b for(unsigned c_block_data_begin = 0; c_block_data_begin + CPerBlock < C; c_block_data_begin += CPerBlock, p_in_global_block_offset += CPerBlock * in_cb_global_desc.GetStride(I0), - p_wei_global_block_offset += CPerBlock * wei_csrk_global_desc.GetStride(I0), + p_wei_global_block_offset += CPerBlock * wei_cyxk_global_desc.GetStride(I0), even_loop = !even_loop) { Float* p_in_block_now = even_loop ? p_in_block_0 : p_in_block_1; @@ -275,9 +275,9 @@ __global__ void gridwise_implicit_gemm_convolution_2_chwn_csrk_khwn_lds_double_b // compute on current data // a series of GEMM - for(unsigned s = 0; s < Y; ++s) + for(unsigned y = 0; y < Y; ++y) { - for(unsigned r = 0; r < X; ++r) + for(unsigned x = 0; x < X; ++x) { auto f_accum = [](auto& acc, const auto&& v) { acc += v; }; #if 1 @@ -285,8 +285,8 @@ __global__ void gridwise_implicit_gemm_convolution_2_chwn_csrk_khwn_lds_double_b #else blockwise_gemm.Run_RegisterDoubleBuffer #endif - (p_wei_block_now + wei_csrk_block_desc.Get1dIndex(0, s, r, 0), - p_in_block_now + s * Wi + r, + (p_wei_block_now + wei_cyxk_block_desc.Get1dIndex(0, y, x, 0), + p_in_block_now + y * Wi + x, p_out_thread, f_accum); } @@ -305,9 +305,9 @@ __global__ void gridwise_implicit_gemm_convolution_2_chwn_csrk_khwn_lds_double_b __syncthreads(); - for(unsigned s = 0; s < Y; ++s) + for(unsigned y = 0; y < Y; ++y) { - for(unsigned r = 0; r < X; ++r) + for(unsigned x = 0; x < X; ++x) { auto f_accum = [](auto& acc, const auto&& v) { acc += v; }; #if 0 @@ -315,8 +315,8 @@ __global__ void gridwise_implicit_gemm_convolution_2_chwn_csrk_khwn_lds_double_b #else blockwise_gemm.Run_RegisterDoubleBuffer #endif - (p_wei_block_now + wei_csrk_block_desc.Get1dIndex(0, s, r, 0), - p_in_block_now + s * Wi + r, + (p_wei_block_now + wei_cyxk_block_desc.Get1dIndex(0, y, x, 0), + p_in_block_now + y * Wi + x, p_out_thread, f_accum); } diff --git a/src/include/tensor.hpp b/src/include/tensor.hpp index 09ac224007..d0c785c16e 100644 --- a/src/include/tensor.hpp +++ b/src/include/tensor.hpp @@ -8,16 +8,16 @@ #include template -std::ostream& LogRange(std::ostream& os, Range&& r, std::string delim) +std::ostream& LogRange(std::ostream& os, Range&& range, std::string delim) { bool first = true; - for(auto&& x : r) + for(auto&& v : range) { if(first) first = false; else os << delim; - os << x; + os << v; } return os; } diff --git a/src/include/threadwise_direct_convolution.hip.hpp b/src/include/threadwise_direct_convolution.hip.hpp index 04ea7b1506..32d446491b 100644 --- a/src/include/threadwise_direct_convolution.hip.hpp +++ b/src/include/threadwise_direct_convolution.hip.hpp @@ -38,16 +38,16 @@ __device__ void threadwise_direct_convolution_1(InDesc, { for(unsigned c = 0; c < wei_desc.GetLength(I1); ++c) { - for(unsigned s = 0; s < wei_desc.GetLength(I2); ++s) + for(unsigned y = 0; y < wei_desc.GetLength(I2); ++y) { - for(unsigned r = 0; r < wei_desc.GetLength(I3); ++r) + for(unsigned x = 0; x < wei_desc.GetLength(I3); ++x) { - const unsigned hi = ho + s; - const unsigned wi = wo + r; + const unsigned hi = ho + y; + const unsigned wi = wo + x; const unsigned in_index = in_desc.Get1dIndex(n, c, hi, wi); - const unsigned wei_index = wei_desc.Get1dIndex(k, c, s, r); + const unsigned wei_index = wei_desc.Get1dIndex(k, c, y, x); const unsigned out_index = out_desc.Get1dIndex(n, k, ho, wo); @@ -153,18 +153,18 @@ __device__ void threadwise_direct_convolution_3(InDesc, #if 0 // this verison reused old input data in register, and read new data from LDS // loop over vertical direction - for(unsigned s = 0; s < wei_desc.GetLength(I2); ++s) + for(unsigned y = 0; y < wei_desc.GetLength(I2); ++y) { // read first input threadwise_4d_tensor_copy(in_desc, - p_in + in_desc.Get1dIndex(0, 0, s, 0), + p_in + in_desc.Get1dIndex(0, 0, y, 0), in_reg_desc, p_in_reg, in_reg_desc.GetLengths()); // read first 1x1 weight threadwise_4d_tensor_copy(wei_desc, - p_wei + wei_desc.Get1dIndex(0, 0, s, 0), + p_wei + wei_desc.Get1dIndex(0, 0, y, 0), wei_reg_desc, p_wei_reg, wei_reg_desc.GetLengths()); @@ -174,11 +174,11 @@ __device__ void threadwise_direct_convolution_3(InDesc, in_reg_desc, p_in_reg, wei_reg_desc, p_wei_reg, out_desc, p_out); // loop over horizontal direction - for(unsigned r = 1; r < wei_desc.GetLength(I3); ++r) + for(unsigned x = 1; x < wei_desc.GetLength(I3); ++x) { // read new weight threadwise_4d_tensor_copy(wei_desc, - p_wei + wei_desc.Get1dIndex(0, 0, s, r), + p_wei + wei_desc.Get1dIndex(0, 0, y, x), wei_reg_desc, p_wei_reg, wei_reg_desc.GetLengths()); @@ -189,7 +189,7 @@ __device__ void threadwise_direct_convolution_3(InDesc, // read new input threadwise_4d_tensor_copy( in_desc, - p_in + in_desc.Get1dIndex(0, 0, s, r + in_reg_desc.GetLength(I3) - 1), + p_in + in_desc.Get1dIndex(0, 0, y, x + in_reg_desc.GetLength(I3) - 1), in_reg_desc, p_in_reg + in_reg_desc.Get1dIndex(0, 0, 0, in_reg_desc.GetLength(I3) - in_w_new_read), @@ -203,21 +203,21 @@ __device__ void threadwise_direct_convolution_3(InDesc, #elif 1 // this version read all input from LDS when filter moves // loop over vertical direction - for(unsigned s = 0; s < wei_desc.GetLength(I2); ++s) + for(unsigned y = 0; y < wei_desc.GetLength(I2); ++y) { // loop over horizontal direction - for(unsigned r = 0; r < wei_desc.GetLength(I3); ++r) + for(unsigned x = 0; x < wei_desc.GetLength(I3); ++x) { // read new weight threadwise_4d_tensor_copy(wei_desc, - p_wei + wei_desc.Get1dIndex(0, 0, s, r), + p_wei + wei_desc.Get1dIndex(0, 0, y, x), wei_reg_desc, p_wei_reg, wei_reg_desc.GetLengths()); // read new input threadwise_4d_tensor_copy(in_desc, - p_in + in_desc.Get1dIndex(0, 0, s, r), + p_in + in_desc.Get1dIndex(0, 0, y, x), in_reg_desc, p_in_reg, in_reg_desc.GetLengths());