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
d13dcab5
Commit
d13dcab5
authored
Jun 28, 2019
by
Shucai Xiao
Browse files
clang format
parent
ce649bb5
Changes
3
Hide whitespace changes
Inline
Side-by-side
Showing
3 changed files
with
11 additions
and
18 deletions
+11
-18
src/targets/gpu/device/argmax.cpp
src/targets/gpu/device/argmax.cpp
+1
-1
src/targets/gpu/device/argmin.cpp
src/targets/gpu/device/argmin.cpp
+1
-1
src/targets/gpu/include/migraphx/gpu/device/arg_op.hpp
src/targets/gpu/include/migraphx/gpu/device/arg_op.hpp
+9
-16
No files found.
src/targets/gpu/device/argmax.cpp
View file @
d13dcab5
...
...
@@ -15,7 +15,7 @@ namespace device {
void
argmax
(
hipStream_t
stream
,
const
argument
&
result
,
const
argument
&
arg
,
int
axis
)
{
arg
.
visit
([
&
](
auto
input
)
{
using
type
=
device_type
<
std
::
remove_cv_t
<
typename
decltype
(
input
)
::
value_type
>>
;
using
type
=
device_type
<
std
::
remove_cv_t
<
typename
decltype
(
input
)
::
value_type
>>
;
arg_op
<
pair_max
<
type
,
int64_t
>>
(
pair_max
<
type
,
int64_t
>
{},
stream
,
result
,
arg
,
axis
);
});
}
...
...
src/targets/gpu/device/argmin.cpp
View file @
d13dcab5
...
...
@@ -15,7 +15,7 @@ namespace device {
void
argmin
(
hipStream_t
stream
,
const
argument
&
result
,
const
argument
&
arg
,
int
axis
)
{
arg
.
visit
([
&
](
auto
input
)
{
using
type
=
device_type
<
std
::
remove_cv_t
<
typename
decltype
(
input
)
::
value_type
>>
;
using
type
=
device_type
<
std
::
remove_cv_t
<
typename
decltype
(
input
)
::
value_type
>>
;
arg_op
<
pair_min
<
type
,
int64_t
>>
(
pair_min
<
type
,
int64_t
>
{},
stream
,
result
,
arg
,
axis
);
});
}
...
...
src/targets/gpu/include/migraphx/gpu/device/arg_op.hpp
View file @
d13dcab5
...
...
@@ -41,12 +41,12 @@ struct pair_min
template
<
class
T
,
class
Op
>
inline
__device__
void
block_reduce_arg
(
T
*
data_ptr
,
int64_t
*
index_ptr
,
Op
op
,
std
::
size_t
block_size
,
std
::
size_t
thr_idx
,
std
::
size_t
item_num
,
std
::
size_t
output_index
)
int64_t
*
index_ptr
,
Op
op
,
std
::
size_t
block_size
,
std
::
size_t
thr_idx
,
std
::
size_t
item_num
,
std
::
size_t
output_index
)
{
while
(
true
)
{
...
...
@@ -77,8 +77,7 @@ inline __device__ void block_reduce_arg(T* data_ptr,
__syncthreads
();
}
template
<
class
Op
>
template
<
class
Op
>
void
arg_op
(
Op
op
,
hipStream_t
stream
,
const
argument
&
result
,
const
argument
&
arg
,
int
axis
)
{
auto
arg_shape
=
arg
.
get_shape
();
...
...
@@ -124,13 +123,8 @@ void arg_op(Op op, hipStream_t stream, const argument& result, const argument& a
__syncthreads
();
auto
item_num
=
(
remaining_item_num
>
block_size
)
?
block_size
:
remaining_item_num
;
block_reduce_arg
<
type
,
Op
>
(
lds_data
,
lds_index
,
op
,
block_size
,
thr_idx
,
item_num
,
max_block_size
);
block_reduce_arg
<
type
,
Op
>
(
lds_data
,
lds_index
,
op
,
block_size
,
thr_idx
,
item_num
,
max_block_size
);
remaining_item_num
-=
block_size
;
}
...
...
@@ -149,4 +143,3 @@ void arg_op(Op op, hipStream_t stream, const argument& result, const argument& a
}
// namespace migraphx
#endif
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