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
9c3c435a
Commit
9c3c435a
authored
Jan 18, 2023
by
aska-0096
Browse files
groupconv: Sanity check[OK], Performance[Bad]
parent
abfc94b2
Changes
2
Show whitespace changes
Inline
Side-by-side
Showing
2 changed files
with
7 additions
and
25 deletions
+7
-25
example/30_grouped_conv_fwd_multiple_d/run_grouped_conv_fwd_bias_relu_add_wmma_example.inc
...ple_d/run_grouped_conv_fwd_bias_relu_add_wmma_example.inc
+1
-1
include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_d_wmma_cshuffle.hpp
...impl/device_grouped_conv_fwd_multiple_d_wmma_cshuffle.hpp
+6
-24
No files found.
example/30_grouped_conv_fwd_multiple_d/run_grouped_conv_fwd_bias_relu_add_wmma_example.inc
View file @
9c3c435a
...
@@ -54,7 +54,7 @@ using DeviceConvFwdInstance =
...
@@ -54,7 +54,7 @@ using DeviceConvFwdInstance =
256
,
// BlockSize
256
,
// BlockSize
128
,
// MPerBlock
128
,
// MPerBlock
256
,
// NPerBlock
256
,
// NPerBlock
8
,
// K0PerBlock
4
,
// K0PerBlock
8
,
// K1
8
,
// K1
16
,
// MPerWMMA
16
,
// MPerWMMA
16
,
// NPerWMMA
16
,
// NPerWMMA
...
...
include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_d_wmma_cshuffle.hpp
View file @
9c3c435a
...
@@ -435,20 +435,12 @@ struct DeviceGroupedConvFwdMultipleD_Wmma_CShuffle
...
@@ -435,20 +435,12 @@ struct DeviceGroupedConvFwdMultipleD_Wmma_CShuffle
ds_grid_desc_m_n_
=
DeviceOp
::
MakeDsGridDescriptor_M_N
(
ds_g_n_k_wos_lengths
,
ds_g_n_k_wos_strides
);
ds_grid_desc_m_n_
=
DeviceOp
::
MakeDsGridDescriptor_M_N
(
ds_g_n_k_wos_lengths
,
ds_g_n_k_wos_strides
);
// populate desc for Ds/E
// populate desc for Ds/E
if
(
GridwiseOp
::
CheckValidity
(
a_grid_desc_ak0_m_ak1_
,
e_grid_desc_mblock_mperblock_nblock_nperblock_
=
b_grid_desc_bk0_n_bk1_
,
GridwiseOp
::
MakeEGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock
(
ds_grid_desc_m_n_
,
e_grid_desc_m_n_
);
e_grid_desc_m_n_
,
ds_grid_desc_mblock_mperblock_nblock_nperblock_
=
block_2_etile_map_
))
GridwiseOp
::
MakeDsGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock
(
{
ds_grid_desc_m_n_
);
// e_grid_desc_mblock_mperblock_nblock_nperblock_ =
// GridwiseOp::MakeEGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock(
// e_grid_desc_m_n_);
// ds_grid_desc_mblock_mperblock_nblock_nperblock_ =
// GridwiseOp::MakeDsGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock(
// ds_grid_desc_m_n_);
}
}
}
void
Print
()
const
void
Print
()
const
...
@@ -520,16 +512,6 @@ struct DeviceGroupedConvFwdMultipleD_Wmma_CShuffle
...
@@ -520,16 +512,6 @@ struct DeviceGroupedConvFwdMultipleD_Wmma_CShuffle
arg
.
Print
();
arg
.
Print
();
}
}
if
(
!
GridwiseOp
::
CheckValidity
(
arg
.
a_grid_desc_ak0_m_ak1_
,
arg
.
b_grid_desc_bk0_n_bk1_
,
arg
.
ds_grid_desc_m_n_
,
arg
.
e_grid_desc_m_n_
,
arg
.
block_2_etile_map_
))
{
throw
std
::
runtime_error
(
"wrong! GridwiseGemmMultipleD_wmma_cshuffle has invalid setting"
);
}
const
index_t
grid_size
=
const
index_t
grid_size
=
arg
.
block_2_etile_map_
.
CalculateGridSize
(
arg
.
e_grid_desc_m_n_
)
*
arg
.
num_group_
;
arg
.
block_2_etile_map_
.
CalculateGridSize
(
arg
.
e_grid_desc_m_n_
)
*
arg
.
num_group_
;
...
...
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