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
735e102a
"mmdet3d/vscode:/vscode.git/clone" did not exist on "49a1e5557672a4f2fc6c5c28e46751d3281e089d"
Commit
735e102a
authored
Aug 09, 2018
by
Paul
Browse files
Formatting
parent
e63c09c5
Changes
3
Hide whitespace changes
Inline
Side-by-side
Showing
3 changed files
with
13 additions
and
23 deletions
+13
-23
src/include/migraph/check_shapes.hpp
src/include/migraph/check_shapes.hpp
+1
-1
src/targets/gpu/hip_contiguous.cpp
src/targets/gpu/hip_contiguous.cpp
+9
-14
test/gpu/miopen.cpp
test/gpu/miopen.cpp
+3
-8
No files found.
src/include/migraph/check_shapes.hpp
View file @
735e102a
...
...
@@ -98,7 +98,7 @@ struct check_shapes
const
check_shapes
&
not_broadcasted
()
const
{
// if(!this->all_of([](const shape& s) { return not s.broadcasted(); }))
// MIGRAPH_THROW(prefix() + "Shapes are broadcasted");
// MIGRAPH_THROW(prefix() + "Shapes are broadcasted");
return
*
this
;
}
...
...
src/targets/gpu/hip_contiguous.cpp
View file @
735e102a
...
...
@@ -12,7 +12,7 @@ struct index
std
::
size_t
group
;
};
template
<
class
F
>
template
<
class
F
>
__global__
void
launcher
(
F
f
)
{
index
idx
{
blockIdx
.
x
*
blockDim
.
x
+
threadIdx
.
x
,
threadIdx
.
x
,
blockIdx
.
x
};
...
...
@@ -27,12 +27,7 @@ auto launch(std::size_t global, std::size_t local)
using
f_type
=
decltype
(
f
);
dim3
nblocks
(
global
/
local
);
dim3
nthreads
(
local
);
hipLaunchKernelGGL
((
launcher
<
f_type
>
),
nblocks
,
nthreads
,
0
,
nullptr
,
f
);
hipLaunchKernelGGL
((
launcher
<
f_type
>
),
nblocks
,
nthreads
,
0
,
nullptr
,
f
);
};
}
...
...
@@ -135,17 +130,17 @@ void hip_contiguous(migraph::shape output_shape, migraph::argument arg, migraph:
const
auto
&
s
=
arg
.
get_shape
();
hip_tensor_descriptor
<
ndim
>
a_desc
(
s
.
lens
(),
s
.
strides
());
hip_tensor_descriptor
<
ndim
>
at_desc
(
output_shape
.
lens
(),
output_shape
.
strides
());
auto
*
a
=
input
.
data
();
auto
*
at
=
output
.
data
();
auto
nelements
=
s
.
elements
();
std
::
size_t
nlocal
=
512
;
std
::
size_t
nglobal
=
512
*
nlocal
;
auto
*
a
=
input
.
data
();
auto
*
at
=
output
.
data
();
auto
nelements
=
s
.
elements
();
std
::
size_t
nlocal
=
512
;
std
::
size_t
nglobal
=
512
*
nlocal
;
launch
(
nglobal
,
nlocal
)([
=
](
auto
idx
)
mutable
{
for
(
size_t
i
=
idx
.
global
;
i
<
nelements
;
i
+=
nglobal
)
{
size_t
lidx
=
a_desc
.
linear
(
at_desc
.
multi
(
i
));
at
[
i
]
=
a
[
lidx
];
size_t
lidx
=
a_desc
.
linear
(
at_desc
.
multi
(
i
));
at
[
i
]
=
a
[
lidx
];
}
});
});
...
...
test/gpu/miopen.cpp
View file @
735e102a
...
...
@@ -17,17 +17,12 @@
struct
auto_eval
{
migraph
::
program
*
p
;
migraph
::
program
::
parameter_map
*
m
;
migraph
::
program
::
parameter_map
*
m
;
migraph
::
argument
result
;
auto_eval
(
migraph
::
program
&
pp
,
migraph
::
program
::
parameter_map
&
pm
)
:
p
(
&
pp
),
m
(
&
pm
)
{}
auto_eval
(
migraph
::
program
&
pp
,
migraph
::
program
::
parameter_map
&
pm
)
:
p
(
&
pp
),
m
(
&
pm
)
{}
migraph
::
argument
operator
()()
const
{
return
p
->
eval
(
*
m
);
}
migraph
::
argument
operator
()()
const
{
return
p
->
eval
(
*
m
);
}
~
auto_eval
()
{
...
...
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