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
4511f877
Commit
4511f877
authored
May 09, 2022
by
Chao Liu
Browse files
refactor profiler
parent
519b6aaf
Changes
69
Expand all
Hide whitespace changes
Inline
Side-by-side
Showing
20 changed files
with
196 additions
and
705 deletions
+196
-705
library/src/tensor_operation_instance/gpu/gemm_splitk/device_gemm_xdl_splitk_c_shuffle_f16_f16_f16_mk_kn_mn_instance.cpp
...mm_xdl_splitk_c_shuffle_f16_f16_f16_mk_kn_mn_instance.cpp
+0
-0
library/src/tensor_operation_instance/gpu/gemm_splitk/device_gemm_xdl_splitk_c_shuffle_f16_f16_f16_mk_nk_mn_instance.cpp
...mm_xdl_splitk_c_shuffle_f16_f16_f16_mk_nk_mn_instance.cpp
+0
-0
library/src/tensor_operation_instance/gpu/gemm_splitk/device_gemm_xdl_splitk_f32_f32_f32_km_kn_mn_instance.cpp
.../device_gemm_xdl_splitk_f32_f32_f32_km_kn_mn_instance.cpp
+0
-0
library/src/tensor_operation_instance/gpu/gemm_splitk/device_gemm_xdl_splitk_f32_f32_f32_km_nk_mn_instance.cpp
.../device_gemm_xdl_splitk_f32_f32_f32_km_nk_mn_instance.cpp
+0
-0
library/src/tensor_operation_instance/gpu/gemm_splitk/device_gemm_xdl_splitk_f32_f32_f32_mk_kn_mn_instance.cpp
.../device_gemm_xdl_splitk_f32_f32_f32_mk_kn_mn_instance.cpp
+0
-0
library/src/tensor_operation_instance/gpu/gemm_splitk/device_gemm_xdl_splitk_f32_f32_f32_mk_nk_mn_instance.cpp
.../device_gemm_xdl_splitk_f32_f32_f32_mk_nk_mn_instance.cpp
+0
-0
profiler/CMakeLists.txt
profiler/CMakeLists.txt
+3
-3
profiler/include/profile_batched_gemm_impl.hpp
profiler/include/profile_batched_gemm_impl.hpp
+13
-34
profiler/include/profile_batched_gemm_reduce_impl.hpp
profiler/include/profile_batched_gemm_reduce_impl.hpp
+7
-8
profiler/include/profile_conv_bwd_data_impl.hpp
profiler/include/profile_conv_bwd_data_impl.hpp
+10
-3
profiler/include/profile_conv_bwd_weight_impl.hpp
profiler/include/profile_conv_bwd_weight_impl.hpp
+8
-12
profiler/include/profile_conv_fwd_bias_relu_add_impl.hpp
profiler/include/profile_conv_fwd_bias_relu_add_impl.hpp
+7
-3
profiler/include/profile_conv_fwd_bias_relu_atomic_add_impl.hpp
...er/include/profile_conv_fwd_bias_relu_atomic_add_impl.hpp
+0
-330
profiler/include/profile_conv_fwd_bias_relu_impl.hpp
profiler/include/profile_conv_fwd_bias_relu_impl.hpp
+7
-3
profiler/include/profile_convnd_bwd_data_impl.hpp
profiler/include/profile_convnd_bwd_data_impl.hpp
+27
-55
profiler/include/profile_convnd_fwd.hpp
profiler/include/profile_convnd_fwd.hpp
+0
-9
profiler/include/profile_gemm_bias_2d_impl.hpp
profiler/include/profile_gemm_bias_2d_impl.hpp
+15
-14
profiler/include/profile_gemm_bias_relu_add_impl.hpp
profiler/include/profile_gemm_bias_relu_add_impl.hpp
+15
-13
profiler/include/profile_gemm_bias_relu_impl.hpp
profiler/include/profile_gemm_bias_relu_impl.hpp
+11
-9
profiler/include/profile_gemm_impl.hpp
profiler/include/profile_gemm_impl.hpp
+73
-209
No files found.
library/src/tensor_operation_instance/gpu/gemm/device_gemm_xdl_splitk_c_shuffle_f16_f16_f16_mk_kn_mn_instance.cpp
→
library/src/tensor_operation_instance/gpu/gemm
_splitk
/device_gemm_xdl_splitk_c_shuffle_f16_f16_f16_mk_kn_mn_instance.cpp
View file @
4511f877
File moved
library/src/tensor_operation_instance/gpu/gemm/device_gemm_xdl_splitk_c_shuffle_f16_f16_f16_mk_nk_mn_instance.cpp
→
library/src/tensor_operation_instance/gpu/gemm
_splitk
/device_gemm_xdl_splitk_c_shuffle_f16_f16_f16_mk_nk_mn_instance.cpp
View file @
4511f877
File moved
library/src/tensor_operation_instance/gpu/gemm/device_gemm_xdl_splitk_f32_f32_f32_km_kn_mn_instance.cpp
→
library/src/tensor_operation_instance/gpu/gemm
_splitk
/device_gemm_xdl_splitk_f32_f32_f32_km_kn_mn_instance.cpp
View file @
4511f877
File moved
library/src/tensor_operation_instance/gpu/gemm/device_gemm_xdl_splitk_f32_f32_f32_km_nk_mn_instance.cpp
→
library/src/tensor_operation_instance/gpu/gemm
_splitk
/device_gemm_xdl_splitk_f32_f32_f32_km_nk_mn_instance.cpp
View file @
4511f877
File moved
library/src/tensor_operation_instance/gpu/gemm/device_gemm_xdl_splitk_f32_f32_f32_mk_kn_mn_instance.cpp
→
library/src/tensor_operation_instance/gpu/gemm
_splitk
/device_gemm_xdl_splitk_f32_f32_f32_mk_kn_mn_instance.cpp
View file @
4511f877
File moved
library/src/tensor_operation_instance/gpu/gemm/device_gemm_xdl_splitk_f32_f32_f32_mk_nk_mn_instance.cpp
→
library/src/tensor_operation_instance/gpu/gemm
_splitk
/device_gemm_xdl_splitk_f32_f32_f32_mk_nk_mn_instance.cpp
View file @
4511f877
File moved
profiler/CMakeLists.txt
View file @
4511f877
...
...
@@ -24,6 +24,7 @@ include_directories(BEFORE
set
(
PROFILER_SOURCE
src/profiler.cpp
src/profile_gemm.cpp
src/profile_gemm_splitk.cpp
src/profile_gemm_bias_2d.cpp
src/profile_gemm_bias_relu.cpp
src/profile_gemm_bias_relu_add.cpp
...
...
@@ -31,7 +32,6 @@ set(PROFILER_SOURCE
src/profile_batched_gemm.cpp
src/profile_conv_fwd_bias_relu.cpp
src/profile_conv_fwd_bias_relu_add.cpp
src/profile_conv_fwd_bias_relu_atomic_add.cpp
src/profile_convnd_fwd.cpp
src/profile_convnd_bwd_data.cpp
src/profile_reduce.cpp
...
...
@@ -44,8 +44,9 @@ add_executable(ckProfiler ${PROFILER_SOURCE})
target_link_libraries
(
ckProfiler PRIVATE host_tensor
)
target_link_libraries
(
ckProfiler PRIVATE conv_fwd_util
)
target_link_libraries
(
ckProfiler PRIVATE device_gemm_reduce_instance
)
target_link_libraries
(
ckProfiler PRIVATE device_gemm_instance
)
target_link_libraries
(
ckProfiler PRIVATE device_gemm_splitk_instance
)
target_link_libraries
(
ckProfiler PRIVATE device_gemm_reduce_instance
)
target_link_libraries
(
ckProfiler PRIVATE device_gemm_bias2d_instance
)
target_link_libraries
(
ckProfiler PRIVATE device_gemm_bias_relu_instance
)
target_link_libraries
(
ckProfiler PRIVATE device_gemm_bias_relu_add_instance
)
...
...
@@ -55,7 +56,6 @@ target_link_libraries(ckProfiler PRIVATE device_conv2d_fwd_instance)
target_link_libraries
(
ckProfiler PRIVATE device_conv3d_fwd_instance
)
target_link_libraries
(
ckProfiler PRIVATE device_conv2d_fwd_bias_relu_instance
)
target_link_libraries
(
ckProfiler PRIVATE device_conv2d_fwd_bias_relu_add_instance
)
target_link_libraries
(
ckProfiler PRIVATE device_conv2d_fwd_bias_relu_atomic_add_instance
)
target_link_libraries
(
ckProfiler PRIVATE device_convnd_bwd_data_instance
)
target_link_libraries
(
ckProfiler PRIVATE device_reduce_instance
)
target_link_libraries
(
ckProfiler PRIVATE device_grouped_gemm_instance
)
...
...
profiler/include/profile_batched_gemm_impl.hpp
View file @
4511f877
...
...
@@ -37,14 +37,10 @@ void add_device_batched_gemm_xdl_f32_f32_f32_gmk_gkn_gmn_instances(std::vector<D
void
add_device_batched_gemm_xdl_f32_f32_f32_gmk_gnk_gmn_instances
(
std
::
vector
<
DeviceGemmNoOpPtr
>&
);
void
add_device_batched_gemm_xdl_f32_f32_f32_gkm_gkn_gmn_instances
(
std
::
vector
<
DeviceGemmNoOpPtr
>&
);
void
add_device_batched_gemm_xdl_f32_f32_f32_gkm_gnk_gmn_instances
(
std
::
vector
<
DeviceGemmNoOpPtr
>&
);
void
add_device_batched_gemm_xdl_int8_int8_int8_gmk_gkn_gmn_instances
(
std
::
vector
<
DeviceGemmNoOpPtr
>&
);
void
add_device_batched_gemm_xdl_int8_int8_int8_gmk_gnk_gmn_instances
(
std
::
vector
<
DeviceGemmNoOpPtr
>&
);
void
add_device_batched_gemm_xdl_int8_int8_int8_gkm_gkn_gmn_instances
(
std
::
vector
<
DeviceGemmNoOpPtr
>&
);
void
add_device_batched_gemm_xdl_int8_int8_int8_gkm_gnk_gmn_instances
(
std
::
vector
<
DeviceGemmNoOpPtr
>&
);
void
add_device_batched_gemm_xdl_i8_i8_i8_gmk_gkn_gmn_instances
(
std
::
vector
<
DeviceGemmNoOpPtr
>&
);
void
add_device_batched_gemm_xdl_i8_i8_i8_gmk_gnk_gmn_instances
(
std
::
vector
<
DeviceGemmNoOpPtr
>&
);
void
add_device_batched_gemm_xdl_i8_i8_i8_gkm_gkn_gmn_instances
(
std
::
vector
<
DeviceGemmNoOpPtr
>&
);
void
add_device_batched_gemm_xdl_i8_i8_i8_gkm_gnk_gmn_instances
(
std
::
vector
<
DeviceGemmNoOpPtr
>&
);
}
// namespace device_batched_gemm_instance
}
// namespace device
...
...
@@ -72,8 +68,6 @@ bool profile_batched_gemm_impl(int do_verification,
int
StrideC
,
int
BatchCount
)
{
bool
pass
=
true
;
auto
f_host_tensor_descriptor
=
[](
std
::
size_t
batch_count
,
std
::
size_t
row
,
std
::
size_t
col
,
...
...
@@ -297,40 +291,38 @@ bool profile_batched_gemm_impl(int do_verification,
is_same
<
CLayout
,
tensor_layout
::
gemm
::
RowMajor
>::
value
)
{
ck
::
tensor_operation
::
device
::
device_batched_gemm_instance
::
add_device_batched_gemm_xdl_i
nt
8_i
nt
8_i
nt
8_gmk_gkn_gmn_instances
(
gemm_ptrs
);
add_device_batched_gemm_xdl_i8_i8_i8_gmk_gkn_gmn_instances
(
gemm_ptrs
);
}
else
if
constexpr
(
is_same
<
ALayout
,
tensor_layout
::
gemm
::
RowMajor
>::
value
&&
is_same
<
BLayout
,
tensor_layout
::
gemm
::
ColumnMajor
>::
value
&&
is_same
<
CLayout
,
tensor_layout
::
gemm
::
RowMajor
>::
value
)
{
ck
::
tensor_operation
::
device
::
device_batched_gemm_instance
::
add_device_batched_gemm_xdl_i
nt
8_i
nt
8_i
nt
8_gmk_gnk_gmn_instances
(
gemm_ptrs
);
add_device_batched_gemm_xdl_i8_i8_i8_gmk_gnk_gmn_instances
(
gemm_ptrs
);
}
else
if
constexpr
(
is_same
<
ALayout
,
tensor_layout
::
gemm
::
ColumnMajor
>::
value
&&
is_same
<
BLayout
,
tensor_layout
::
gemm
::
RowMajor
>::
value
&&
is_same
<
CLayout
,
tensor_layout
::
gemm
::
RowMajor
>::
value
)
{
ck
::
tensor_operation
::
device
::
device_batched_gemm_instance
::
add_device_batched_gemm_xdl_i
nt
8_i
nt
8_i
nt
8_gkm_gkn_gmn_instances
(
gemm_ptrs
);
add_device_batched_gemm_xdl_i8_i8_i8_gkm_gkn_gmn_instances
(
gemm_ptrs
);
}
else
if
constexpr
(
is_same
<
ALayout
,
tensor_layout
::
gemm
::
ColumnMajor
>::
value
&&
is_same
<
BLayout
,
tensor_layout
::
gemm
::
ColumnMajor
>::
value
&&
is_same
<
CLayout
,
tensor_layout
::
gemm
::
RowMajor
>::
value
)
{
ck
::
tensor_operation
::
device
::
device_batched_gemm_instance
::
add_device_batched_gemm_xdl_i
nt
8_i
nt
8_i
nt
8_gkm_gnk_gmn_instances
(
gemm_ptrs
);
add_device_batched_gemm_xdl_i8_i8_i8_gkm_gnk_gmn_instances
(
gemm_ptrs
);
}
}
if
(
gemm_ptrs
.
size
()
<=
0
)
{
throw
std
::
runtime_error
(
"wrong! no device GEMM instance found"
);
}
std
::
cout
<<
"found "
<<
gemm_ptrs
.
size
()
<<
" instances"
<<
std
::
endl
;
std
::
string
best_gemm_name
;
float
best_ave_time
=
0
;
float
best_tflops
=
0
;
float
best_gb_per_sec
=
0
;
bool
pass
=
true
;
// profile device GEMM instances
for
(
auto
&
gemm_ptr
:
gemm_ptrs
)
...
...
@@ -383,20 +375,8 @@ bool profile_batched_gemm_impl(int do_verification,
{
c_device_buf
.
FromDevice
(
c_g_m_n_device_result
.
mData
.
data
());
if
constexpr
(
is_same
<
ADataType
,
ck
::
bhalf_t
>::
value
&&
is_same
<
BDataType
,
ck
::
bhalf_t
>::
value
&&
is_same
<
CDataType
,
ck
::
bhalf_t
>::
value
)
{
bf16_to_f32_
(
c_g_m_n_device_result
,
*
c_f32_g_m_n_device_result
);
float
err
=
check_error
(
*
c_f32_g_m_n_host_result
,
*
c_f32_g_m_n_device_result
);
pass
=
pass
&&
(
err
<
1E-6
);
}
else
{
float
err
=
check_error
(
c_g_m_n_host_result
,
c_g_m_n_device_result
);
pass
=
pass
&&
(
err
<
1E-6
);
}
pass
=
pass
&&
ck
::
utils
::
check_err
(
c_g_m_n_device_result
.
mData
,
c_g_m_n_host_result
.
mData
);
if
(
do_log
)
{
...
...
@@ -412,8 +392,7 @@ bool profile_batched_gemm_impl(int do_verification,
}
else
{
std
::
cout
<<
"this device GEMM instance does not support this GEMM problem"
<<
std
::
endl
;
std
::
cout
<<
"does not support this problem"
<<
std
::
endl
;
}
}
...
...
profiler/include/profile_batched_gemm_reduce_impl.hpp
View file @
4511f877
#pragma once
#include "check_err.hpp"
#include "config.hpp"
#include "device.hpp"
#include "host_tensor.hpp"
...
...
@@ -312,13 +313,11 @@ bool profile_batched_gemm_reduce_impl(int do_verification,
d0_device_buf
.
FromDevice
(
d0_g_m_device_result
.
mData
.
data
());
d1_device_buf
.
FromDevice
(
d1_g_m_device_result
.
mData
.
data
());
float
c_error
=
check_error
(
c_g_m_n_host_result
,
c_g_m_n_device_result
);
float
d0_error
=
check_error
(
d0_g_m_host_result
,
d0_g_m_device_result
);
float
d1_error
=
check_error
(
d1_g_m_host_result
,
d1_g_m_device_result
);
pass
=
pass
&&
(
c_error
<
1E-6
);
pass
=
pass
&&
(
d0_error
<
1E-6
);
pass
=
pass
&&
(
d1_error
<
1E-6
);
pass
=
pass
&&
ck
::
utils
::
check_err
(
c_g_m_n_device_result
.
mData
,
c_g_m_n_host_result
.
mData
)
&&
ck
::
utils
::
check_err
(
d0_g_m_device_result
.
mData
,
d0_g_m_host_result
.
mData
)
&&
ck
::
utils
::
check_err
(
d1_g_m_device_result
.
mData
,
d1_g_m_host_result
.
mData
);
if
(
do_log
)
{
...
...
@@ -344,7 +343,7 @@ bool profile_batched_gemm_reduce_impl(int do_verification,
}
else
{
std
::
cout
<<
"does not support this
GEMM
problem"
<<
std
::
endl
;
std
::
cout
<<
"does not support this problem"
<<
std
::
endl
;
}
}
...
...
profiler/include/profile_conv_bwd_data_impl.hpp
View file @
4511f877
...
...
@@ -48,7 +48,7 @@ template <int NDimSpatial,
typename
InLayout
,
typename
WeiLayout
,
typename
OutLayout
>
void
profile_conv_bwd_data_impl
(
int
do_verification
,
bool
profile_conv_bwd_data_impl
(
int
do_verification
,
int
init_method
,
bool
do_log
,
int
nrepeat
,
...
...
@@ -63,6 +63,8 @@ void profile_conv_bwd_data_impl(int do_verification,
std
::
vector
<
ck
::
index_t
>
input_left_pads
,
std
::
vector
<
ck
::
index_t
>
input_right_pads
)
{
bool
pass
=
true
;
const
ck
::
index_t
Y
=
filter_spatial_lengths
[
0
];
const
ck
::
index_t
X
=
filter_spatial_lengths
[
1
];
...
...
@@ -226,6 +228,9 @@ void profile_conv_bwd_data_impl(int do_verification,
if
(
conv_ptr
->
IsSupportedArgument
(
argument_ptr
.
get
()))
{
// re-init to zero before profiling next kernel
in_device_buf
.
SetZero
();
std
::
string
conv_name
=
conv_ptr
->
GetTypeString
();
float
ave_time
=
invoker_ptr
->
Run
(
argument_ptr
.
get
(),
nrepeat
);
...
...
@@ -255,8 +260,8 @@ void profile_conv_bwd_data_impl(int do_verification,
{
in_device_buf
.
FromDevice
(
in_n_c_hi_wi_device_result
.
mData
.
data
());
ck
::
utils
::
check_err
(
in_n_c_hi_wi_device_result
.
mData
,
in_n_c_hi_wi_host_result
.
mData
);
pass
=
pass
&&
ck
::
utils
::
check_err
(
in_n_c_hi_wi_device_result
.
mData
,
in_n_c_hi_wi_host_result
.
mData
);
if
(
do_log
)
{
...
...
@@ -277,6 +282,8 @@ void profile_conv_bwd_data_impl(int do_verification,
std
::
cout
<<
"Best Perf: "
<<
best_ave_time
<<
" ms, "
<<
best_tflops
<<
" TFlops, "
<<
best_gb_per_sec
<<
" GB/s, "
<<
best_conv_name
<<
std
::
endl
;
return
pass
;
}
}
// namespace profiler
...
...
profiler/include/profile_conv_bwd_weight_impl.hpp
View file @
4511f877
#pragma once
#include "check_err.hpp"
#include "config.hpp"
#include "device.hpp"
#include "host_tensor.hpp"
...
...
@@ -56,6 +57,8 @@ bool profile_conv_bwd_weight_impl(int do_verification,
std
::
vector
<
ck
::
index_t
>
input_right_pads
,
ck
::
index_t
split_k
)
{
bool
pass
=
true
;
const
ck
::
index_t
Y
=
filter_spatial_lengths
[
0
];
const
ck
::
index_t
X
=
filter_spatial_lengths
[
1
];
...
...
@@ -181,14 +184,11 @@ bool profile_conv_bwd_weight_impl(int do_verification,
float
best_gb_per_sec
=
0
;
// profile device Conv instances
bool
pass
=
true
;
for
(
auto
&
conv_ptr
:
conv_ptrs
)
{
// using atomic, so need to reset input
if
(
split_k
>
1
)
{
wei_device_buf
.
SetZero
();
}
// using atomic, so need to reset
wei_device_buf
.
SetZero
();
auto
argument_ptr
=
conv_ptr
->
MakeArgumentPointer
(
static_cast
<
InDataType
*>
(
in_device_buf
.
GetDeviceBuffer
()),
static_cast
<
WeiDataType
*>
(
wei_device_buf
.
GetDeviceBuffer
()),
...
...
@@ -241,12 +241,8 @@ bool profile_conv_bwd_weight_impl(int do_verification,
{
wei_device_buf
.
FromDevice
(
wei_k_c_y_x_device_result
.
mData
.
data
());
float
max_error
=
check_error
(
wei_k_c_y_x_host_result
,
wei_k_c_y_x_device_result
);
if
(
max_error
>
8
)
{
pass
=
false
;
std
::
cout
<<
"Fail info:"
<<
conv_ptr
->
GetTypeString
()
<<
std
::
endl
;
}
pass
=
pass
&&
ck
::
utils
::
check_err
(
wei_k_c_y_x_device_result
.
mData
,
wei_k_c_y_x_host_result
.
mData
);
if
(
do_log
)
{
...
...
profiler/include/profile_conv_fwd_bias_relu_add_impl.hpp
View file @
4511f877
...
...
@@ -39,7 +39,7 @@ template <int NDimSpatial,
typename
InLayout
,
typename
WeiLayout
,
typename
OutLayout
>
void
profile_conv_fwd_bias_relu_add_impl
(
int
do_verification
,
bool
profile_conv_fwd_bias_relu_add_impl
(
int
do_verification
,
int
init_method
,
bool
do_log
,
int
nrepeat
,
...
...
@@ -54,6 +54,8 @@ void profile_conv_fwd_bias_relu_add_impl(int do_verification,
std
::
vector
<
ck
::
index_t
>
input_left_pads
,
std
::
vector
<
ck
::
index_t
>
input_right_pads
)
{
bool
pass
=
true
;
const
ck
::
index_t
Y
=
filter_spatial_lengths
[
0
];
const
ck
::
index_t
X
=
filter_spatial_lengths
[
1
];
...
...
@@ -247,8 +249,8 @@ void profile_conv_fwd_bias_relu_add_impl(int do_verification,
{
out_device_buf
.
FromDevice
(
out_n_k_ho_wo_device_result
.
mData
.
data
());
ck
::
utils
::
check_err
(
out_n_k_ho_wo_device_result
.
mData
,
out_n_k_ho_wo_host_result
.
mData
);
pass
=
pass
&&
ck
::
utils
::
check_err
(
out_n_k_ho_wo_device_result
.
mData
,
out_n_k_ho_wo_host_result
.
mData
);
if
(
do_log
)
{
...
...
@@ -269,6 +271,8 @@ void profile_conv_fwd_bias_relu_add_impl(int do_verification,
std
::
cout
<<
"Best Perf: "
<<
best_ave_time
<<
" ms, "
<<
best_tflops
<<
" TFlops, "
<<
best_gb_per_sec
<<
" GB/s, "
<<
best_conv_name
<<
std
::
endl
;
return
pass
;
}
}
// namespace profiler
...
...
profiler/include/profile_conv_fwd_bias_relu_atomic_add_impl.hpp
deleted
100644 → 0
View file @
519b6aaf
#pragma once
#include "check_err.hpp"
#include "config.hpp"
#include "device.hpp"
#include "host_tensor.hpp"
#include "host_tensor_generator.hpp"
#include "host_conv.hpp"
#include "tensor_layout.hpp"
#include "device_tensor.hpp"
#include "device_conv_fwd_bias_activation.hpp"
#include "element_wise_operation.hpp"
namespace
ck
{
namespace
tensor_operation
{
namespace
device
{
namespace
device_conv2d_fwd_bias_activation_atomic_add_instance
{
using
DeviceConvFwdBiasReluPtr
=
DeviceConvFwdBiasActivationPtr
<
ck
::
tensor_operation
::
element_wise
::
PassThrough
,
ck
::
tensor_operation
::
element_wise
::
PassThrough
,
ck
::
tensor_operation
::
element_wise
::
AddRelu
>
;
void
add_device_conv2d_fwd_xdl_c_shuffle_bias_relu_atomic_add_nhwc_kyxc_nhwk_f16_instances
(
std
::
vector
<
DeviceConvFwdBiasReluPtr
>&
);
}
// namespace device_conv2d_fwd_bias_activation_atomic_add_instance
}
// namespace device
}
// namespace tensor_operation
}
// namespace ck
namespace
ck
{
namespace
profiler
{
void
cpu_conv_bias_relu_atomic_add
(
ck
::
half_t
*
in_ptr
,
ck
::
half_t
*
weight_ptr
,
ck
::
half_t
*
output_ptr
,
ck
::
half_t
*
bias_ptr
,
const
ck
::
index_t
N
,
const
ck
::
index_t
K
,
const
ck
::
index_t
C
,
const
ck
::
index_t
Y
,
const
ck
::
index_t
X
,
const
ck
::
index_t
Hi
,
const
ck
::
index_t
Wi
,
const
ck
::
index_t
Ho
,
const
ck
::
index_t
Wo
,
const
ck
::
index_t
Stride
,
const
ck
::
index_t
Dilation
,
const
ck
::
index_t
Pad
)
{
const
auto
in_desc
=
HostTensorDescriptor
(
std
::
vector
<
std
::
size_t
>
{
static_cast
<
std
::
size_t
>
(
N
),
static_cast
<
std
::
size_t
>
(
Hi
),
static_cast
<
std
::
size_t
>
(
Wi
),
static_cast
<
std
::
size_t
>
(
C
)});
const
auto
wei_desc
=
HostTensorDescriptor
(
std
::
vector
<
std
::
size_t
>
{
static_cast
<
std
::
size_t
>
(
K
),
static_cast
<
std
::
size_t
>
(
Y
),
static_cast
<
std
::
size_t
>
(
X
),
static_cast
<
std
::
size_t
>
(
C
)});
const
auto
out_desc
=
HostTensorDescriptor
(
std
::
vector
<
std
::
size_t
>
{
static_cast
<
std
::
size_t
>
(
N
),
static_cast
<
std
::
size_t
>
(
Ho
),
static_cast
<
std
::
size_t
>
(
Wo
),
static_cast
<
std
::
size_t
>
(
K
)});
const
auto
bias_desc
=
HostTensorDescriptor
(
std
::
vector
<
std
::
size_t
>
{
static_cast
<
std
::
size_t
>
(
K
)});
auto
f_k
=
[
&
](
auto
k
)
{
for
(
int
n
=
0
;
n
<
N
;
++
n
)
{
for
(
int
ho
=
0
;
ho
<
Ho
;
++
ho
)
{
for
(
int
wo
=
0
;
wo
<
Wo
;
++
wo
)
{
double
v
=
0
;
for
(
int
c
=
0
;
c
<
C
;
++
c
)
{
for
(
int
y
=
0
;
y
<
Y
;
++
y
)
{
int
hi
=
ho
*
Stride
+
y
*
Dilation
-
Pad
;
for
(
int
x
=
0
;
x
<
X
;
++
x
)
{
int
wi
=
wo
*
Stride
+
x
*
Dilation
-
Pad
;
if
(
hi
>=
0
&&
hi
<
Hi
&&
wi
>=
0
&&
wi
<
Wi
)
{
double
in
=
in_ptr
[
in_desc
.
GetOffsetFromMultiIndex
(
n
,
hi
,
wi
,
c
)];
double
wei
=
weight_ptr
[
wei_desc
.
GetOffsetFromMultiIndex
(
k
,
y
,
x
,
c
)];
v
+=
in
*
wei
;
}
}
}
}
v
+=
bias_ptr
[
bias_desc
.
GetOffsetFromMultiIndex
(
k
)];
v
=
v
>
0
?
v
:
0
;
output_ptr
[
out_desc
.
GetOffsetFromMultiIndex
(
n
,
ho
,
wo
,
k
)]
=
v
;
}
}
}
};
make_ParallelTensorFunctor
(
f_k
,
K
)(
std
::
thread
::
hardware_concurrency
());
}
template
<
int
NDimSpatial
,
typename
InDataType
,
typename
WeiDataType
,
typename
OutDataType
,
typename
InLayout
,
typename
WeiLayout
,
typename
OutLayout
>
void
profile_conv_fwd_bias_relu_atomic_add_impl
(
int
do_verification
,
int
init_method
,
bool
do_log
,
int
nrepeat
,
ck
::
index_t
N
,
ck
::
index_t
K
,
ck
::
index_t
C
,
std
::
vector
<
ck
::
index_t
>
input_spatial_lengths
,
std
::
vector
<
ck
::
index_t
>
filter_spatial_lengths
,
std
::
vector
<
ck
::
index_t
>
output_spatial_lengths
,
std
::
vector
<
ck
::
index_t
>
conv_filter_strides
,
std
::
vector
<
ck
::
index_t
>
conv_filter_dilations
,
std
::
vector
<
ck
::
index_t
>
input_left_pads
,
std
::
vector
<
ck
::
index_t
>
input_right_pads
)
{
const
ck
::
index_t
Y
=
filter_spatial_lengths
[
0
];
const
ck
::
index_t
X
=
filter_spatial_lengths
[
1
];
const
ck
::
index_t
Hi
=
input_spatial_lengths
[
0
];
const
ck
::
index_t
Wi
=
input_spatial_lengths
[
1
];
const
ck
::
index_t
Ho
=
output_spatial_lengths
[
0
];
const
ck
::
index_t
Wo
=
output_spatial_lengths
[
1
];
auto
f_host_tensor_descriptor
=
[](
std
::
size_t
N_
,
std
::
size_t
C_
,
std
::
size_t
H
,
std
::
size_t
W
,
auto
layout
)
{
if
constexpr
(
is_same
<
decltype
(
layout
),
ck
::
tensor_layout
::
convolution
::
NCHW
>::
value
||
is_same
<
decltype
(
layout
),
ck
::
tensor_layout
::
convolution
::
KCYX
>::
value
||
is_same
<
decltype
(
layout
),
ck
::
tensor_layout
::
convolution
::
NKHW
>::
value
)
{
return
HostTensorDescriptor
(
std
::
vector
<
std
::
size_t
>
({
N_
,
C_
,
H
,
W
}),
std
::
vector
<
std
::
size_t
>
({
C_
*
H
*
W
,
H
*
W
,
W
,
1
}));
}
else
if
constexpr
(
is_same
<
decltype
(
layout
),
tensor_layout
::
convolution
::
NHWC
>::
value
||
is_same
<
decltype
(
layout
),
tensor_layout
::
convolution
::
KYXC
>::
value
||
is_same
<
decltype
(
layout
),
tensor_layout
::
convolution
::
NHWK
>::
value
)
{
return
HostTensorDescriptor
(
std
::
vector
<
std
::
size_t
>
({
N_
,
C_
,
H
,
W
}),
std
::
vector
<
std
::
size_t
>
({
C_
*
H
*
W
,
1
,
W
*
C_
,
C_
}));
}
};
Tensor
<
InDataType
>
in_n_c_hi_wi
(
f_host_tensor_descriptor
(
N
,
C
,
Hi
,
Wi
,
InLayout
{}));
Tensor
<
WeiDataType
>
wei_k_c_y_x
(
f_host_tensor_descriptor
(
K
,
C
,
Y
,
X
,
WeiLayout
{}));
Tensor
<
OutDataType
>
out_n_k_ho_wo_host_result
(
f_host_tensor_descriptor
(
N
,
K
,
Ho
,
Wo
,
OutLayout
{}));
Tensor
<
OutDataType
>
out_n_k_ho_wo_device_result
(
f_host_tensor_descriptor
(
N
,
K
,
Ho
,
Wo
,
OutLayout
{}));
// bias: assume contiguous 1d vector
Tensor
<
OutDataType
>
bias_k
(
HostTensorDescriptor
(
std
::
vector
<
std
::
size_t
>
({
static_cast
<
std
::
size_t
>
(
K
)})));
std
::
cout
<<
"in_n_c_hi_wi: "
<<
in_n_c_hi_wi
.
mDesc
<<
std
::
endl
;
std
::
cout
<<
"wei_k_c_y_x: "
<<
wei_k_c_y_x
.
mDesc
<<
std
::
endl
;
std
::
cout
<<
"out_n_k_ho_wo: "
<<
out_n_k_ho_wo_host_result
.
mDesc
<<
std
::
endl
;
std
::
cout
<<
"bias_k: "
<<
bias_k
.
mDesc
<<
std
::
endl
;
switch
(
init_method
)
{
case
0
:
break
;
case
1
:
in_n_c_hi_wi
.
GenerateTensorValue
(
GeneratorTensor_2
<
InDataType
>
{
-
5
,
5
});
wei_k_c_y_x
.
GenerateTensorValue
(
GeneratorTensor_2
<
WeiDataType
>
{
-
5
,
5
});
bias_k
.
GenerateTensorValue
(
GeneratorTensor_2
<
OutDataType
>
{
-
5
,
5
});
break
;
default:
in_n_c_hi_wi
.
GenerateTensorValue
(
GeneratorTensor_3
<
InDataType
>
{
0.0
,
1.0
});
wei_k_c_y_x
.
GenerateTensorValue
(
GeneratorTensor_3
<
WeiDataType
>
{
-
0.5
,
0.5
});
bias_k
.
GenerateTensorValue
(
GeneratorTensor_3
<
OutDataType
>
{
0.0
,
1.0
});
}
using
InElementOp
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
using
WeiElementOp
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
using
OutElementOp
=
ck
::
tensor_operation
::
element_wise
::
AddRelu
;
if
(
do_verification
)
{
cpu_conv_bias_relu_atomic_add
(
in_n_c_hi_wi
.
mData
.
data
(),
wei_k_c_y_x
.
mData
.
data
(),
out_n_k_ho_wo_host_result
.
mData
.
data
(),
bias_k
.
mData
.
data
(),
N
,
K
,
C
,
Y
,
X
,
Hi
,
Wi
,
Ho
,
Wo
,
conv_filter_strides
[
0
],
conv_filter_dilations
[
0
],
input_left_pads
[
0
]);
}
DeviceMem
in_device_buf
(
sizeof
(
InDataType
)
*
in_n_c_hi_wi
.
mDesc
.
GetElementSpace
());
DeviceMem
wei_device_buf
(
sizeof
(
WeiDataType
)
*
wei_k_c_y_x
.
mDesc
.
GetElementSpace
());
DeviceMem
out_device_buf
(
sizeof
(
OutDataType
)
*
out_n_k_ho_wo_device_result
.
mDesc
.
GetElementSpace
());
DeviceMem
bias_device_buf
(
sizeof
(
OutDataType
)
*
bias_k
.
mDesc
.
GetElementSpace
());
in_device_buf
.
ToDevice
(
in_n_c_hi_wi
.
mData
.
data
());
wei_device_buf
.
ToDevice
(
wei_k_c_y_x
.
mData
.
data
());
bias_device_buf
.
ToDevice
(
bias_k
.
mData
.
data
());
using
DeviceConvFwdBiasReluPtr
=
ck
::
tensor_operation
::
device
::
DeviceConvFwdBiasActivationPtr
<
InElementOp
,
WeiElementOp
,
OutElementOp
>
;
// add device operator instances
std
::
vector
<
DeviceConvFwdBiasReluPtr
>
op_ptrs
;
if
constexpr
(
ck
::
is_same_v
<
ck
::
remove_cv_t
<
InDataType
>
,
ck
::
half_t
>
&&
ck
::
is_same_v
<
ck
::
remove_cv_t
<
WeiDataType
>
,
ck
::
half_t
>
&&
ck
::
is_same_v
<
ck
::
remove_cv_t
<
OutDataType
>
,
ck
::
half_t
>
)
{
ck
::
tensor_operation
::
device
::
device_conv2d_fwd_bias_activation_atomic_add_instance
::
add_device_conv2d_fwd_xdl_c_shuffle_bias_relu_atomic_add_nhwc_kyxc_nhwk_f16_instances
(
op_ptrs
);
}
if
(
op_ptrs
.
size
()
<=
0
)
{
throw
std
::
runtime_error
(
"wrong! no device Conv instance found"
);
}
std
::
string
best_conv_name
;
float
best_ave_time
=
0
;
float
best_tflops
=
0
;
float
best_gb_per_sec
=
0
;
// profile device Conv instances
for
(
auto
&
op_ptr
:
op_ptrs
)
{
auto
argument_ptr
=
op_ptr
->
MakeArgumentPointer
(
static_cast
<
const
InDataType
*>
(
in_device_buf
.
GetDeviceBuffer
()),
static_cast
<
const
WeiDataType
*>
(
wei_device_buf
.
GetDeviceBuffer
()),
static_cast
<
OutDataType
*>
(
out_device_buf
.
GetDeviceBuffer
()),
static_cast
<
const
OutDataType
*>
(
bias_device_buf
.
GetDeviceBuffer
()),
N
,
K
,
C
,
input_spatial_lengths
,
filter_spatial_lengths
,
output_spatial_lengths
,
conv_filter_strides
,
conv_filter_dilations
,
input_left_pads
,
input_right_pads
,
InElementOp
{},
WeiElementOp
{},
OutElementOp
{});
auto
invoker_ptr
=
op_ptr
->
MakeInvokerPointer
();
if
(
op_ptr
->
IsSupportedArgument
(
argument_ptr
.
get
()))
{
std
::
string
conv_name
=
op_ptr
->
GetTypeString
();
float
ave_time
=
invoker_ptr
->
Run
(
argument_ptr
.
get
(),
nrepeat
);
std
::
size_t
flop
=
std
::
size_t
(
2
)
*
N
*
K
*
Ho
*
Wo
*
C
*
Y
*
X
;
std
::
size_t
num_btype
=
sizeof
(
InDataType
)
*
(
N
*
C
*
Hi
*
Wi
)
+
sizeof
(
WeiDataType
)
*
(
K
*
C
*
Y
*
X
)
+
sizeof
(
OutDataType
)
*
(
N
*
K
*
Ho
*
Wo
)
+
sizeof
(
OutDataType
)
*
(
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, "
<<
conv_name
<<
std
::
endl
;
if
(
tflops
>
best_tflops
)
{
best_conv_name
=
conv_name
;
best_tflops
=
tflops
;
best_ave_time
=
ave_time
;
best_gb_per_sec
=
gb_per_sec
;
}
if
(
do_verification
)
{
out_device_buf
.
FromDevice
(
out_n_k_ho_wo_device_result
.
mData
.
data
());
ck
::
utils
::
check_err
(
out_n_k_ho_wo_device_result
.
mData
,
out_n_k_ho_wo_host_result
.
mData
);
if
(
do_log
)
{
LogRangeAsType
<
float
>
(
std
::
cout
<<
"in : "
,
in_n_c_hi_wi
.
mData
,
","
)
<<
std
::
endl
;
LogRangeAsType
<
float
>
(
std
::
cout
<<
"wei: "
,
wei_k_c_y_x
.
mData
,
","
)
<<
std
::
endl
;
LogRangeAsType
<
float
>
(
std
::
cout
<<
"out_host : "
,
out_n_k_ho_wo_host_result
.
mData
,
","
)
<<
std
::
endl
;
LogRangeAsType
<
float
>
(
std
::
cout
<<
"out_device: "
,
out_n_k_ho_wo_device_result
.
mData
,
","
)
<<
std
::
endl
;
}
}
}
}
std
::
cout
<<
"Best Perf: "
<<
best_ave_time
<<
" ms, "
<<
best_tflops
<<
" TFlops, "
<<
best_gb_per_sec
<<
" GB/s, "
<<
best_conv_name
<<
std
::
endl
;
}
}
// namespace profiler
}
// namespace ck
profiler/include/profile_conv_fwd_bias_relu_impl.hpp
View file @
4511f877
...
...
@@ -38,7 +38,7 @@ template <int NDimSpatial,
typename
InLayout
,
typename
WeiLayout
,
typename
OutLayout
>
void
profile_conv_fwd_bias_relu_impl
(
int
do_verification
,
bool
profile_conv_fwd_bias_relu_impl
(
int
do_verification
,
int
init_method
,
bool
do_log
,
int
nrepeat
,
...
...
@@ -53,6 +53,8 @@ void profile_conv_fwd_bias_relu_impl(int do_verification,
std
::
vector
<
ck
::
index_t
>
input_left_pads
,
std
::
vector
<
ck
::
index_t
>
input_right_pads
)
{
bool
pass
=
true
;
const
ck
::
index_t
Y
=
filter_spatial_lengths
[
0
];
const
ck
::
index_t
X
=
filter_spatial_lengths
[
1
];
...
...
@@ -234,8 +236,8 @@ void profile_conv_fwd_bias_relu_impl(int do_verification,
{
out_device_buf
.
FromDevice
(
out_n_k_ho_wo_device_result
.
mData
.
data
());
ck
::
utils
::
check_err
(
out_n_k_ho_wo_device_result
.
mData
,
out_n_k_ho_wo_host_result
.
mData
);
pass
=
pass
&&
ck
::
utils
::
check_err
(
out_n_k_ho_wo_device_result
.
mData
,
out_n_k_ho_wo_host_result
.
mData
);
if
(
do_log
)
{
...
...
@@ -256,6 +258,8 @@ void profile_conv_fwd_bias_relu_impl(int do_verification,
std
::
cout
<<
"Best Perf: "
<<
best_ave_time
<<
" ms, "
<<
best_tflops
<<
" TFlops, "
<<
best_gb_per_sec
<<
" GB/s, "
<<
best_conv_name
<<
std
::
endl
;
return
pass
;
}
}
// namespace profiler
...
...
profiler/include/profile_convnd_bwd_data_impl.hpp
View file @
4511f877
...
...
@@ -23,6 +23,7 @@ using DeviceConvBwdDataNoOpPtr =
DeviceConvBwdDataPtr
<
ck
::
tensor_operation
::
element_wise
::
PassThrough
,
ck
::
tensor_operation
::
element_wise
::
PassThrough
,
ck
::
tensor_operation
::
element_wise
::
PassThrough
>
;
void
add_device_conv1d_bwd_data_xdl_nwc_kxc_nwk_f32_instances
(
std
::
vector
<
DeviceConvBwdDataNoOpPtr
>&
);
void
add_device_conv1d_bwd_data_xdl_nwc_kxc_nwk_f16_instances
(
...
...
@@ -49,6 +50,7 @@ void add_device_conv3d_bwd_data_xdl_ndhwc_kzyxc_ndhwk_bf16_instances(
std
::
vector
<
DeviceConvBwdDataNoOpPtr
>&
);
void
add_device_conv3d_bwd_data_xdl_ndhwc_kzyxc_ndhwk_int8_instances
(
std
::
vector
<
DeviceConvBwdDataNoOpPtr
>&
);
}
// namespace device_conv2d_bwd_data_instance
}
// namespace device
}
// namespace tensor_operation
...
...
@@ -217,21 +219,6 @@ void get_device_conv_bwd_data_op_ptr(
}
}
template
<
typename
T
>
static
bool
check_out
(
const
Tensor
<
T
>&
ref
,
const
Tensor
<
T
>&
result
)
{
float
max_diff
=
1e-6
;
for
(
int
i
=
0
;
i
<
ref
.
mData
.
size
();
++
i
)
{
float
diff
=
std
::
abs
(
double
(
ref
.
mData
[
i
])
-
double
(
result
.
mData
[
i
]));
if
(
max_diff
<
diff
)
{
return
false
;
}
}
return
true
;
}
template
<
typename
DataType
>
void
show_data_nhwc_layout
(
Tensor
<
DataType
>&
nhwc
)
{
...
...
@@ -281,6 +268,8 @@ bool profile_convnd_bwd_data_impl(int do_verification,
const
std
::
vector
<
ck
::
index_t
>&
input_left_pads
,
const
std
::
vector
<
ck
::
index_t
>&
input_right_pads
)
{
bool
pass
=
true
;
using
InElementOp
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
using
WeiElementOp
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
using
OutElementOp
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
...
...
@@ -335,28 +324,10 @@ bool profile_convnd_bwd_data_impl(int do_verification,
out_device_buf
.
ToDevice
(
output
.
mData
.
data
());
wei_device_buf
.
ToDevice
(
weights
.
mData
.
data
());
// reset input to zero
in_device_buf
.
SetZero
();
// reference calculation
if
(
do_verification
)
{
auto
RunReference
=
[
&
](
auto
&
ref_conv
)
{
auto
ref_invoker
=
ref_conv
.
MakeInvoker
();
auto
ref_argument
=
ref_conv
.
MakeArgument
(
input_host_result
,
weights
,
output
,
conv_filter_strides
,
conv_filter_dilations
,
input_left_pads
,
input_right_pads
,
InElementOp
{},
WeiElementOp
{},
OutElementOp
{});
ref_invoker
.
Run
(
ref_argument
);
};
auto
ref_conv
=
ck
::
tensor_operation
::
host
::
ReferenceConvBwdData
<
InDataType
,
auto
ref_conv
=
ck
::
tensor_operation
::
host
::
ReferenceConvBwdData
<
InDataType
,
WeiDataType
,
OutDataType
,
AccDataType
,
...
...
@@ -364,7 +335,19 @@ bool profile_convnd_bwd_data_impl(int do_verification,
WeiElementOp
,
OutElementOp
,
NDimSpatial
>
();
RunReference
(
ref_conv
);
auto
ref_invoker
=
ref_conv
.
MakeInvoker
();
auto
ref_argument
=
ref_conv
.
MakeArgument
(
input_host_result
,
weights
,
output
,
conv_filter_strides
,
conv_filter_dilations
,
input_left_pads
,
input_right_pads
,
InElementOp
{},
WeiElementOp
{},
OutElementOp
{});
ref_invoker
.
Run
(
ref_argument
);
}
// add device Conv instances
...
...
@@ -372,10 +355,7 @@ bool profile_convnd_bwd_data_impl(int do_verification,
get_device_conv_bwd_data_op_ptr
(
InDataType
{},
WeiDataType
{},
OutDataType
{},
conv_ptrs
,
NDimSpatial
);
if
(
conv_ptrs
.
size
()
<=
0
)
{
throw
std
::
runtime_error
(
"wrong! no device Conv instance found"
);
}
std
::
cout
<<
"found "
<<
conv_ptrs
.
size
()
<<
" instances"
<<
std
::
endl
;
std
::
string
best_conv_name
;
float
best_ave_time
=
0
;
...
...
@@ -383,7 +363,6 @@ bool profile_convnd_bwd_data_impl(int do_verification,
float
best_gb_per_sec
=
0
;
// profile device Conv instances
bool
success
=
true
;
for
(
auto
&
conv_ptr
:
conv_ptrs
)
{
auto
argument_ptr
=
conv_ptr
->
MakeArgumentPointer
(
...
...
@@ -408,6 +387,9 @@ bool profile_convnd_bwd_data_impl(int do_verification,
if
(
conv_ptr
->
IsSupportedArgument
(
argument_ptr
.
get
()))
{
// re-init to zero before profiling next kernel
in_device_buf
.
SetZero
();
std
::
string
conv_name
=
conv_ptr
->
GetTypeString
();
float
ave_time
=
invoker_ptr
->
Run
(
argument_ptr
.
get
(),
nrepeat
);
...
...
@@ -436,18 +418,8 @@ bool profile_convnd_bwd_data_impl(int do_verification,
{
in_device_buf
.
FromDevice
(
input_device_result
.
mData
.
data
());
if
(
!
check_out
(
input_host_result
,
input_device_result
))
{
std
::
cout
<<
"Fail Info: "
<<
conv_ptr
->
GetTypeString
()
<<
std
::
endl
;
success
=
false
;
}
else
{
std
::
cout
<<
"Pass Info: "
<<
conv_ptr
->
GetTypeString
()
<<
std
::
endl
;
}
check_error
(
input_host_result
,
input_device_result
);
pass
=
pass
&&
ck
::
utils
::
check_err
(
input_device_result
.
mData
,
input_host_result
.
mData
);
if
(
do_log
)
{
...
...
@@ -473,8 +445,8 @@ bool profile_convnd_bwd_data_impl(int do_verification,
std
::
cout
<<
"Best Perf: "
<<
best_ave_time
<<
" ms, "
<<
best_tflops
<<
" TFlops, "
<<
best_gb_per_sec
<<
" GB/s, "
<<
best_conv_name
<<
std
::
endl
;
return
success
;
}
return
pass
;
}
}
// namespace profiler
}
// namespace ck
profiler/include/profile_convnd_fwd.hpp
deleted
100644 → 0
View file @
519b6aaf
#pragma once
namespace
ck
{
namespace
profiler
{
int
profile_convnd_fwd
(
int
argc
,
char
*
argv
[]);
}
// namespace profiler
}
// namespace ck
profiler/include/profile_gemm_bias_2d_impl.hpp
View file @
4511f877
...
...
@@ -62,7 +62,7 @@ template <typename ADataType,
typename
ALayout
,
typename
BLayout
,
typename
CLayout
>
void
profile_gemm_bias_2d_impl
(
int
do_verification
,
bool
profile_gemm_bias_2d_impl
(
int
do_verification
,
int
init_method
,
bool
do_log
,
int
nrepeat
,
...
...
@@ -75,6 +75,8 @@ void profile_gemm_bias_2d_impl(int do_verification,
float
alpha
,
float
beta
)
{
bool
pass
=
true
;
auto
f_host_tensor_descriptor
=
[](
std
::
size_t
row
,
std
::
size_t
col
,
std
::
size_t
stride
,
auto
layout
)
{
if
(
is_same
<
decltype
(
layout
),
tensor_layout
::
gemm
::
RowMajor
>::
value
)
...
...
@@ -115,9 +117,6 @@ void profile_gemm_bias_2d_impl(int do_verification,
c0_m_n
.
GenerateTensorValue
(
GeneratorTensor_3
<
C0DataType
>
{
-
0.5
,
0.5
},
num_thread
);
}
// set zero to c_device_buf
c_m_n_device_result
.
GenerateTensorValue
(
GeneratorTensor_0
<
CDataType
>
{},
num_thread
);
using
AElementOp
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
using
BElementOp
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
using
CElementOp
=
ck
::
tensor_operation
::
element_wise
::
AlphaBetaAdd
;
...
...
@@ -137,9 +136,8 @@ void profile_gemm_bias_2d_impl(int do_verification,
BElementOp
,
CElementOp
>
;
auto
ref_gemm
=
ReferenceGemmInstance
{};
auto
ref_invoker
=
ref_gemm
.
MakeInvoker
();
auto
ref_gemm
=
ReferenceGemmInstance
{};
auto
ref_invoker
=
ref_gemm
.
MakeInvoker
();
auto
ref_argument
=
ref_gemm
.
MakeArgument
(
a_m_k
,
b_k_n
,
c0_m_n
,
c_m_n_host_result
,
a_element_op
,
b_element_op
,
c_element_op
);
...
...
@@ -225,10 +223,7 @@ void profile_gemm_bias_2d_impl(int do_verification,
}
}
if
(
gemm_ptrs
.
size
()
<=
0
)
{
throw
std
::
runtime_error
(
"wrong! no device GEMM instance found"
);
}
std
::
cout
<<
"found "
<<
gemm_ptrs
.
size
()
<<
" instances"
<<
std
::
endl
;
std
::
string
best_gemm_name
;
float
best_ave_time
=
0
;
...
...
@@ -257,6 +252,9 @@ void profile_gemm_bias_2d_impl(int do_verification,
if
(
gemm_ptr
->
IsSupportedArgument
(
argument_ptr
.
get
()))
{
// re-init C to zero before profiling next kernel
c_device_buf
.
SetZero
();
std
::
string
gemm_name
=
gemm_ptr
->
GetTypeString
();
float
ave_time
=
invoker_ptr
->
Run
(
argument_ptr
.
get
(),
nrepeat
);
...
...
@@ -264,7 +262,7 @@ void profile_gemm_bias_2d_impl(int do_verification,
std
::
size_t
flop
=
std
::
size_t
(
2
)
*
M
*
N
*
K
;
std
::
size_t
num_btype
=
sizeof
(
ADataType
)
*
M
*
K
+
sizeof
(
BDataType
)
*
K
*
M
+
sizeof
(
CDataType
)
*
M
*
N
;
sizeof
(
ADataType
)
*
M
*
K
+
sizeof
(
BDataType
)
*
K
*
N
+
sizeof
(
CDataType
)
*
M
*
N
;
float
tflops
=
static_cast
<
float
>
(
flop
)
/
1.E9
/
ave_time
;
...
...
@@ -285,7 +283,8 @@ void profile_gemm_bias_2d_impl(int do_verification,
{
c_device_buf
.
FromDevice
(
c_m_n_device_result
.
mData
.
data
());
ck
::
utils
::
check_err
(
c_m_n_device_result
.
mData
,
c_m_n_host_result
.
mData
);
pass
=
pass
&&
ck
::
utils
::
check_err
(
c_m_n_device_result
.
mData
,
c_m_n_host_result
.
mData
);
if
(
do_log
)
{
...
...
@@ -301,12 +300,14 @@ void profile_gemm_bias_2d_impl(int do_verification,
}
else
{
std
::
cout
<<
"does not support this
GEMM
problem"
<<
std
::
endl
;
std
::
cout
<<
"does not support this problem"
<<
std
::
endl
;
}
}
std
::
cout
<<
"Best Perf: "
<<
best_ave_time
<<
" ms, "
<<
best_tflops
<<
" TFlops, "
<<
best_gb_per_sec
<<
" GB/s, "
<<
best_gemm_name
<<
std
::
endl
;
return
pass
;
}
}
// namespace profiler
...
...
profiler/include/profile_gemm_bias_relu_add_impl.hpp
View file @
4511f877
...
...
@@ -45,7 +45,7 @@ template <typename ADataType,
typename
ALayout
,
typename
BLayout
,
typename
CLayout
>
void
profile_gemm_bias_relu_add_impl
(
int
do_verification
,
bool
profile_gemm_bias_relu_add_impl
(
int
do_verification
,
int
init_method
,
bool
do_log
,
int
nrepeat
,
...
...
@@ -58,6 +58,8 @@ void profile_gemm_bias_relu_add_impl(int do_verification,
int
StrideC1
,
int
KBatch
=
1
)
{
bool
pass
=
true
;
auto
f_host_tensor_descriptor
=
[](
std
::
size_t
row
,
std
::
size_t
col
,
std
::
size_t
stride
,
auto
layout
)
{
if
(
is_same
<
decltype
(
layout
),
tensor_layout
::
gemm
::
RowMajor
>::
value
)
...
...
@@ -74,16 +76,13 @@ void profile_gemm_bias_relu_add_impl(int do_verification,
Tensor
<
ADataType
>
a_m_k
(
f_host_tensor_descriptor
(
M
,
K
,
StrideA
,
ALayout
{}));
Tensor
<
BDataType
>
b_k_n
(
f_host_tensor_descriptor
(
K
,
N
,
StrideB
,
BLayout
{}));
Tensor
<
CDataType
>
c_m_n_host_result
(
f_host_tensor_descriptor
(
M
,
N
,
StrideC
,
CLayout
{}));
Tensor
<
CDataType
>
c_m_n_device_result
(
f_host_tensor_descriptor
(
M
,
N
,
StrideC
,
CLayout
{}));
// c0_n[n]
Tensor
<
CDataType
>
c0_n
(
HostTensorDescriptor
(
std
::
vector
<
std
::
size_t
>
({
static_cast
<
std
::
size_t
>
(
N
)}),
std
::
vector
<
std
::
size_t
>
({
1
})));
// c1_m_n[m ,n]
Tensor
<
BDataType
>
c1_m_n
(
f_host_tensor_descriptor
(
M
,
N
,
StrideC
,
CLayout
{}));
Tensor
<
CDataType
>
c_m_n_host_result
(
f_host_tensor_descriptor
(
M
,
N
,
StrideC
,
CLayout
{}));
Tensor
<
CDataType
>
c_m_n_device_result
(
f_host_tensor_descriptor
(
M
,
N
,
StrideC
,
CLayout
{}));
std
::
cout
<<
"a_m_k: "
<<
a_m_k
.
mDesc
<<
std
::
endl
;
std
::
cout
<<
"b_k_n: "
<<
b_k_n
.
mDesc
<<
std
::
endl
;
std
::
cout
<<
"c_m_n: "
<<
c_m_n_host_result
.
mDesc
<<
std
::
endl
;
...
...
@@ -106,9 +105,6 @@ void profile_gemm_bias_relu_add_impl(int do_verification,
c1_m_n
.
GenerateTensorValue
(
GeneratorTensor_3
<
CDataType
>
{
0.0
,
1.0
});
}
// set zero to c_device_buf
c_m_n_device_result
.
GenerateTensorValue
(
GeneratorTensor_0
<
CDataType
>
{});
using
AElementOp
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
using
BElementOp
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
using
CElementOp
=
ck
::
tensor_operation
::
element_wise
::
AddReluAdd
;
...
...
@@ -230,13 +226,16 @@ void profile_gemm_bias_relu_add_impl(int do_verification,
if
(
gemm_ptr
->
IsSupportedArgument
(
argument_ptr
.
get
()))
{
// re-init C to zero before profiling next kernel
c_device_buf
.
SetZero
();
std
::
string
gemm_name
=
gemm_ptr
->
GetTypeString
();
float
ave_time
=
invoker_ptr
->
Run
(
argument_ptr
.
get
(),
nrepeat
);
std
::
size_t
flop
=
std
::
size_t
(
2
)
*
M
*
N
*
K
;
std
::
size_t
num_btype
=
sizeof
(
ADataType
)
*
M
*
K
+
sizeof
(
BDataType
)
*
K
*
M
+
std
::
size_t
num_btype
=
sizeof
(
ADataType
)
*
M
*
K
+
sizeof
(
BDataType
)
*
K
*
N
+
sizeof
(
CDataType
)
*
M
*
N
+
sizeof
(
CDataType
)
*
N
+
sizeof
(
CDataType
)
*
M
*
N
;
...
...
@@ -259,7 +258,8 @@ void profile_gemm_bias_relu_add_impl(int do_verification,
{
c_device_buf
.
FromDevice
(
c_m_n_device_result
.
mData
.
data
());
ck
::
utils
::
check_err
(
c_m_n_device_result
.
mData
,
c_m_n_host_result
.
mData
);
pass
=
pass
&&
ck
::
utils
::
check_err
(
c_m_n_device_result
.
mData
,
c_m_n_host_result
.
mData
);
if
(
do_log
)
{
...
...
@@ -276,12 +276,14 @@ void profile_gemm_bias_relu_add_impl(int do_verification,
}
else
{
std
::
cout
<<
"does not support this
GEMM
problem"
<<
std
::
endl
;
std
::
cout
<<
"does not support this problem"
<<
std
::
endl
;
}
}
std
::
cout
<<
"Best Perf: "
<<
best_ave_time
<<
" ms, "
<<
best_tflops
<<
" TFlops, "
<<
best_gb_per_sec
<<
" GB/s, "
<<
best_gemm_name
<<
std
::
endl
;
return
pass
;
}
}
// namespace profiler
...
...
profiler/include/profile_gemm_bias_relu_impl.hpp
View file @
4511f877
...
...
@@ -45,7 +45,7 @@ template <typename ADataType,
typename
ALayout
,
typename
BLayout
,
typename
CLayout
>
void
profile_gemm_bias_relu_impl
(
int
do_verification
,
bool
profile_gemm_bias_relu_impl
(
int
do_verification
,
int
init_method
,
bool
do_log
,
int
nrepeat
,
...
...
@@ -57,6 +57,8 @@ void profile_gemm_bias_relu_impl(int do_verification,
int
StrideC
,
int
KBatch
=
1
)
{
bool
pass
=
true
;
auto
f_host_tensor_descriptor
=
[](
std
::
size_t
row
,
std
::
size_t
col
,
std
::
size_t
stride
,
auto
layout
)
{
if
(
is_same
<
decltype
(
layout
),
tensor_layout
::
gemm
::
RowMajor
>::
value
)
...
...
@@ -73,13 +75,13 @@ void profile_gemm_bias_relu_impl(int do_verification,
Tensor
<
ADataType
>
a_m_k
(
f_host_tensor_descriptor
(
M
,
K
,
StrideA
,
ALayout
{}));
Tensor
<
BDataType
>
b_k_n
(
f_host_tensor_descriptor
(
K
,
N
,
StrideB
,
BLayout
{}));
Tensor
<
CDataType
>
c_m_n_host_result
(
f_host_tensor_descriptor
(
M
,
N
,
StrideC
,
CLayout
{}));
Tensor
<
CDataType
>
c_m_n_device_result
(
f_host_tensor_descriptor
(
M
,
N
,
StrideC
,
CLayout
{}));
// c0_n[n]
Tensor
<
CDataType
>
c0_n
(
HostTensorDescriptor
(
std
::
vector
<
std
::
size_t
>
({
static_cast
<
std
::
size_t
>
(
N
)}),
std
::
vector
<
std
::
size_t
>
({
1
})));
Tensor
<
CDataType
>
c_m_n_host_result
(
f_host_tensor_descriptor
(
M
,
N
,
StrideC
,
CLayout
{}));
Tensor
<
CDataType
>
c_m_n_device_result
(
f_host_tensor_descriptor
(
M
,
N
,
StrideC
,
CLayout
{}));
std
::
cout
<<
"a_m_k: "
<<
a_m_k
.
mDesc
<<
std
::
endl
;
std
::
cout
<<
"b_k_n: "
<<
b_k_n
.
mDesc
<<
std
::
endl
;
std
::
cout
<<
"c_m_n: "
<<
c_m_n_host_result
.
mDesc
<<
std
::
endl
;
...
...
@@ -100,9 +102,6 @@ void profile_gemm_bias_relu_impl(int do_verification,
c0_n
.
GenerateTensorValue
(
GeneratorTensor_3
<
CDataType
>
{
0.0
,
1.0
});
}
// set zero to c_device_buf
c_m_n_device_result
.
GenerateTensorValue
(
GeneratorTensor_0
<
CDataType
>
{},
num_thread
);
using
AElementOp
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
using
BElementOp
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
using
CElementOp
=
ck
::
tensor_operation
::
element_wise
::
AddRelu
;
...
...
@@ -238,7 +237,8 @@ void profile_gemm_bias_relu_impl(int do_verification,
{
c_device_buf
.
FromDevice
(
c_m_n_device_result
.
mData
.
data
());
ck
::
utils
::
check_err
(
c_m_n_device_result
.
mData
,
c_m_n_host_result
.
mData
);
pass
=
pass
&&
ck
::
utils
::
check_err
(
c_m_n_device_result
.
mData
,
c_m_n_host_result
.
mData
);
if
(
do_log
)
{
...
...
@@ -254,12 +254,14 @@ void profile_gemm_bias_relu_impl(int do_verification,
}
else
{
std
::
cout
<<
"does not support this
GEMM
problem"
<<
std
::
endl
;
std
::
cout
<<
"does not support this problem"
<<
std
::
endl
;
}
}
std
::
cout
<<
"Best Perf: "
<<
best_ave_time
<<
" ms, "
<<
best_tflops
<<
" TFlops, "
<<
best_gb_per_sec
<<
" GB/s, "
<<
best_gemm_name
<<
std
::
endl
;
return
pass
;
}
}
// namespace profiler
...
...
profiler/include/profile_gemm_impl.hpp
View file @
4511f877
This diff is collapsed.
Click to expand it.
Prev
1
2
3
4
Next
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