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
7da35f54
Commit
7da35f54
authored
Jun 26, 2019
by
Shucai Xiao
Browse files
clang format
parent
3e70d01b
Changes
2
Hide whitespace changes
Inline
Side-by-side
Showing
2 changed files
with
11 additions
and
10 deletions
+11
-10
src/targets/gpu/device/logsoftmax.cpp
src/targets/gpu/device/logsoftmax.cpp
+5
-3
src/targets/gpu/include/migraphx/gpu/device/reduce_opers.hpp
src/targets/gpu/include/migraphx/gpu/device/reduce_opers.hpp
+6
-7
No files found.
src/targets/gpu/device/logsoftmax.cpp
View file @
7da35f54
...
...
@@ -54,8 +54,9 @@ void logsoftmax(hipStream_t stream, const argument& result, const argument& arg,
auto
item_num
=
(
remaining_item_num
>
block_size
)
?
block_size
:
remaining_item_num
;
// reduce_max(lds_data, block_size, thr_idx, item_num, max_block_size);
block_reduce
<
type
,
max_op
<
type
>>
(
lds_data
,
max_op
<
type
>
{},
block_size
,
thr_idx
,
item_num
,
max_block_size
);
block_reduce
<
type
,
max_op
<
type
>>
(
lds_data
,
max_op
<
type
>
{},
block_size
,
thr_idx
,
item_num
,
max_block_size
);
remaining_item_num
-=
block_size
;
}
...
...
@@ -77,7 +78,8 @@ void logsoftmax(hipStream_t stream, const argument& result, const argument& arg,
auto
item_num
=
(
remaining_item_num
>
block_size
)
?
block_size
:
remaining_item_num
;
// reduce_sum(lds_data, block_size, thr_idx, item_num, max_block_size);
block_reduce
<
type
,
sum_op
<
type
>>
(
lds_data
,
sum_op
<
type
>
{},
block_size
,
thr_idx
,
item_num
,
max_block_size
);
block_reduce
<
type
,
sum_op
<
type
>>
(
lds_data
,
sum_op
<
type
>
{},
block_size
,
thr_idx
,
item_num
,
max_block_size
);
remaining_item_num
-=
block_size
;
}
...
...
src/targets/gpu/include/migraphx/gpu/device/reduce_opers.hpp
View file @
7da35f54
...
...
@@ -30,11 +30,11 @@ struct sum_op
template
<
class
T
,
class
Op
>
inline
__device__
void
block_reduce
(
T
*
data_ptr
,
Op
op
,
std
::
size_t
block_size
,
std
::
size_t
thr_idx
,
std
::
size_t
item_num
,
std
::
size_t
max_index
)
Op
op
,
std
::
size_t
block_size
,
std
::
size_t
thr_idx
,
std
::
size_t
item_num
,
std
::
size_t
max_index
)
{
while
(
true
)
{
...
...
@@ -42,7 +42,7 @@ inline __device__ void block_reduce(T* data_ptr,
auto
size
=
item_num
/
2
;
for
(
std
::
size_t
i
=
thr_idx
;
i
<
size
;
i
+=
block_size
)
{
//data_ptr[i] = ::max(to_hip_type(data_ptr[i]), to_hip_type(data_ptr[i + stride]));
//
data_ptr[i] = ::max(to_hip_type(data_ptr[i]), to_hip_type(data_ptr[i + stride]));
data_ptr
[
i
]
=
op
(
data_ptr
[
i
],
data_ptr
[
i
+
stride
]);
}
__syncthreads
();
...
...
@@ -62,7 +62,6 @@ inline __device__ void block_reduce(T* data_ptr,
__syncthreads
();
}
template
<
class
T
>
inline
__device__
void
reduce_max
(
T
*
data_ptr
,
std
::
size_t
block_size
,
...
...
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