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
08818705
Commit
08818705
authored
Mar 08, 2022
by
Shucai Xiao
Browse files
clang format
parent
1d9d1e49
Changes
1
Hide whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
13 additions
and
11 deletions
+13
-11
src/targets/gpu/device/contiguous.cpp
src/targets/gpu/device/contiguous.cpp
+13
-11
No files found.
src/targets/gpu/device/contiguous.cpp
View file @
08818705
...
...
@@ -9,18 +9,19 @@ inline namespace MIGRAPHX_INLINE_NS {
namespace
gpu
{
namespace
device
{
__global__
void
cont_kernel
(
void
*
in
,
void
*
out
,
int
os1
,
int
os2
,
int
os3
,
int
is1
,
int
is2
,
int
is3
)
__global__
void
cont_kernel
(
void
*
in
,
void
*
out
,
int
os1
,
int
os2
,
int
os3
,
int
is1
,
int
is2
,
int
is3
)
{
int
i1
=
blockIdx
.
x
;
int
i2
=
blockIdx
.
y
;
int
i3
=
blockIdx
.
z
;
int
i4
=
threadIdx
.
x
;
__half
*
in_ptr
=
reinterpret_cast
<
__half
*>
(
in
);
__half
*
in_ptr
=
reinterpret_cast
<
__half
*>
(
in
);
__half
*
out_ptr
=
reinterpret_cast
<
__half
*>
(
out
);
int
out_idx
=
i1
*
os1
+
i2
*
os2
+
i3
*
os3
+
i4
;
int
in_idx
=
i1
*
is1
+
i2
*
is2
+
i3
*
is3
+
i4
;
int
out_idx
=
i1
*
os1
+
i2
*
os2
+
i3
*
os3
+
i4
;
int
in_idx
=
i1
*
is1
+
i2
*
is2
+
i3
*
is3
+
i4
;
out_ptr
[
out_idx
]
=
in_ptr
[
in_idx
];
}
...
...
@@ -39,16 +40,17 @@ void contiguous_nonstandard(hipStream_t stream, const argument& result, const ar
// auto in_stride = in_s.strides();
// auto out_stride = s.strides();
// cont_kernel<<<grid, block, 0, stream>>>(arg.data(), result.data(), out_stride[0], out_stride[1], out_stride[2], in_stride[0], in_stride[1], in_stride[2]);
// cont_kernel<<<grid, block, 0, stream>>>(arg.data(), result.data(), out_stride[0],
// out_stride[1], out_stride[2], in_stride[0], in_stride[1], in_stride[2]);
// }
// else
// {
visit_all
(
result
,
arg
)([
&
](
auto
output_v
,
auto
input_v
)
{
hip_visit_views
(
output_v
,
input_v
,
s
)([
&
](
auto
output
,
auto
input
,
auto
standard_shape
)
{
mi_gs_launch
(
stream
,
standard_shape
)([
=
](
auto
idx
)
__device__
{
output
[
idx
]
=
input
[
idx
];
});
});
});
visit_all
(
result
,
arg
)([
&
](
auto
output_v
,
auto
input_v
)
{
hip_visit_views
(
output_v
,
input_v
,
s
)([
&
](
auto
output
,
auto
input
,
auto
standard_shape
)
{
mi_gs_launch
(
stream
,
standard_shape
)([
=
](
auto
idx
)
__device__
{
output
[
idx
]
=
input
[
idx
];
});
});
});
// }
}
...
...
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