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
MIGraphX
Commits
0e237605
Commit
0e237605
authored
Sep 13, 2022
by
turneram
Browse files
Formatting
parent
985fb0dd
Changes
2
Hide whitespace changes
Inline
Side-by-side
Showing
2 changed files
with
19 additions
and
18 deletions
+19
-18
src/targets/gpu/kernels/include/migraphx/kernels/ck_elementwise.hpp
...s/gpu/kernels/include/migraphx/kernels/ck_elementwise.hpp
+17
-17
test/verify/0ck_elementwise_half_test.cpp
test/verify/0ck_elementwise_half_test.cpp
+2
-1
No files found.
src/targets/gpu/kernels/include/migraphx/kernels/ck_elementwise.hpp
View file @
0e237605
...
...
@@ -50,9 +50,9 @@ struct CKBinaryElementwise
template
<
class
Desc_M
>
constexpr
auto
PadDescriptor_M_1d
(
Desc_M
desc_m
)
{
auto
gridSize
=
72
;
auto
blockSize
=
1024
;
auto
MPerThread
=
8
;
auto
gridSize
=
72
;
auto
blockSize
=
1024
;
auto
MPerThread
=
8
;
const
auto
M
=
desc_m
.
GetLength
(
I0
);
const
ck
::
index_t
loop_step
=
gridSize
*
blockSize
*
MPerThread
;
const
auto
pad
=
ck
::
math
::
integer_least_multiple
(
M
,
loop_step
)
-
M
;
...
...
@@ -67,8 +67,8 @@ struct CKBinaryElementwise
template
<
class
L
,
class
S
>
constexpr
auto
MakeDescriptor_M
(
const
L
&
lengths
,
const
S
&
strides
)
{
auto
tupleOfShape
=
generate_tuple
(
[
&
](
auto
I
)
{
return
static_cast
<
ck
::
index_t
>
(
lengths
[
I
]);
},
ck
::
Number
<
ndim
>
{});
auto
tupleOfShape
=
generate_tuple
(
[
&
](
auto
I
)
{
return
static_cast
<
ck
::
index_t
>
(
lengths
[
I
]);
},
ck
::
Number
<
ndim
>
{});
auto
tupleOfStride
=
generate_tuple
(
[
&
](
auto
I
)
{
return
static_cast
<
ck
::
index_t
>
(
strides
[
I
]);
},
ck
::
Number
<
ndim
>
{});
const
auto
desc
=
make_naive_tensor_descriptor
(
tupleOfShape
,
tupleOfStride
);
...
...
@@ -101,23 +101,23 @@ struct Add
template
<
class
T
,
class
U
,
class
V
>
__device__
void
ck_elementwise
(
const
T
&
a_t
,
const
U
&
b_t
,
const
V
&
c_t
)
{
constexpr
auto
a_lens
=
get_shape_c
<
T
>
{}.
lens
;
constexpr
auto
a_strides
=
get_shape_c
<
T
>
{}.
strides
;
constexpr
auto
a_lens
=
get_shape_c
<
T
>
{}.
lens
;
constexpr
auto
a_strides
=
get_shape_c
<
T
>
{}.
strides
;
constexpr
ck
::
index_t
a_ndim
=
decltype
(
a_lens
.
size
()){};
auto
a_bin_op
=
CKBinaryElementwise
<
a_ndim
>
{};
constexpr
auto
a_desc
=
a_bin_op
.
MakeDescriptor_M
(
a_lens
,
a_strides
);
auto
a_bin_op
=
CKBinaryElementwise
<
a_ndim
>
{};
constexpr
auto
a_desc
=
a_bin_op
.
MakeDescriptor_M
(
a_lens
,
a_strides
);
constexpr
auto
b_lens
=
get_shape_c
<
U
>
{}.
lens
;
constexpr
auto
b_strides
=
get_shape_c
<
U
>
{}.
strides
;
constexpr
auto
b_lens
=
get_shape_c
<
U
>
{}.
lens
;
constexpr
auto
b_strides
=
get_shape_c
<
U
>
{}.
strides
;
constexpr
ck
::
index_t
b_ndim
=
decltype
(
b_lens
.
size
()){};
auto
b_bin_op
=
CKBinaryElementwise
<
b_ndim
>
{};
constexpr
auto
b_desc
=
b_bin_op
.
MakeDescriptor_M
(
b_lens
,
b_strides
);
auto
b_bin_op
=
CKBinaryElementwise
<
b_ndim
>
{};
constexpr
auto
b_desc
=
b_bin_op
.
MakeDescriptor_M
(
b_lens
,
b_strides
);
constexpr
auto
c_lens
=
get_shape_c
<
V
>
{}.
lens
;
constexpr
auto
c_strides
=
get_shape_c
<
V
>
{}.
strides
;
constexpr
auto
c_lens
=
get_shape_c
<
V
>
{}.
lens
;
constexpr
auto
c_strides
=
get_shape_c
<
V
>
{}.
strides
;
constexpr
ck
::
index_t
c_ndim
=
decltype
(
c_lens
.
size
()){};
auto
c_bin_op
=
CKBinaryElementwise
<
c_ndim
>
{};
constexpr
auto
c_desc
=
c_bin_op
.
MakeDescriptor_M
(
c_lens
,
c_strides
);
auto
c_bin_op
=
CKBinaryElementwise
<
c_ndim
>
{};
constexpr
auto
c_desc
=
c_bin_op
.
MakeDescriptor_M
(
c_lens
,
c_strides
);
using
AGridDesc_M
=
decltype
(
a_desc
);
using
BGridDesc_M
=
decltype
(
b_desc
);
...
...
test/verify/0ck_elementwise_half_test.cpp
View file @
0e237605
...
...
@@ -37,7 +37,8 @@ struct ck_elementwise_half : verify_program<ck_elementwise_half>
migraphx
::
shape
m2_shape
{
migraphx
::
shape
::
half_type
,
{
3072
}};
auto
l1
=
mm
->
add_parameter
(
"1"
,
m1_shape
);
auto
l2
=
mm
->
add_parameter
(
"2"
,
m2_shape
);
l2
=
mm
->
add_instruction
(
migraphx
::
make_op
(
"multibroadcast"
,
{{
"out_lens"
,
{
2
,
384
,
3072
}}}),
l2
);
l2
=
mm
->
add_instruction
(
migraphx
::
make_op
(
"multibroadcast"
,
{{
"out_lens"
,
{
2
,
384
,
3072
}}}),
l2
);
mm
->
add_instruction
(
migraphx
::
make_op
(
"ck_elementwise"
),
l1
,
l2
);
...
...
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