Skip to content
GitLab
Menu
Projects
Groups
Snippets
Loading...
Help
Help
Support
Community forum
Keyboard shortcuts
?
Submit feedback
Contribute to GitLab
Sign in / Register
Toggle navigation
Menu
Open sidebar
gaoqiong
composable_kernel
Commits
fe797cf2
"test/git@developer.sourcefind.cn:OpenDAS/torchaudio.git" did not exist on "26941fa3777618dfc659e638e524b65f22dd32a6"
Commit
fe797cf2
authored
Sep 07, 2021
by
ltqin
Browse files
rename device convolution file and function name
parent
ad7bd495
Changes
3
Hide whitespace changes
Inline
Side-by-side
Showing
3 changed files
with
22 additions
and
24 deletions
+22
-24
host/driver_offline/include/device_convolution_backward_weight_implicit_gemm_v4r4r2_xdlops_atomic_nchw_kcyx_nkhw.hpp
...ght_implicit_gemm_v4r4r2_xdlops_atomic_nchw_kcyx_nkhw.hpp
+1
-1
host/driver_offline/include/driver_gemm_xdlops_v2r4.hpp
host/driver_offline/include/driver_gemm_xdlops_v2r4.hpp
+2
-4
host/driver_offline/src/conv_wrw_driver_offline.cpp
host/driver_offline/src/conv_wrw_driver_offline.cpp
+19
-19
No files found.
host/driver_offline/include/device_convolution_backward_weight_implicit_gemm_v4r4r
3
_xdlops_nchw_kcyx_nkhw.hpp
→
host/driver_offline/include/device_convolution_backward_weight_implicit_gemm_v4r4r
2
_xdlops_
atomic_
nchw_kcyx_nkhw.hpp
View file @
fe797cf2
...
@@ -14,7 +14,7 @@ template <typename TInWei,
...
@@ -14,7 +14,7 @@ template <typename TInWei,
typename
ConvDilations
,
typename
ConvDilations
,
typename
InLeftPads
,
typename
InLeftPads
,
typename
InRightPads
>
typename
InRightPads
>
void
device_convolution_backward_weight_implicit_gemm_v4r4r
3
_xdlops_nchw_kcyx_nkhw
(
void
device_convolution_backward_weight_implicit_gemm_v4r4r
2
_xdlops_
atomic_
nchw_kcyx_nkhw
(
const
InLengths
&
in_n_c_hi_wi_lengths
,
const
InLengths
&
in_n_c_hi_wi_lengths
,
const
WeiLengths
&
wei_k_c_y_x_lengths
,
const
WeiLengths
&
wei_k_c_y_x_lengths
,
const
OutLengths
&
out_n_k_ho_wo_lengths
,
const
OutLengths
&
out_n_k_ho_wo_lengths
,
...
...
host/driver_offline/include/driver_gemm_xdlops_v2r4.hpp
View file @
fe797cf2
...
@@ -128,9 +128,7 @@ __host__ float driver_gemm_xdlops_v2r4(const FloatAB* p_a_grid,
...
@@ -128,9 +128,7 @@ __host__ float driver_gemm_xdlops_v2r4(const FloatAB* p_a_grid,
// const auto kbatch = GridwiseGemm::CalculateKBatch(c_m_n_grid_desc, b_k0_n_k1_grid_desc);
// const auto kbatch = GridwiseGemm::CalculateKBatch(c_m_n_grid_desc, b_k0_n_k1_grid_desc);
const
auto
a_b_k0_m_k1_grid_desc
=
GridwiseGemm
::
MakeABK0MK1GridDescriptor
(
a_k0_m_k1_grid_desc
);
const
auto
a_b_k0_m_k1_grid_desc
=
GridwiseGemm
::
MakeABK0MK1GridDescriptor
(
a_k0_m_k1_grid_desc
);
const
auto
b_b_k0_n_k1_grid_desc
=
GridwiseGemm
::
MakeBBK0NK1GridDescriptor
(
b_k0_n_k1_grid_desc
);
const
auto
b_b_k0_n_k1_grid_desc
=
GridwiseGemm
::
MakeBBK0NK1GridDescriptor
(
b_k0_n_k1_grid_desc
);
{
// std::cout << "k batch number is: " << kbatch << std::endl;
}
if
(
!
GridwiseGemm
::
CheckValidity
(
a_k0_m_k1_grid_desc
,
b_k0_n_k1_grid_desc
,
c_m_n_grid_desc
))
if
(
!
GridwiseGemm
::
CheckValidity
(
a_k0_m_k1_grid_desc
,
b_k0_n_k1_grid_desc
,
c_m_n_grid_desc
))
{
{
throw
std
::
runtime_error
(
throw
std
::
runtime_error
(
...
@@ -150,7 +148,7 @@ __host__ float driver_gemm_xdlops_v2r4(const FloatAB* p_a_grid,
...
@@ -150,7 +148,7 @@ __host__ float driver_gemm_xdlops_v2r4(const FloatAB* p_a_grid,
const
index_t
grid_size
=
GridwiseGemm
::
CalculateGridSize
(
c_m_n_grid_desc
);
const
index_t
grid_size
=
GridwiseGemm
::
CalculateGridSize
(
c_m_n_grid_desc
);
{
{
std
::
cout
<<
"gridSize : "
<<
grid_size
<<
grid_size
<<
std
::
endl
;
std
::
cout
<<
"gridSize : "
<<
grid_size
<<
std
::
endl
;
}
}
const
auto
kernel
=
kernel_gemm_xdlops_v2r4
<
GridwiseGemm
,
const
auto
kernel
=
kernel_gemm_xdlops_v2r4
<
GridwiseGemm
,
FloatAB
,
FloatAB
,
...
...
host/driver_offline/src/conv_wrw_driver_offline.cpp
View file @
fe797cf2
...
@@ -13,16 +13,16 @@
...
@@ -13,16 +13,16 @@
#include "host_conv_bwd_weight.hpp"
#include "host_conv_bwd_weight.hpp"
#include "device_tensor.hpp"
#include "device_tensor.hpp"
#include "device_convolution_backward_weight_implicit_gemm_v4r4r2_xdlops_nchw_kcyx_nkhw.hpp"
#include "device_convolution_backward_weight_implicit_gemm_v4r4r2_xdlops_nchw_kcyx_nkhw.hpp"
#include "device_convolution_backward_weight_implicit_gemm_v4r4r
3
_xdlops_nchw_kcyx_nkhw.hpp"
#include "device_convolution_backward_weight_implicit_gemm_v4r4r
2
_xdlops_
atomic_
nchw_kcyx_nkhw.hpp"
#define USE_DYNAMIC_MODE 1
#define USE_DYNAMIC_MODE 1
#define USE_CONV_WRW_V4R4R2_XDL_NCHW 1
#define USE_CONV_WRW_V4R4R2_XDL_NCHW 1
#define USE_CONV_WRW_V4R4R
3
_XDL_NCHW 1
#define USE_CONV_WRW_V4R4R
2
_XDL_
ATOMIC_
NCHW 1
enum
ConvBackwardWeightAlgo
enum
ConvBackwardWeightAlgo
{
{
V4R4R2XDLNCHW
,
V4R4R2XDLNCHW
,
V4R4R
3
XDLNCHW
,
V4R4R
2
XDL
ATOMIC
NCHW
,
};
};
int
main
(
int
argc
,
char
*
argv
[])
int
main
(
int
argc
,
char
*
argv
[])
...
@@ -262,8 +262,8 @@ int main(int argc, char* argv[])
...
@@ -262,8 +262,8 @@ int main(int argc, char* argv[])
}
}
#endif
#endif
#if USE_CONV_WRW_V4R4R
3
_XDL_NCHW
#if USE_CONV_WRW_V4R4R
2
_XDL_
ATOMIC_
NCHW
if
(
algo
==
ConvBackwardWeightAlgo
::
V4R4R
3
XDLNCHW
)
if
(
algo
==
ConvBackwardWeightAlgo
::
V4R4R
2
XDL
ATOMIC
NCHW
)
{
{
if
(
layout
!=
ConvTensorLayout
::
NCHW
)
if
(
layout
!=
ConvTensorLayout
::
NCHW
)
{
{
...
@@ -272,20 +272,20 @@ int main(int argc, char* argv[])
...
@@ -272,20 +272,20 @@ int main(int argc, char* argv[])
const
auto
tmp
=
f_make_for_device_nchw
();
const
auto
tmp
=
f_make_for_device_nchw
();
device_convolution_backward_weight_implicit_gemm_v4r4r
3
_xdlops_nchw_kcyx_nkhw
<
in_data_t
,
device_convolution_backward_weight_implicit_gemm_v4r4r
2
_xdlops_
atomic_
nchw_kcyx_nkhw
<
acc
_data_t
,
in
_data_t
,
out
_data_t
>
(
acc
_data_t
,
tmp
[
I0
],
out_data_t
>
(
tmp
[
I0
],
tmp
[
I1
],
tmp
[
I1
],
tmp
[
I2
],
tmp
[
I2
],
tmp
[
I3
],
tmp
[
I3
],
tmp
[
I4
],
tmp
[
I4
],
tmp
[
I5
],
tmp
[
I5
],
tmp
[
I6
],
tmp
[
I6
],
in
,
in
,
wei_device
,
wei_device
,
out
,
out
,
nrepeat
);
nrepeat
);
}
}
#endif
#endif
...
...
Write
Preview
Markdown
is supported
0%
Try again
or
attach a new file
.
Attach a file
Cancel
You are about to add
0
people
to the discussion. Proceed with caution.
Finish editing this message first!
Cancel
Please
register
or
sign in
to comment