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
e84c2a33
Unverified
Commit
e84c2a33
authored
Oct 31, 2023
by
Rostyslav Geyyer
Committed by
GitHub
Oct 31, 2023
Browse files
Merge branch 'develop' into lwpck-987
parents
ca47e0c0
db4461c1
Changes
44
Show whitespace changes
Inline
Side-by-side
Showing
4 changed files
with
100 additions
and
26 deletions
+100
-26
script/redis-cli.conf
script/redis-cli.conf
+10
-0
script/sccache_wrapper.sh
script/sccache_wrapper.sh
+56
-0
test/conv_tensor_rearrange/test_conv_tensor_rearrange.cpp
test/conv_tensor_rearrange/test_conv_tensor_rearrange.cpp
+26
-20
test/conv_tensor_rearrange/test_conv_tensor_rearrange_interface.cpp
...tensor_rearrange/test_conv_tensor_rearrange_interface.cpp
+8
-6
No files found.
script/redis-cli.conf
0 → 100644
View file @
e84c2a33
fips
=
no
setuid
=
root
setgid
=
root
pid
= /
var
/
run
/
stunnel
.
pid
debug
=
7
options
=
NO_SSLv2
options
=
NO_SSLv3
[
redis
-
cli
]
client
=
yes
accept
=
127
.
0
.
0
.
1
:
6379
script/sccache_wrapper.sh
0 → 100755
View file @
e84c2a33
#!/bin/bash
set
-e
COMPILERS_HASH_DIR
=
${
COMPILERS_HASH_DIR
:-
"/tmp/.sccache"
}
SCCACHE_EXTRAFILES
=
${
SCCACHE_EXTRAFILES
:-
"
${
COMPILERS_HASH_DIR
}
/rocm_compilers_hash_file"
}
SCCACHE_BIN
=
${
SCCACHE_BIN
:-
"
${
SCCACHE_INSTALL_LOCATION
}
/sccache"
}
ENFORCE_REDIS
=
"false"
while
[
"
$1
"
!=
""
]
;
do
case
$1
in
--enforce_redis
)
shift
;
ENFORCE_REDIS
=
"true"
;;
--no-hipcc
)
shift
;;
*
)
break
;;
esac
done
setup_rocm_compilers_hash_file
()
{
mkdir
-p
"
$COMPILERS_HASH_DIR
"
HIPCC_MD5
=
"
$(
md5sum
"
${
ROCM_PATH
}
/bin/hipcc"
)
"
pushd
"
${
ROCM_PATH
}
/amdgcn/bitcode"
DEVICELIBS_BITCODES_MD5
=
"
$(
find
.
-type
f
-exec
md5sum
{}
\;
|
sort
|
md5sum
)
"
popd
HIPCC_HASH_VALUE
=
"
${
HIPCC_MD5
%% *
}
"
DEVICELIBS_BITCODES_HASH_VALUE
=
"
${
DEVICELIBS_BITCODES_MD5
%% *
}
"
# MD5 checksums of clang and clang-offload-bundler cannot be used since they will keep changing
# if the ROCM_PATH changes, ie; for every mainline build.
# This is because ROCM_PATH gets encoded into the clang/clang-offload-bundler binaries as part
# of RPATH.
# The versions themselves contain the commit hash of the compiler repo at the time of building.
# Hence, this should be a viable alternative to using the binary checksum itself.
CLANG_VERSION
=
"
$(
"
${
ROCM_PATH
}
/llvm/bin/clang"
--version
|
head
-n
1
)
"
CLANG_OFFLOAD_BUNDLER_VERSION
=
"
$(
"
${
ROCM_PATH
}
/llvm/bin/clang-offload-bundler"
--version
|
head
-n
1
)
"
printf
'%s: %s\n'
'clang version'
"
${
CLANG_VERSION
}
"
|
tee
-a
"
$SCCACHE_EXTRAFILES
"
printf
'%s: %s\n'
'clang-offload-bundler version'
"
${
CLANG_OFFLOAD_BUNDLER_VERSION
}
"
|
tee
-a
"
$SCCACHE_EXTRAFILES
"
printf
'%s: %s\n'
'hipcc md5sum'
"
${
HIPCC_HASH_VALUE
}
"
|
tee
-a
"
$SCCACHE_EXTRAFILES
"
printf
'%s: %s\n'
'devicelibs bitcode md5sum'
"
${
DEVICELIBS_BITCODES_HASH_VALUE
}
"
|
tee
-a
"
$SCCACHE_EXTRAFILES
"
echo
"sccache-wrapper: compilers hash file set up at
${
SCCACHE_EXTRAFILES
}
"
cat
"
$SCCACHE_EXTRAFILES
"
}
if
[
"
${
ENFORCE_REDIS
}
"
==
"true"
]
;
then
if
[
-z
"
${
SCCACHE_REDIS
}
"
]
;
then
echo
"SCCACHE_REDIS not set. Not wrapping compilers with sccache."
exit
10
else
response
=
$(
redis-cli
-u
${
SCCACHE_REDIS
}
ping
)
||
true
if
[
"
${
response
}
"
!=
"PONG"
]
;
then
echo
"Redis server unreachable. Not wrapping compilers with sccache."
exit
20
fi
fi
fi
setup_rocm_compilers_hash_file
$SCCACHE_BIN
--version
$SCCACHE_BIN
--start-server
test/conv_tensor_rearrange/test_conv_tensor_rearrange.cpp
View file @
e84c2a33
...
@@ -45,14 +45,20 @@ class TestConvTensorRearrange : public ::testing::Test
...
@@ -45,14 +45,20 @@ class TestConvTensorRearrange : public ::testing::Test
using
namespace
ck
::
tensor_layout
::
convolution
;
using
namespace
ck
::
tensor_layout
::
convolution
;
using
namespace
ck
::
conv_tensor_rearrange_op
;
using
namespace
ck
::
conv_tensor_rearrange_op
;
using
KernelTypes1d
=
using
KernelTypes1d
=
::
testing
::
Types
<
std
::
tuple
<
GNWC
,
ImageToColumn
>
,
::
testing
::
Types
<
std
::
tuple
<
GNWC
,
ImageToColumn
>
,
std
::
tuple
<
GNWC
,
ColumnToImage
>>
;
std
::
tuple
<
GNWC
,
ColumnToImage
>
,
std
::
tuple
<
NWGC
,
ImageToColumn
>
,
std
::
tuple
<
NWGC
,
ColumnToImage
>>
;
using
KernelTypes2d
=
using
KernelTypes2d
=
::
testing
::
Types
<
std
::
tuple
<
GNHWC
,
ImageToColumn
>
,
::
testing
::
Types
<
std
::
tuple
<
GNHWC
,
ImageToColumn
>
,
std
::
tuple
<
GNHWC
,
ColumnToImage
>>
;
std
::
tuple
<
GNHWC
,
ColumnToImage
>
,
std
::
tuple
<
NHWGC
,
ImageToColumn
>
,
std
::
tuple
<
NHWGC
,
ColumnToImage
>>
;
using
KernelTypes3d
=
using
KernelTypes3d
=
::
testing
::
Types
<
std
::
tuple
<
GNDHWC
,
ImageToColumn
>
,
::
testing
::
Types
<
std
::
tuple
<
GNDHWC
,
ImageToColumn
>
,
std
::
tuple
<
GNDHWC
,
ColumnToImage
>>
;
std
::
tuple
<
GNDHWC
,
ColumnToImage
>
,
std
::
tuple
<
NDHWGC
,
ImageToColumn
>
,
std
::
tuple
<
NDHWGC
,
ColumnToImage
>>
;
template
<
typename
Tuple
>
template
<
typename
Tuple
>
class
TestConvTensorRearrange1d
:
public
TestConvTensorRearrange
<
Tuple
>
class
TestConvTensorRearrange1d
:
public
TestConvTensorRearrange
<
Tuple
>
...
@@ -77,16 +83,16 @@ TYPED_TEST(TestConvTensorRearrange1d, Test1D)
...
@@ -77,16 +83,16 @@ TYPED_TEST(TestConvTensorRearrange1d, Test1D)
{
{
this
->
conv_params
.
clear
();
this
->
conv_params
.
clear
();
this
->
conv_params
.
push_back
({
1
,
1
,
4
,
1
,
192
,
{
3
},
{
28
},
{
1
},
{
1
},
{
1
},
{
1
}});
this
->
conv_params
.
push_back
({
1
,
2
,
4
,
1
,
192
,
{
3
},
{
28
},
{
1
},
{
1
},
{
1
},
{
1
}});
this
->
conv_params
.
push_back
({
1
,
1
,
64
,
1
,
64
,
{
3
},
{
14
},
{
1
},
{
1
},
{
1
},
{
1
}});
this
->
conv_params
.
push_back
({
1
,
2
,
64
,
1
,
64
,
{
3
},
{
14
},
{
1
},
{
1
},
{
1
},
{
1
}});
this
->
conv_params
.
push_back
({
1
,
1
,
64
,
1
,
64
,
{
1
},
{
7
},
{
3
},
{
1
},
{
0
},
{
0
}});
this
->
conv_params
.
push_back
({
1
,
2
,
64
,
1
,
64
,
{
1
},
{
7
},
{
3
},
{
1
},
{
0
},
{
0
}});
this
->
conv_params
.
push_back
({
1
,
1
,
64
,
1
,
64
,
{
1
},
{
3
},
{
1
},
{
1
},
{
0
},
{
0
}});
this
->
conv_params
.
push_back
({
1
,
2
,
64
,
1
,
64
,
{
1
},
{
3
},
{
1
},
{
1
},
{
0
},
{
0
}});
// ScalarPerVector should be 1
// ScalarPerVector should be 1
this
->
conv_params
.
push_back
({
1
,
1
,
4
,
1
,
1
,
{
3
},
{
28
},
{
1
},
{
1
},
{
1
},
{
1
}});
this
->
conv_params
.
push_back
({
1
,
2
,
4
,
1
,
1
,
{
3
},
{
28
},
{
1
},
{
1
},
{
1
},
{
1
}});
// stride != 1
// stride != 1
this
->
conv_params
.
push_back
({
1
,
1
,
1
,
1
,
4
,
{
3
},
{
28
},
{
2
},
{
1
},
{
1
},
{
1
}});
this
->
conv_params
.
push_back
({
1
,
2
,
1
,
1
,
4
,
{
3
},
{
28
},
{
2
},
{
1
},
{
1
},
{
1
}});
// dilation != 1
// dilation != 1
this
->
conv_params
.
push_back
({
1
,
1
,
1
,
1
,
4
,
{
3
},
{
28
},
{
1
},
{
2
},
{
1
},
{
1
}});
this
->
conv_params
.
push_back
({
1
,
2
,
1
,
1
,
4
,
{
3
},
{
28
},
{
1
},
{
2
},
{
1
},
{
1
}});
#ifdef CK_ENABLE_FP32
#ifdef CK_ENABLE_FP32
this
->
template
Run
<
1
,
float
,
float
>();
this
->
template
Run
<
1
,
float
,
float
>();
#endif
#endif
...
@@ -106,13 +112,13 @@ TYPED_TEST(TestConvTensorRearrange2d, Test2D)
...
@@ -106,13 +112,13 @@ TYPED_TEST(TestConvTensorRearrange2d, Test2D)
this
->
conv_params
.
clear
();
this
->
conv_params
.
clear
();
this
->
conv_params
.
push_back
(
this
->
conv_params
.
push_back
(
{
2
,
1
,
4
,
1
,
192
,
{
3
,
3
},
{
28
,
28
},
{
1
,
1
},
{
1
,
1
},
{
1
,
1
},
{
1
,
1
}});
{
2
,
2
,
4
,
1
,
192
,
{
3
,
3
},
{
28
,
28
},
{
1
,
1
},
{
1
,
1
},
{
1
,
1
},
{
1
,
1
}});
this
->
conv_params
.
push_back
(
this
->
conv_params
.
push_back
(
{
2
,
1
,
64
,
1
,
64
,
{
3
,
3
},
{
14
,
14
},
{
1
,
1
},
{
1
,
1
},
{
1
,
1
},
{
1
,
1
}});
{
2
,
2
,
64
,
1
,
64
,
{
3
,
3
},
{
14
,
14
},
{
1
,
1
},
{
1
,
1
},
{
1
,
1
},
{
1
,
1
}});
this
->
conv_params
.
push_back
({
2
,
1
,
64
,
1
,
64
,
{
1
,
1
},
{
7
,
7
},
{
3
,
3
},
{
1
,
1
},
{
0
,
0
},
{
0
,
0
}});
this
->
conv_params
.
push_back
({
2
,
1
,
64
,
1
,
64
,
{
1
,
1
},
{
7
,
7
},
{
3
,
3
},
{
1
,
1
},
{
0
,
0
},
{
0
,
0
}});
this
->
conv_params
.
push_back
({
2
,
1
,
64
,
1
,
64
,
{
1
,
1
},
{
3
,
3
},
{
1
,
1
},
{
1
,
1
},
{
0
,
0
},
{
0
,
0
}});
this
->
conv_params
.
push_back
({
2
,
1
,
64
,
1
,
64
,
{
1
,
1
},
{
3
,
3
},
{
1
,
1
},
{
1
,
1
},
{
0
,
0
},
{
0
,
0
}});
this
->
conv_params
.
push_back
(
this
->
conv_params
.
push_back
(
{
2
,
1
,
64
,
1
,
64
,
{
3
,
3
},
{
28
,
28
},
{
2
,
2
},
{
2
,
2
},
{
1
,
1
},
{
1
,
1
}});
{
2
,
2
,
64
,
1
,
64
,
{
3
,
3
},
{
28
,
28
},
{
2
,
2
},
{
2
,
2
},
{
1
,
1
},
{
1
,
1
}});
#ifdef CK_ENABLE_FP32
#ifdef CK_ENABLE_FP32
this
->
template
Run
<
2
,
float
,
float
>();
this
->
template
Run
<
2
,
float
,
float
>();
#endif
#endif
...
@@ -131,13 +137,13 @@ TYPED_TEST(TestConvTensorRearrange3d, Test3D)
...
@@ -131,13 +137,13 @@ TYPED_TEST(TestConvTensorRearrange3d, Test3D)
{
{
this
->
conv_params
.
clear
();
this
->
conv_params
.
clear
();
this
->
conv_params
.
push_back
(
this
->
conv_params
.
push_back
(
{
3
,
1
,
16
,
1
,
64
,
{
1
,
1
,
1
},
{
7
,
7
,
7
},
{
2
,
2
,
2
},
{
3
,
3
,
3
},
{
0
,
0
,
0
},
{
0
,
0
,
0
}});
{
3
,
2
,
16
,
1
,
64
,
{
1
,
1
,
1
},
{
7
,
7
,
7
},
{
2
,
2
,
2
},
{
3
,
3
,
3
},
{
0
,
0
,
0
},
{
0
,
0
,
0
}});
this
->
conv_params
.
push_back
(
this
->
conv_params
.
push_back
(
{
3
,
1
,
2
,
1
,
64
,
{
3
,
3
,
3
},
{
14
,
14
,
3
},
{
1
,
1
,
1
},
{
1
,
1
,
1
},
{
1
,
1
,
1
},
{
1
,
1
,
1
}});
{
3
,
2
,
2
,
1
,
64
,
{
3
,
3
,
3
},
{
14
,
14
,
3
},
{
1
,
1
,
1
},
{
1
,
1
,
1
},
{
1
,
1
,
1
},
{
1
,
1
,
1
}});
this
->
conv_params
.
push_back
(
this
->
conv_params
.
push_back
(
{
3
,
1
,
32
,
1
,
64
,
{
1
,
1
,
1
},
{
3
,
3
,
3
},
{
1
,
1
,
1
},
{
1
,
1
,
1
},
{
0
,
0
,
0
},
{
0
,
0
,
0
}});
{
3
,
2
,
32
,
1
,
64
,
{
1
,
1
,
1
},
{
3
,
3
,
3
},
{
1
,
1
,
1
},
{
1
,
1
,
1
},
{
0
,
0
,
0
},
{
0
,
0
,
0
}});
this
->
conv_params
.
push_back
(
this
->
conv_params
.
push_back
(
{
3
,
1
,
64
,
1
,
64
,
{
3
,
3
,
3
},
{
14
,
14
,
14
},
{
2
,
2
,
2
},
{
2
,
2
,
2
},
{
1
,
1
,
1
},
{
1
,
1
,
1
}});
{
3
,
2
,
64
,
1
,
64
,
{
3
,
3
,
3
},
{
14
,
14
,
14
},
{
2
,
2
,
2
},
{
2
,
2
,
2
},
{
1
,
1
,
1
},
{
1
,
1
,
1
}});
#ifdef CK_ENABLE_FP32
#ifdef CK_ENABLE_FP32
this
->
template
Run
<
3
,
float
,
float
>();
this
->
template
Run
<
3
,
float
,
float
>();
#endif
#endif
...
...
test/conv_tensor_rearrange/test_conv_tensor_rearrange_interface.cpp
View file @
e84c2a33
...
@@ -53,7 +53,7 @@ class TestConvTensorRearrangeInterface : public ::testing::Test
...
@@ -53,7 +53,7 @@ class TestConvTensorRearrangeInterface : public ::testing::Test
template
<
typename
ConvTensorRearrangeOp
>
template
<
typename
ConvTensorRearrangeOp
>
bool
Run
()
bool
Run
()
{
{
const
auto
G
=
conv_param
.
G_
;
const
auto
N
=
conv_param
.
N_
;
const
auto
N
=
conv_param
.
N_
;
const
auto
C
=
conv_param
.
C_
;
const
auto
C
=
conv_param
.
C_
;
const
auto
FakeC
=
const
auto
FakeC
=
...
@@ -71,13 +71,13 @@ class TestConvTensorRearrangeInterface : public ::testing::Test
...
@@ -71,13 +71,13 @@ class TestConvTensorRearrangeInterface : public ::testing::Test
const
auto
image_desc
=
const
auto
image_desc
=
ck
::
utils
::
conv
::
make_input_host_tensor_descriptor_g_n_c_wis_packed
<
ImLayout
>
(
ck
::
utils
::
conv
::
make_input_host_tensor_descriptor_g_n_c_wis_packed
<
ImLayout
>
(
conv_param
);
conv_param
);
const
auto
gemm_desc
=
HostTensorDescriptor
({
NDoHoWo
,
CZYX
});
const
auto
gemm_desc
=
HostTensorDescriptor
({
G
,
NDoHoWo
,
CZYX
});
std
::
array
<
ck
::
index_t
,
NDimSpatial
>
input_spatial_lengths
{};
std
::
array
<
ck
::
index_t
,
NDimSpatial
>
input_spatial_lengths
{};
std
::
array
<
ck
::
index_t
,
NDimSpatial
>
filter_spatial_lengths
{};
std
::
array
<
ck
::
index_t
,
NDimSpatial
>
filter_spatial_lengths
{};
std
::
array
<
ck
::
index_t
,
NDimSpatial
>
output_spatial_lengths
{};
std
::
array
<
ck
::
index_t
,
NDimSpatial
>
output_spatial_lengths
{};
std
::
array
<
ck
::
index_t
,
NDimSpatial
+
3
>
input_g_n_c_wis_strides
{};
std
::
array
<
ck
::
index_t
,
NDimSpatial
+
3
>
input_g_n_c_wis_strides
{};
std
::
array
<
ck
::
index_t
,
2
>
output_m_k_strides
{};
std
::
array
<
ck
::
index_t
,
3
>
output_
g_
m_k_strides
{};
std
::
array
<
ck
::
index_t
,
NDimSpatial
>
conv_filter_strides
{};
std
::
array
<
ck
::
index_t
,
NDimSpatial
>
conv_filter_strides
{};
std
::
array
<
ck
::
index_t
,
NDimSpatial
>
conv_filter_dilations
{};
std
::
array
<
ck
::
index_t
,
NDimSpatial
>
conv_filter_dilations
{};
std
::
array
<
ck
::
index_t
,
NDimSpatial
>
input_left_pads
{};
std
::
array
<
ck
::
index_t
,
NDimSpatial
>
input_left_pads
{};
...
@@ -89,7 +89,7 @@ class TestConvTensorRearrangeInterface : public ::testing::Test
...
@@ -89,7 +89,7 @@ class TestConvTensorRearrangeInterface : public ::testing::Test
copy
(
conv_param
.
filter_spatial_lengths_
,
filter_spatial_lengths
);
copy
(
conv_param
.
filter_spatial_lengths_
,
filter_spatial_lengths
);
copy
(
conv_param
.
output_spatial_lengths_
,
output_spatial_lengths
);
copy
(
conv_param
.
output_spatial_lengths_
,
output_spatial_lengths
);
copy
(
image_desc
.
GetStrides
(),
input_g_n_c_wis_strides
);
copy
(
image_desc
.
GetStrides
(),
input_g_n_c_wis_strides
);
copy
(
gemm_desc
.
GetStrides
(),
output_m_k_strides
);
copy
(
gemm_desc
.
GetStrides
(),
output_
g_
m_k_strides
);
copy
(
conv_param
.
conv_filter_strides_
,
conv_filter_strides
);
copy
(
conv_param
.
conv_filter_strides_
,
conv_filter_strides
);
copy
(
conv_param
.
conv_filter_dilations_
,
conv_filter_dilations
);
copy
(
conv_param
.
conv_filter_dilations_
,
conv_filter_dilations
);
copy
(
conv_param
.
input_left_pads_
,
input_left_pads
);
copy
(
conv_param
.
input_left_pads_
,
input_left_pads
);
...
@@ -100,13 +100,14 @@ class TestConvTensorRearrangeInterface : public ::testing::Test
...
@@ -100,13 +100,14 @@ class TestConvTensorRearrangeInterface : public ::testing::Test
auto
img2col
=
DeviceImgToColInstance
{};
auto
img2col
=
DeviceImgToColInstance
{};
auto
argument
=
img2col
.
MakeArgument
(
nullptr
,
auto
argument
=
img2col
.
MakeArgument
(
nullptr
,
nullptr
,
nullptr
,
G
,
N
,
N
,
IsCPacked
?
C
:
FakeC
,
IsCPacked
?
C
:
FakeC
,
input_spatial_lengths
,
input_spatial_lengths
,
filter_spatial_lengths
,
filter_spatial_lengths
,
output_spatial_lengths
,
output_spatial_lengths
,
input_g_n_c_wis_strides
,
input_g_n_c_wis_strides
,
output_m_k_strides
,
output_
g_
m_k_strides
,
conv_filter_strides
,
conv_filter_strides
,
conv_filter_dilations
,
conv_filter_dilations
,
input_left_pads
,
input_left_pads
,
...
@@ -119,13 +120,14 @@ class TestConvTensorRearrangeInterface : public ::testing::Test
...
@@ -119,13 +120,14 @@ class TestConvTensorRearrangeInterface : public ::testing::Test
auto
col2img
=
DeviceColToimgInstance
{};
auto
col2img
=
DeviceColToimgInstance
{};
auto
argument
=
col2img
.
MakeArgument
(
nullptr
,
auto
argument
=
col2img
.
MakeArgument
(
nullptr
,
nullptr
,
nullptr
,
G
,
N
,
N
,
IsCPacked
?
C
:
FakeC
,
IsCPacked
?
C
:
FakeC
,
input_spatial_lengths
,
input_spatial_lengths
,
filter_spatial_lengths
,
filter_spatial_lengths
,
output_spatial_lengths
,
output_spatial_lengths
,
input_g_n_c_wis_strides
,
input_g_n_c_wis_strides
,
output_m_k_strides
,
output_
g_
m_k_strides
,
conv_filter_strides
,
conv_filter_strides
,
conv_filter_dilations
,
conv_filter_dilations
,
input_left_pads
,
input_left_pads
,
...
...
Prev
1
2
3
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