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
9d12476e
Commit
9d12476e
authored
Sep 08, 2022
by
turneram
Browse files
Formatting
parent
cc2535e0
Changes
2
Hide whitespace changes
Inline
Side-by-side
Showing
2 changed files
with
8 additions
and
10 deletions
+8
-10
src/targets/gpu/jit/ck_gemm.cpp
src/targets/gpu/jit/ck_gemm.cpp
+0
-1
src/targets/gpu/kernels/include/migraphx/kernels/ck_gemm.hpp
src/targets/gpu/kernels/include/migraphx/kernels/ck_gemm.hpp
+8
-9
No files found.
src/targets/gpu/jit/ck_gemm.cpp
View file @
9d12476e
...
@@ -39,7 +39,6 @@ namespace migraphx {
...
@@ -39,7 +39,6 @@ namespace migraphx {
inline
namespace
MIGRAPHX_INLINE_NS
{
inline
namespace
MIGRAPHX_INLINE_NS
{
namespace
gpu
{
namespace
gpu
{
// NOLINTNEXTLINE
// NOLINTNEXTLINE
static
const
char
*
const
ck_gemm_kernel
=
R"__migraphx__(
static
const
char
*
const
ck_gemm_kernel
=
R"__migraphx__(
#include <migraphx/kernels/ck_gemm.hpp>
#include <migraphx/kernels/ck_gemm.hpp>
...
...
src/targets/gpu/kernels/include/migraphx/kernels/ck_gemm.hpp
View file @
9d12476e
...
@@ -37,7 +37,6 @@
...
@@ -37,7 +37,6 @@
#include "ck/device_utility/device_prop.hpp"
#include "ck/device_utility/device_prop.hpp"
#include "ck/device_utility/kernel_launch.hpp"
#include "ck/device_utility/kernel_launch.hpp"
namespace
migraphx
{
namespace
migraphx
{
// static constexpr auto I0 = Number<0>{};
// static constexpr auto I0 = Number<0>{};
...
@@ -165,17 +164,17 @@ namespace migraphx {
...
@@ -165,17 +164,17 @@ namespace migraphx {
template
<
class
T
,
class
U
,
class
V
>
template
<
class
T
,
class
U
,
class
V
>
__device__
void
ck_gemm
(
const
T
&
/* a_t */
,
const
U
&
/* b_t */
,
const
V
&
/* c_t */
)
__device__
void
ck_gemm
(
const
T
&
/* a_t */
,
const
U
&
/* b_t */
,
const
V
&
/* c_t */
)
{
{
constexpr
auto
alens
=
get_shape_c
<
T
>
{}.
lens
;
constexpr
auto
alens
=
get_shape_c
<
T
>
{}.
lens
;
constexpr
auto
m
=
alens
[
0
];
constexpr
auto
m
=
alens
[
0
];
constexpr
auto
k
=
alens
[
1
];
constexpr
auto
k
=
alens
[
1
];
constexpr
auto
alens
=
get_shape_c
<
U
>
{}.
lens
;
constexpr
auto
alens
=
get_shape_c
<
U
>
{}.
lens
;
constexpr
auto
n
=
alens
[
1
];
constexpr
auto
n
=
alens
[
1
];
constexpr
auto
astrides
=
get_shape_c
<
T
>
{}.
strides
;
constexpr
auto
astrides
=
get_shape_c
<
T
>
{}.
strides
;
constexpr
auto
as
=
astrides
[
1
];
constexpr
auto
as
=
astrides
[
1
];
constexpr
auto
bstrides
=
get_shape_c
<
U
>
{}.
strides
;
constexpr
auto
bstrides
=
get_shape_c
<
U
>
{}.
strides
;
constexpr
auto
bs
=
bstrides
[
1
];
constexpr
auto
bs
=
bstrides
[
1
];
constexpr
auto
cstrides
=
get_shape_c
<
V
>
{}.
strides
;
constexpr
auto
cstrides
=
get_shape_c
<
V
>
{}.
strides
;
constexpr
auto
cs
=
cstrides
[
1
];
constexpr
auto
cs
=
cstrides
[
1
];
printf
(
"%i %i %i, %i %i %i
\n
"
,
int
(
m
),
int
(
n
),
int
(
k
),
int
(
as
),
int
(
bs
),
int
(
cs
));
printf
(
"%i %i %i, %i %i %i
\n
"
,
int
(
m
),
int
(
n
),
int
(
k
),
int
(
as
),
int
(
bs
),
int
(
cs
));
}
}
...
...
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