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
0aeeb4bb
"git@developer.sourcefind.cn:gaoqiong/migraphx.git" did not exist on "d478675c6628534dddd1a222ee62bd9b5c0b5211"
Commit
0aeeb4bb
authored
Jul 13, 2018
by
wsttiger
Browse files
Fixed up for clang-tidy and cppcheck and formatting
parent
7b55afee
Changes
1
Show whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
16 additions
and
15 deletions
+16
-15
src/targets/miopen/kernels.cpp
src/targets/miopen/kernels.cpp
+16
-15
No files found.
src/targets/miopen/kernels.cpp
View file @
0aeeb4bb
...
@@ -6,7 +6,7 @@ namespace migraph {
...
@@ -6,7 +6,7 @@ namespace migraph {
namespace
miopen
{
namespace
miopen
{
template
<
int
NDIM
>
template
<
int
NDIM
>
struct
HIPT
ensor
D
escriptor
struct
hip_t
ensor
_d
escriptor
{
{
size_t
lens
[
NDIM
];
size_t
lens
[
NDIM
];
size_t
strides
[
NDIM
];
size_t
strides
[
NDIM
];
...
@@ -24,21 +24,21 @@ __host__ __device__ void multiindex(size_t (&strides)[NDIM], size_t idx, size_t*
...
@@ -24,21 +24,21 @@ __host__ __device__ void multiindex(size_t (&strides)[NDIM], size_t idx, size_t*
}
}
template
<
typename
T
,
int
NDIM
>
template
<
typename
T
,
int
NDIM
>
__global__
void
contiguous_gpu
(
const
T
*
A
,
__global__
void
contiguous_gpu
(
const
T
*
a
,
HIPT
ensor
D
escriptor
<
NDIM
>
td_a
,
hip_t
ensor
_d
escriptor
<
NDIM
>
a_desc
,
T
*
A
t
,
T
*
a
t
,
HIPT
ensor
D
escriptor
<
NDIM
>
td_at
,
hip_t
ensor
_d
escriptor
<
NDIM
>
at_desc
,
size_t
nelements
)
size_t
nelements
)
{
{
for
(
size_t
i
=
blockIdx
.
x
*
blockDim
.
x
+
threadIdx
.
x
;
i
<
nelements
;
for
(
size_t
i
=
blockIdx
.
x
*
blockDim
.
x
+
threadIdx
.
x
;
i
<
nelements
;
i
+=
blockDim
.
x
*
gridDim
.
x
)
i
+=
blockDim
.
x
*
gridDim
.
x
)
{
{
size_t
s
[
NDIM
];
size_t
s
[
NDIM
];
multiindex
<
NDIM
>
(
td_at
.
strides
,
i
,
s
);
multiindex
<
NDIM
>
(
at_desc
.
strides
,
i
,
s
);
size_t
lidx
=
0
;
size_t
lidx
=
0
;
for
(
size_t
j
=
0
;
j
<
NDIM
;
j
++
)
for
(
size_t
j
=
0
;
j
<
NDIM
;
j
++
)
lidx
+=
s
[
j
]
*
td_a
.
strides
[
j
];
lidx
+=
s
[
j
]
*
a_desc
.
strides
[
j
];
A
t
[
i
]
=
A
[
lidx
];
a
t
[
i
]
=
a
[
lidx
];
}
}
}
}
...
@@ -48,12 +48,13 @@ void hip_contiguous(migraph::shape output_shape, migraph::argument arg, migraph:
...
@@ -48,12 +48,13 @@ void hip_contiguous(migraph::shape output_shape, migraph::argument arg, migraph:
visit_all
(
result
,
arg
)([
&
](
auto
output
,
auto
input
)
{
visit_all
(
result
,
arg
)([
&
](
auto
output
,
auto
input
)
{
if
(
ndim
==
4
)
if
(
ndim
==
4
)
{
{
HIPTensorDescriptor
<
4
>
td_a
,
td_at
;
hip_tensor_descriptor
<
4
>
a_desc
{};
auto
s
=
arg
.
get_shape
();
hip_tensor_descriptor
<
4
>
at_desc
{};
const
auto
&
s
=
arg
.
get_shape
();
for
(
int
i
=
0
;
i
<
ndim
;
i
++
)
for
(
int
i
=
0
;
i
<
ndim
;
i
++
)
{
{
td_a
.
strides
[
i
]
=
s
.
strides
().
at
(
i
);
a_desc
.
strides
[
i
]
=
s
.
strides
().
at
(
i
);
td_at
.
strides
[
i
]
=
output_shape
.
strides
().
at
(
i
);
at_desc
.
strides
[
i
]
=
output_shape
.
strides
().
at
(
i
);
}
}
dim3
nblocks
(
512
);
dim3
nblocks
(
512
);
dim3
nthreads
(
512
);
dim3
nthreads
(
512
);
...
@@ -61,11 +62,11 @@ void hip_contiguous(migraph::shape output_shape, migraph::argument arg, migraph:
...
@@ -61,11 +62,11 @@ void hip_contiguous(migraph::shape output_shape, migraph::argument arg, migraph:
nblocks
,
nblocks
,
nthreads
,
nthreads
,
0
,
0
,
0
,
nullptr
,
input
.
data
(),
input
.
data
(),
td_a
,
a_desc
,
output
.
data
(),
output
.
data
(),
td_at
,
at_desc
,
s
.
elements
());
s
.
elements
());
}
}
else
else
...
...
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