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
4f940d01
Commit
4f940d01
authored
Mar 02, 2022
by
ltqin
Browse files
first version that can run
parent
baf405cd
Changes
2
Show whitespace changes
Inline
Side-by-side
Showing
2 changed files
with
41 additions
and
36 deletions
+41
-36
example/14_conv2d_backward_weight_xdl/README.md
example/14_conv2d_backward_weight_xdl/README.md
+15
-14
example/14_conv2d_backward_weight_xdl/main.cpp
example/14_conv2d_backward_weight_xdl/main.cpp
+26
-22
No files found.
example/14_conv2d_backward_weight_xdl/README.md
View file @
4f940d01
# Instructions for ```conv2d_
fwd
_xdl``` Example
# Instructions for ```conv2d_
wrw
_xdl``` Example
## Docker script
## Docker script
```
bash
```
bash
...
@@ -13,7 +13,7 @@ rocm/tensorflow:rocm4.3.1-tf2.6-dev \
...
@@ -13,7 +13,7 @@ rocm/tensorflow:rocm4.3.1-tf2.6-dev \
/bin/bash
/bin/bash
```
```
## Build ```conv2d_
fwd
_xdl```
## Build ```conv2d_
wrw
_xdl```
```
bash
```
bash
mkdir
build
&&
cd
build
mkdir
build
&&
cd
build
```
```
...
@@ -30,28 +30,29 @@ cmake \
...
@@ -30,28 +30,29 @@ cmake \
```
```
```
bash
```
bash
make
-j
conv2d_
fwd
_xdl
make
-j
conv2d_
wrw
_xdl
```
```
## Run ```conv2d_
fwd
_xdl```
## Run ```conv2d_
wrw
_xdl```
```
bash
```
bash
#arg1: verification (0=no, 1=yes)
#arg1: verification (0=no, 1=yes)
#arg2: initialization (0=no init, 1=integer value, 2=decimal value)
#arg2: initialization (0=no init, 1=integer value, 2=decimal value)
#arg3: run kernel # of times (>1)
#arg3: run kernel # of times (>1)
#arg4 to 18: N, K, C, Y, X, Hi, Wi, Sy, Sx, Dy, Dx, LeftPy, LeftPx, RightPy, RightPx
#arg4: is show log (0=no, 1=yes)
./example/conv2d_fwd_xdl 0 1 5
#arg5 to 19: N, K, C, Y, X, Hi, Wi, Sy, Sx, Dy, Dx, LeftPy, LeftPx, RightPy, RightPx, split-k
./example/conv2d_fwd_xdl 0 1 5 1
```
```
Result
(MI100 @ 1087Mhz, 133.5TFlops peak FP16)
Result
```
```
in_n_c_hi_wi: dim 4, lengths {128, 1
9
2, 71, 71}, strides {
967872, 1, 13632
, 1
9
2}
in_n_c_hi_wi: dim 4, lengths {128, 12
8
, 71, 71}, strides {
645248, 1, 9088
, 12
8
}
wei_k_c_y_x: dim 4, lengths {256, 1
9
2, 3, 3}, strides {1
728
, 1,
576
, 1
9
2}
wei_k_c_y_x: dim 4, lengths {256, 12
8
, 3, 3}, strides {1
152
, 1,
384
, 12
8
}
out_n_k_ho_wo: dim 4, lengths {128, 256, 36, 36}, strides {331776, 1, 9216, 256}
out_n_k_ho_wo: dim 4, lengths {128, 256, 36, 36}, strides {331776, 1, 9216, 256}
arg.a_grid_desc_k0_m_k1_{
216, 165888, 8
}
arg.a_grid_desc_
kbatch_
k0_m_k1_{
1, 20736, 256
}
arg.b_grid_desc_k0_n_k1_{
216, 256, 8
}
arg.b_grid_desc_
kbatch_
k0_n_k1_{
1, 20736, 1152
}
arg.c_grid_desc_m_n_{
165888, 256
}
arg.c_grid_desc_m_n_{
256, 1152
}
launch_and_time_kernel: grid_dim {1
296
, 1, 1}, block_dim {256, 1, 1}
launch_and_time_kernel: grid_dim {1
8
, 1, 1}, block_dim {256, 1, 1}
Warm up
Warm up
Start running 5 times...
Start running 5 times...
Perf: 1
.43206 ms, 102.486
TFlops, 2
32.947
GB/s
Perf: 1
2.0997 ms, 8.08653
TFlops, 2
0.7201
GB/s
```
```
example/14_conv2d_backward_weight_xdl/main.cpp
View file @
4f940d01
...
@@ -78,6 +78,7 @@ int main(int argc, char* argv[])
...
@@ -78,6 +78,7 @@ int main(int argc, char* argv[])
bool
do_verification
=
0
;
bool
do_verification
=
0
;
int
init_method
=
0
;
int
init_method
=
0
;
int
nrepeat
=
5
;
int
nrepeat
=
5
;
int
do_log
=
0
;
// Conv shape
// Conv shape
ck
::
index_t
N
=
128
;
ck
::
index_t
N
=
128
;
...
@@ -97,42 +98,45 @@ int main(int argc, char* argv[])
...
@@ -97,42 +98,45 @@ int main(int argc, char* argv[])
ck
::
index_t
in_right_pad_w
=
1
;
ck
::
index_t
in_right_pad_w
=
1
;
ck
::
index_t
split_k
=
1
;
ck
::
index_t
split_k
=
1
;
if
(
argc
==
4
)
if
(
argc
==
5
)
{
{
do_verification
=
std
::
stoi
(
argv
[
1
]);
do_verification
=
std
::
stoi
(
argv
[
1
]);
init_method
=
std
::
stoi
(
argv
[
2
]);
init_method
=
std
::
stoi
(
argv
[
2
]);
nrepeat
=
std
::
stoi
(
argv
[
3
]);
nrepeat
=
std
::
stoi
(
argv
[
3
]);
do_log
=
std
::
stoi
(
argv
[
4
]);
}
}
else
if
(
argc
==
2
0
)
else
if
(
argc
==
2
1
)
{
{
do_verification
=
std
::
stoi
(
argv
[
1
]);
do_verification
=
std
::
stoi
(
argv
[
1
]);
init_method
=
std
::
stoi
(
argv
[
2
]);
init_method
=
std
::
stoi
(
argv
[
2
]);
nrepeat
=
std
::
stoi
(
argv
[
3
]);
nrepeat
=
std
::
stoi
(
argv
[
3
]);
do_log
=
std
::
stoi
(
argv
[
4
]);
N
=
std
::
stoi
(
argv
[
4
]);
K
=
std
::
stoi
(
argv
[
5
]);
N
=
std
::
stoi
(
argv
[
5
]);
C
=
std
::
stoi
(
argv
[
6
]);
K
=
std
::
stoi
(
argv
[
6
]);
Y
=
std
::
stoi
(
argv
[
7
]);
C
=
std
::
stoi
(
argv
[
7
]);
X
=
std
::
stoi
(
argv
[
8
]);
Y
=
std
::
stoi
(
argv
[
8
]);
Hi
=
std
::
stoi
(
argv
[
9
]);
X
=
std
::
stoi
(
argv
[
9
]);
Wi
=
std
::
stoi
(
argv
[
10
]);
Hi
=
std
::
stoi
(
argv
[
10
]);
conv_stride_h
=
std
::
stoi
(
argv
[
11
]);
Wi
=
std
::
stoi
(
argv
[
11
]);
conv_stride_w
=
std
::
stoi
(
argv
[
12
]);
conv_stride_h
=
std
::
stoi
(
argv
[
12
]);
conv_dilation_h
=
std
::
stoi
(
argv
[
13
]);
conv_stride_w
=
std
::
stoi
(
argv
[
13
]);
conv_dilation_w
=
std
::
stoi
(
argv
[
14
]);
conv_dilation_h
=
std
::
stoi
(
argv
[
14
]);
in_left_pad_h
=
std
::
stoi
(
argv
[
15
]);
conv_dilation_w
=
std
::
stoi
(
argv
[
15
]);
in_left_pad_w
=
std
::
stoi
(
argv
[
16
]);
in_left_pad_h
=
std
::
stoi
(
argv
[
16
]);
in_right_pad_h
=
std
::
stoi
(
argv
[
17
]);
in_left_pad_w
=
std
::
stoi
(
argv
[
17
]);
in_right_pad_w
=
std
::
stoi
(
argv
[
18
]);
in_right_pad_h
=
std
::
stoi
(
argv
[
18
]);
split_k
=
std
::
stoi
(
argv
[
19
]);
in_right_pad_w
=
std
::
stoi
(
argv
[
19
]);
split_k
=
std
::
stoi
(
argv
[
20
]);
}
}
else
else
{
{
printf
(
"arg1: verification (0=no, 1=yes)
\n
"
);
printf
(
"arg1: verification (0=no, 1=yes)
\n
"
);
printf
(
"arg2: initialization (0=no init, 1=integer value, 2=decimal value)
\n
"
);
printf
(
"arg2: initialization (0=no init, 1=integer value, 2=decimal value)
\n
"
);
printf
(
"arg3: run kernel # of times (>1)
\n
"
);
printf
(
"arg3: run kernel # of times (>1)
\n
"
);
printf
(
"arg4 to 18: N, K, C, Y, X, Hi, Wi, Sy, Sx, Dy, Dx, LeftPy, LeftPx, RightPy, "
printf
(
"arg4: is show log (0=no, 1=yes)
\n
"
);
"RightPx
\n
"
);
printf
(
"arg5 to 19: N, K, C, Y, X, Hi, Wi, Sy, Sx, Dy, Dx, LeftPy, LeftPx, RightPy, "
"RightPx, split-k
\n
"
);
exit
(
0
);
exit
(
0
);
}
}
...
@@ -267,7 +271,7 @@ int main(int argc, char* argv[])
...
@@ -267,7 +271,7 @@ int main(int argc, char* argv[])
wei_device_buf
.
FromDevice
(
wei_k_c_y_x_device_result
.
mData
.
data
());
wei_device_buf
.
FromDevice
(
wei_k_c_y_x_device_result
.
mData
.
data
());
if
(
1
)
if
(
do_log
)
{
{
LogRangeAsType
<
float
>
(
std
::
cout
<<
"out: "
,
out_n_k_ho_wo
.
mData
,
","
)
<<
std
::
endl
;
LogRangeAsType
<
float
>
(
std
::
cout
<<
"out: "
,
out_n_k_ho_wo
.
mData
,
","
)
<<
std
::
endl
;
LogRangeAsType
<
float
>
(
std
::
cout
<<
"in : "
,
in_n_c_hi_wi
.
mData
,
","
)
<<
std
::
endl
;
LogRangeAsType
<
float
>
(
std
::
cout
<<
"in : "
,
in_n_c_hi_wi
.
mData
,
","
)
<<
std
::
endl
;
...
...
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