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
3cc34db2
Commit
3cc34db2
authored
Jul 20, 2018
by
wsttiger
Browse files
Fixed analyzer warnings
parent
76f68df4
Changes
1
Show whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
33 additions
and
21 deletions
+33
-21
src/targets/miopen/hip_contiguous.cpp
src/targets/miopen/hip_contiguous.cpp
+33
-21
No files found.
src/targets/miopen/hip_contiguous.cpp
View file @
3cc34db2
...
...
@@ -44,32 +44,47 @@ void visit_tensor_size(std::size_t n, F f)
}
}
template
<
size_t
NDim
>
struct
hip_index
{
size_t
d
[
NDim
];
size_t
&
operator
[](
size_t
i
)
{
return
d
[
i
];
}
size_t
operator
[](
size_t
i
)
const
{
return
d
[
i
];
}
};
template
<
size_t
NDim
>
struct
hip_tensor_descriptor
{
hip_tensor_descriptor
()
=
default
;
template
<
typename
T
,
typename
V
>
hip_tensor_descriptor
(
const
T
&
lens_
,
const
V
&
strides_
)
hip_tensor_descriptor
(
const
T
&
lens_
ext
,
const
V
&
strides_
ext
)
{
for
(
size_t
i
=
0
;
i
<
NDim
;
i
++
)
lens
[
i
]
=
lens_
[
i
];
lens
[
i
]
=
lens_
ext
[
i
];
for
(
size_t
i
=
0
;
i
<
NDim
;
i
++
)
strides
[
i
]
=
strides_
[
i
];
strides
[
i
]
=
strides_
ext
[
i
];
}
size_t
lens
[
NDim
];
size_t
strides
[
NDim
];
};
template
<
size_t
NDim
>
__host__
__device__
void
multiindex
(
size_t
(
&
strides
)[
NDim
],
size_t
idx
,
size_t
*
result
)
{
hip_index
<
NDim
>
multi
(
size_t
idx
)
{
hip_index
<
NDim
>
result
{};
size_t
tidx
=
idx
;
for
(
size_t
is
=
0
;
is
<
NDim
;
is
++
)
{
result
[
is
]
=
tidx
/
strides
[
is
];
tidx
=
tidx
%
strides
[
is
];
}
}
return
result
;
}
size_t
linear
(
hip_index
<
NDim
>
s
)
{
size_t
idx
=
0
;
for
(
size_t
i
=
0
;
i
<
NDim
;
i
++
)
idx
+=
s
[
i
]
*
strides
[
i
];
return
idx
;
}
size_t
lens
[
NDim
]
=
{};
size_t
strides
[
NDim
]
=
{};
};
template
<
typename
T
,
size_t
NDim
>
__global__
void
contiguous_gpu
(
const
T
*
a
,
...
...
@@ -81,11 +96,8 @@ __global__ void contiguous_gpu(const T* a,
for
(
size_t
i
=
blockIdx
.
x
*
blockDim
.
x
+
threadIdx
.
x
;
i
<
nelements
;
i
+=
blockDim
.
x
*
gridDim
.
x
)
{
size_t
s
[
NDim
];
multiindex
<
NDim
>
(
at_desc
.
strides
,
i
,
s
);
size_t
lidx
=
0
;
for
(
size_t
j
=
0
;
j
<
NDim
;
j
++
)
lidx
+=
s
[
j
]
*
a_desc
.
strides
[
j
];
hip_index
<
NDim
>
s
=
at_desc
.
multi
(
i
);
size_t
lidx
=
a_desc
.
linear
(
s
);
at
[
i
]
=
a
[
lidx
];
}
}
...
...
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