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
93ce856f
"git@developer.sourcefind.cn:gaoqiong/composable_kernel.git" did not exist on "341ad956657a0ad3501af0dcceeddb5018449de6"
Commit
93ce856f
authored
May 09, 2023
by
Bartlomiej Kocot
Browse files
Build and style fixes
parent
0245330a
Changes
8
Hide whitespace changes
Inline
Side-by-side
Showing
8 changed files
with
42 additions
and
47 deletions
+42
-47
example/26_contraction/contraction_bilinear_xdl_fp32.cpp
example/26_contraction/contraction_bilinear_xdl_fp32.cpp
+0
-1
example/26_contraction/contraction_bilinear_xdl_fp64.cpp
example/26_contraction/contraction_bilinear_xdl_fp64.cpp
+0
-1
example/26_contraction/contraction_scale_xdl_fp32.cpp
example/26_contraction/contraction_scale_xdl_fp32.cpp
+0
-1
example/26_contraction/contraction_scale_xdl_fp64.cpp
example/26_contraction/contraction_scale_xdl_fp64.cpp
+0
-1
library/include/ck/library/reference_tensor_operation/cpu/reference_contraction.hpp
.../reference_tensor_operation/cpu/reference_contraction.hpp
+5
-7
profiler/include/profiler/profile_contraction_impl.hpp
profiler/include/profiler/profile_contraction_impl.hpp
+9
-8
profiler/src/profile_contraction.cpp
profiler/src/profile_contraction.cpp
+12
-12
test/contraction/test_contraction.cpp
test/contraction/test_contraction.cpp
+16
-16
No files found.
example/26_contraction/contraction_bilinear_xdl_fp32.cpp
View file @
93ce856f
...
@@ -16,7 +16,6 @@
...
@@ -16,7 +16,6 @@
#include "ck/library/utility/host_tensor.hpp"
#include "ck/library/utility/host_tensor.hpp"
#include "ck/library/utility/host_tensor_generator.hpp"
#include "ck/library/utility/host_tensor_generator.hpp"
#include "ck/library/utility/numeric.hpp"
#include "ck/library/utility/numeric.hpp"
#include "ck/library/reference_tensor_operation/cpu/reference_contraction.hpp"
#include "ck/library/reference_tensor_operation/cpu/reference_contraction.hpp"
template
<
ck
::
index_t
...
Is
>
template
<
ck
::
index_t
...
Is
>
...
...
example/26_contraction/contraction_bilinear_xdl_fp64.cpp
View file @
93ce856f
...
@@ -16,7 +16,6 @@
...
@@ -16,7 +16,6 @@
#include "ck/library/utility/host_tensor.hpp"
#include "ck/library/utility/host_tensor.hpp"
#include "ck/library/utility/host_tensor_generator.hpp"
#include "ck/library/utility/host_tensor_generator.hpp"
#include "ck/library/utility/numeric.hpp"
#include "ck/library/utility/numeric.hpp"
#include "ck/library/reference_tensor_operation/cpu/reference_contraction.hpp"
#include "ck/library/reference_tensor_operation/cpu/reference_contraction.hpp"
template
<
ck
::
index_t
...
Is
>
template
<
ck
::
index_t
...
Is
>
...
...
example/26_contraction/contraction_scale_xdl_fp32.cpp
View file @
93ce856f
...
@@ -16,7 +16,6 @@
...
@@ -16,7 +16,6 @@
#include "ck/library/utility/host_tensor.hpp"
#include "ck/library/utility/host_tensor.hpp"
#include "ck/library/utility/host_tensor_generator.hpp"
#include "ck/library/utility/host_tensor_generator.hpp"
#include "ck/library/utility/numeric.hpp"
#include "ck/library/utility/numeric.hpp"
#include "ck/library/reference_tensor_operation/cpu/reference_contraction.hpp"
#include "ck/library/reference_tensor_operation/cpu/reference_contraction.hpp"
template
<
ck
::
index_t
...
Is
>
template
<
ck
::
index_t
...
Is
>
...
...
example/26_contraction/contraction_scale_xdl_fp64.cpp
View file @
93ce856f
...
@@ -16,7 +16,6 @@
...
@@ -16,7 +16,6 @@
#include "ck/library/utility/host_tensor.hpp"
#include "ck/library/utility/host_tensor.hpp"
#include "ck/library/utility/host_tensor_generator.hpp"
#include "ck/library/utility/host_tensor_generator.hpp"
#include "ck/library/utility/numeric.hpp"
#include "ck/library/utility/numeric.hpp"
#include "ck/library/reference_tensor_operation/cpu/reference_contraction.hpp"
#include "ck/library/reference_tensor_operation/cpu/reference_contraction.hpp"
template
<
ck
::
index_t
...
Is
>
template
<
ck
::
index_t
...
Is
>
...
...
library/include/ck/library/reference_tensor_operation/cpu/reference_contraction.hpp
View file @
93ce856f
...
@@ -11,9 +11,7 @@
...
@@ -11,9 +11,7 @@
#include "ck/tensor_operation/gpu/element/element_wise_operation.hpp"
#include "ck/tensor_operation/gpu/element/element_wise_operation.hpp"
using
Bilinear
=
ck
::
tensor_operation
::
element_wise
::
Bilinear
;
using
Bilinear
=
ck
::
tensor_operation
::
element_wise
::
Bilinear
;
using
Scale
=
ck
::
tensor_operation
::
element_wise
::
Scale
;
using
PassThrough
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
namespace
ck
{
namespace
ck
{
namespace
tensor_operation
{
namespace
tensor_operation
{
...
@@ -87,14 +85,14 @@ struct ReferenceContraction_M2_N2_K2 : public ck::tensor_operation::device::Base
...
@@ -87,14 +85,14 @@ struct ReferenceContraction_M2_N2_K2 : public ck::tensor_operation::device::Base
float
Run
(
const
Argument
&
arg
)
float
Run
(
const
Argument
&
arg
)
{
{
auto
f_ms_ns
=
[
&
](
auto
m0
,
auto
m1
,
auto
n0
,
auto
n1
)
{
auto
f_ms_ns
=
[
&
](
auto
m0
,
auto
m1
,
auto
n0
,
auto
n1
)
{
const
in
t
K0
=
arg
.
a_ms_ks_
.
mDesc
.
GetLengths
()[
2
];
const
ck
::
index_
t
K0
=
arg
.
a_ms_ks_
.
mDesc
.
GetLengths
()[
2
];
const
in
t
K1
=
arg
.
a_ms_ks_
.
mDesc
.
GetLengths
()[
3
];
const
ck
::
index_
t
K1
=
arg
.
a_ms_ks_
.
mDesc
.
GetLengths
()[
3
];
AccDataType
v_acc
=
0
;
AccDataType
v_acc
=
0
;
for
(
in
t
k0
=
0
;
k0
<
K0
;
++
k0
)
for
(
ck
::
index_
t
k0
=
0
;
k0
<
K0
;
++
k0
)
{
{
for
(
in
t
k1
=
0
;
k1
<
K1
;
++
k1
)
for
(
ck
::
index_
t
k1
=
0
;
k1
<
K1
;
++
k1
)
{
{
AccDataType
v_a
;
AccDataType
v_a
;
AccDataType
v_b
;
AccDataType
v_b
;
...
...
profiler/include/profiler/profile_contraction_impl.hpp
View file @
93ce856f
...
@@ -16,27 +16,29 @@
...
@@ -16,27 +16,29 @@
#include "ck/library/tensor_operation_instance/gpu/contraction_bilinear.hpp"
#include "ck/library/tensor_operation_instance/gpu/contraction_bilinear.hpp"
#include "ck/library/tensor_operation_instance/gpu/contraction_scale.hpp"
#include "ck/library/tensor_operation_instance/gpu/contraction_scale.hpp"
#include "ck/library/utility/check_err.hpp"
#include "ck/library/utility/check_err.hpp"
#include "ck/library/utility/device_memory.hpp"
#include "ck/library/utility/device_memory.hpp"
#include "ck/library/utility/host_tensor.hpp"
#include "ck/library/utility/host_tensor.hpp"
#include "ck/library/utility/host_tensor_generator.hpp"
#include "ck/library/utility/host_tensor_generator.hpp"
#include "ck/library/utility/literals.hpp"
#include "ck/library/utility/literals.hpp"
#include "ck/library/reference_tensor_operation/cpu/reference_contraction.hpp"
#include "ck/host_utility/io.hpp"
#include "ck/host_utility/io.hpp"
#include "ck/library/reference_tensor_operation/cpu/reference_contraction.hpp"
namespace
ck
{
namespace
ck
{
namespace
profiler
{
namespace
profiler
{
using
Bilinear
=
ck
::
tensor_operation
::
element_wise
::
Bilinear
;
using
Scale
=
ck
::
tensor_operation
::
element_wise
::
Scale
;
template
<
typename
ALayout
,
template
<
typename
ALayout
,
typename
BLayout
,
typename
BLayout
,
typename
CDLayout
,
typename
CDLayout
,
typename
DataType
,
typename
DataType
,
typename
DTupleDataType
,
typename
DTupleDataType
,
typename
CDElementOp
>
typename
CDElementOp
>
int
profile_contraction_impl
(
in
t
do_verification
,
int
profile_contraction_impl
(
ck
::
index_
t
do_verification
,
in
t
init_method
,
ck
::
index_
t
init_method
,
bool
do_log
,
bool
do_log
,
bool
time_kernel
,
bool
time_kernel
,
CDElementOp
cd_element_op
,
CDElementOp
cd_element_op
,
...
@@ -53,8 +55,6 @@ int profile_contraction_impl(int do_verification,
...
@@ -53,8 +55,6 @@ int profile_contraction_impl(int do_verification,
auto
f_host_tensor_descriptor
=
[](
const
std
::
vector
<
ck
::
index_t
>&
dims01
,
auto
f_host_tensor_descriptor
=
[](
const
std
::
vector
<
ck
::
index_t
>&
dims01
,
const
std
::
vector
<
ck
::
index_t
>&
dims23
,
const
std
::
vector
<
ck
::
index_t
>&
dims23
,
const
std
::
vector
<
ck
::
index_t
>&
strides
)
{
const
std
::
vector
<
ck
::
index_t
>&
strides
)
{
using
namespace
ck
::
literals
;
std
::
vector
<
std
::
size_t
>
dims_szt
(
dims01
.
begin
(),
dims01
.
end
());
std
::
vector
<
std
::
size_t
>
dims_szt
(
dims01
.
begin
(),
dims01
.
end
());
dims_szt
.
insert
(
dims_szt
.
end
(),
dims23
.
begin
(),
dims23
.
end
());
dims_szt
.
insert
(
dims_szt
.
end
(),
dims23
.
begin
(),
dims23
.
end
());
std
::
vector
<
std
::
size_t
>
strides_szt
(
strides
.
begin
(),
strides
.
end
());
std
::
vector
<
std
::
size_t
>
strides_szt
(
strides
.
begin
(),
strides
.
end
());
...
@@ -97,6 +97,7 @@ int profile_contraction_impl(int do_verification,
...
@@ -97,6 +97,7 @@ int profile_contraction_impl(int do_verification,
a_device_buf
.
ToDevice
(
a_m_k
.
mData
.
data
());
a_device_buf
.
ToDevice
(
a_m_k
.
mData
.
data
());
b_device_buf
.
ToDevice
(
b_k_n
.
mData
.
data
());
b_device_buf
.
ToDevice
(
b_k_n
.
mData
.
data
());
c_device_buf
.
SetZero
();
d_device_buf
.
ToDevice
(
d_m_n
.
mData
.
data
());
d_device_buf
.
ToDevice
(
d_m_n
.
mData
.
data
());
const
std
::
vector
<
index_t
>
a_ms_ks_lengths
=
{
M
[
0
],
M
[
1
],
K
[
0
],
K
[
1
]};
const
std
::
vector
<
index_t
>
a_ms_ks_lengths
=
{
M
[
0
],
M
[
1
],
K
[
0
],
K
[
1
]};
...
@@ -107,8 +108,8 @@ int profile_contraction_impl(int do_verification,
...
@@ -107,8 +108,8 @@ int profile_contraction_impl(int do_verification,
const
auto
a_element_op
=
AElementOp
{};
const
auto
a_element_op
=
AElementOp
{};
const
auto
b_element_op
=
BElementOp
{};
const
auto
b_element_op
=
BElementOp
{};
constexpr
in
t
NumDim
=
2
;
constexpr
ck
::
index_
t
NumDim
=
2
;
using
DeviceOp
=
ck
::
tensor_operation
::
device
::
DeviceContractionMultipleD
<
NumDim
,
using
DeviceOp
=
ck
::
tensor_operation
::
device
::
DeviceContractionMultipleD
<
NumDim
,
NumDim
,
NumDim
,
NumDim
,
NumDim
,
DataType
,
DataType
,
...
...
profiler/src/profile_contraction.cpp
View file @
93ce856f
...
@@ -59,10 +59,10 @@ static void print_helper_msg()
...
@@ -59,10 +59,10 @@ static void print_helper_msg()
void
collect_index_params
(
char
*
argv
[],
void
collect_index_params
(
char
*
argv
[],
std
::
vector
<
ck
::
index_t
>&
params
,
std
::
vector
<
ck
::
index_t
>&
params
,
const
in
t
from
,
const
ck
::
index_
t
from
,
const
in
t
num
)
const
ck
::
index_
t
num
)
{
{
for
(
in
t
p
=
from
;
p
<
from
+
num
;
p
++
)
for
(
ck
::
index_
t
p
=
from
;
p
<
from
+
num
;
p
++
)
params
.
push_back
(
std
::
stoi
(
argv
[
p
]));
params
.
push_back
(
std
::
stoi
(
argv
[
p
]));
}
}
...
@@ -94,19 +94,19 @@ int profile_contraction(int argc, char* argv[])
...
@@ -94,19 +94,19 @@ int profile_contraction(int argc, char* argv[])
exit
(
1
);
exit
(
1
);
}
}
const
auto
data_type
=
static_cast
<
ContractionDataType
>
(
std
::
stoi
(
argv
[
2
]));
const
auto
data_type
=
static_cast
<
ContractionDataType
>
(
std
::
stoi
(
argv
[
2
]));
const
auto
layout
=
static_cast
<
ContractionMatrixLayout
>
(
std
::
stoi
(
argv
[
3
]));
const
auto
layout
=
static_cast
<
ContractionMatrixLayout
>
(
std
::
stoi
(
argv
[
3
]));
const
bool
do_verification
=
std
::
stoi
(
argv
[
4
]);
const
bool
do_verification
=
std
::
stoi
(
argv
[
4
]);
const
in
t
init_method
=
std
::
stoi
(
argv
[
5
]);
const
ck
::
index_
t
init_method
=
std
::
stoi
(
argv
[
5
]);
const
bool
do_log
=
std
::
stoi
(
argv
[
6
]);
const
bool
do_log
=
std
::
stoi
(
argv
[
6
]);
const
bool
time_kernel
=
std
::
stoi
(
argv
[
7
]);
const
bool
time_kernel
=
std
::
stoi
(
argv
[
7
]);
const
float
alpha
=
std
::
stof
(
argv
[
8
]);
const
float
alpha
=
std
::
stof
(
argv
[
8
]);
const
float
beta
=
with_bilinear
?
std
::
stof
(
argv
[
9
])
:
0
;
const
float
beta
=
with_bilinear
?
std
::
stof
(
argv
[
9
])
:
0
;
std
::
vector
<
ck
::
index_t
>
M
;
std
::
vector
<
ck
::
index_t
>
M
;
std
::
vector
<
ck
::
index_t
>
N
;
std
::
vector
<
ck
::
index_t
>
N
;
std
::
vector
<
ck
::
index_t
>
K
;
std
::
vector
<
ck
::
index_t
>
K
;
const
in
t
dims_arg_num
=
with_bilinear
?
10
:
9
;
const
ck
::
index_
t
dims_arg_num
=
with_bilinear
?
10
:
9
;
collect_index_params
(
argv
,
M
,
dims_arg_num
,
2
);
collect_index_params
(
argv
,
M
,
dims_arg_num
,
2
);
collect_index_params
(
argv
,
N
,
dims_arg_num
+
2
,
2
);
collect_index_params
(
argv
,
N
,
dims_arg_num
+
2
,
2
);
collect_index_params
(
argv
,
K
,
dims_arg_num
+
4
,
2
);
collect_index_params
(
argv
,
K
,
dims_arg_num
+
4
,
2
);
...
...
test/contraction/test_contraction.cpp
View file @
93ce856f
...
@@ -57,13 +57,13 @@ class TestContraction : public ::testing::Test
...
@@ -57,13 +57,13 @@ class TestContraction : public ::testing::Test
{
16384
,
1024
,
32
,
1
},
{
16384
,
1024
,
32
,
1
},
{
16384
,
1024
,
32
,
1
}}};
{
16384
,
1024
,
32
,
1
}}};
std
::
vector
<
in
t
>
init_methods
=
{
0
,
1
,
2
};
std
::
vector
<
ck
::
index_
t
>
init_methods
=
{
0
,
1
,
2
};
std
::
unique_ptr
<
CDElementOp
>
cd_element_op
_ptr
;
std
::
unique_ptr
<
CDElementOp
>
p_
cd_element_op
;
void
Run
()
void
Run
()
{
{
for
(
auto
&
memory_params
:
list_of_memory_params
)
for
(
auto
&
memory_params
:
list_of_memory_params
)
{
{
for
(
const
in
t
init_method
:
init_methods
)
for
(
const
ck
::
index_
t
init_method
:
init_methods
)
{
{
bool
pass
=
bool
pass
=
ck
::
profiler
::
profile_contraction_impl
<
ALayout
,
ck
::
profiler
::
profile_contraction_impl
<
ALayout
,
...
@@ -75,7 +75,7 @@ class TestContraction : public ::testing::Test
...
@@ -75,7 +75,7 @@ class TestContraction : public ::testing::Test
init_method
,
init_method
,
false
/*do_logs*/
,
false
/*do_logs*/
,
false
/*time_kernel*/
,
false
/*time_kernel*/
,
*
cd_element_op
_ptr
,
*
p_
cd_element_op
,
memory_params
.
M
,
memory_params
.
M
,
memory_params
.
N
,
memory_params
.
N
,
memory_params
.
K
,
memory_params
.
K
,
...
@@ -104,35 +104,35 @@ using BilinearKernelTypes =
...
@@ -104,35 +104,35 @@ using BilinearKernelTypes =
std
::
tuple
<
Row
,
Col
,
Row
,
F32
,
ck
::
Tuple
<
F32
>
,
Bilinear
>
,
std
::
tuple
<
Row
,
Col
,
Row
,
F32
,
ck
::
Tuple
<
F32
>
,
Bilinear
>
,
std
::
tuple
<
Col
,
Row
,
Row
,
F32
,
ck
::
Tuple
<
F32
>
,
Bilinear
>
,
std
::
tuple
<
Col
,
Row
,
Row
,
F32
,
ck
::
Tuple
<
F32
>
,
Bilinear
>
,
std
::
tuple
<
Col
,
Col
,
Row
,
F32
,
ck
::
Tuple
<
F32
>
,
Bilinear
>
,
std
::
tuple
<
Col
,
Col
,
Row
,
F32
,
ck
::
Tuple
<
F32
>
,
Bilinear
>
,
std
::
tuple
<
Row
,
Row
,
Row
,
F64
,
Bilinear
>
,
std
::
tuple
<
Row
,
Row
,
Row
,
F64
,
ck
::
Tuple
<
F32
>
,
Bilinear
>
,
std
::
tuple
<
Row
,
Col
,
Row
,
F64
,
Bilinear
>
,
std
::
tuple
<
Row
,
Col
,
Row
,
F64
,
ck
::
Tuple
<
F32
>
,
Bilinear
>
,
std
::
tuple
<
Col
,
Row
,
Row
,
F64
,
Bilinear
>
,
std
::
tuple
<
Col
,
Row
,
Row
,
F64
,
ck
::
Tuple
<
F32
>
,
Bilinear
>
,
std
::
tuple
<
Col
,
Col
,
Row
,
F64
,
Bilinear
>>
;
std
::
tuple
<
Col
,
Col
,
Row
,
F64
,
ck
::
Tuple
<
F32
>
,
Bilinear
>>
;
using
ScaleKernelTypes
=
::
testing
::
Types
<
std
::
tuple
<
Row
,
Row
,
Row
,
F32
,
ck
::
Tuple
<>
,
Scale
>
,
using
ScaleKernelTypes
=
::
testing
::
Types
<
std
::
tuple
<
Row
,
Row
,
Row
,
F32
,
ck
::
Tuple
<>
,
Scale
>
,
std
::
tuple
<
Row
,
Col
,
Row
,
F32
,
ck
::
Tuple
<>
,
Scale
>
,
std
::
tuple
<
Row
,
Col
,
Row
,
F32
,
ck
::
Tuple
<>
,
Scale
>
,
std
::
tuple
<
Col
,
Row
,
Row
,
F32
,
ck
::
Tuple
<>
,
Scale
>
,
std
::
tuple
<
Col
,
Row
,
Row
,
F32
,
ck
::
Tuple
<>
,
Scale
>
,
std
::
tuple
<
Col
,
Col
,
Row
,
F32
,
ck
::
Tuple
<>
,
Scale
>
,
std
::
tuple
<
Col
,
Col
,
Row
,
F32
,
ck
::
Tuple
<>
,
Scale
>
,
std
::
tuple
<
Row
,
Row
,
Row
,
F64
,
Scale
>
,
std
::
tuple
<
Row
,
Row
,
Row
,
F64
,
ck
::
Tuple
<>
,
Scale
>
,
std
::
tuple
<
Row
,
Col
,
Row
,
F64
,
Scale
>
,
std
::
tuple
<
Row
,
Col
,
Row
,
F64
,
ck
::
Tuple
<>
,
Scale
>
,
std
::
tuple
<
Col
,
Row
,
Row
,
F64
,
Scale
>
,
std
::
tuple
<
Col
,
Row
,
Row
,
F64
,
ck
::
Tuple
<>
,
Scale
>
,
std
::
tuple
<
Col
,
Col
,
Row
,
F64
,
Scale
>>
;
std
::
tuple
<
Col
,
Col
,
Row
,
F64
,
ck
::
Tuple
<>
,
Scale
>>
;
TYPED_TEST_SUITE
(
TestContractionBilinear
,
BilinearKernelTypes
);
TYPED_TEST_SUITE
(
TestContractionBilinear
,
BilinearKernelTypes
);
TYPED_TEST_SUITE
(
TestContractionScale
,
ScaleKernelTypes
);
TYPED_TEST_SUITE
(
TestContractionScale
,
ScaleKernelTypes
);
TYPED_TEST
(
TestContractionBilinear
,
bilinear
)
TYPED_TEST
(
TestContractionBilinear
,
bilinear
)
{
{
this
->
cd_element_op
_ptr
=
std
::
make_unique
<
Bilinear
>
(
1.
f
,
1.
f
);
this
->
p_
cd_element_op
=
std
::
make_unique
<
Bilinear
>
(
1.
f
,
1.
f
);
this
->
Run
();
this
->
Run
();
this
->
cd_element_op
_ptr
=
std
::
make_unique
<
Bilinear
>
(
0.5
f
,
0.5
f
);
this
->
p_
cd_element_op
=
std
::
make_unique
<
Bilinear
>
(
0.5
f
,
0.5
f
);
this
->
Run
();
this
->
Run
();
}
}
TYPED_TEST
(
TestContractionScale
,
scale
)
TYPED_TEST
(
TestContractionScale
,
scale
)
{
{
this
->
cd_element_op
_ptr
=
std
::
make_unique
<
Scale
>
(
1.
f
);
this
->
p_
cd_element_op
=
std
::
make_unique
<
Scale
>
(
1.
f
);
this
->
Run
();
this
->
Run
();
this
->
cd_element_op
_ptr
=
std
::
make_unique
<
Scale
>
(
0.5
f
);
this
->
p_
cd_element_op
=
std
::
make_unique
<
Scale
>
(
0.5
f
);
this
->
Run
();
this
->
Run
();
}
}
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