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
bfa455a1
Commit
bfa455a1
authored
Jun 25, 2019
by
Shucai Xiao
Browse files
clang format
parent
605cce41
Changes
3
Hide whitespace changes
Inline
Side-by-side
Showing
3 changed files
with
10 additions
and
7 deletions
+10
-7
src/targets/gpu/device/logsoftmax.cpp
src/targets/gpu/device/logsoftmax.cpp
+2
-2
src/targets/gpu/device/softmax.cpp
src/targets/gpu/device/softmax.cpp
+2
-2
src/targets/gpu/include/migraphx/gpu/device/reduce_opers.hpp
src/targets/gpu/include/migraphx/gpu/device/reduce_opers.hpp
+6
-3
No files found.
src/targets/gpu/device/logsoftmax.cpp
View file @
bfa455a1
...
@@ -41,7 +41,7 @@ void logsoftmax(hipStream_t stream, argument result, argument arg, int axis)
...
@@ -41,7 +41,7 @@ void logsoftmax(hipStream_t stream, argument result, argument arg, int axis)
// load data to lds and compute the batch max
// load data to lds and compute the batch max
size_t
remaining_item_num
=
batch_item_num
;
size_t
remaining_item_num
=
batch_item_num
;
size_t
round_item_num
=
(
batch_item_num
+
block_size
-
1
)
/
block_size
*
block_size
;
size_t
round_item_num
=
(
batch_item_num
+
block_size
-
1
)
/
block_size
*
block_size
;
lds_data
[
max_block_size
]
=
input
[
0
];
lds_data
[
max_block_size
]
=
input
[
0
];
for
(
size_t
i
=
thr_idx
;
i
<
round_item_num
;
i
+=
block_size
)
for
(
size_t
i
=
thr_idx
;
i
<
round_item_num
;
i
+=
block_size
)
{
{
if
(
i
<
batch_item_num
)
if
(
i
<
batch_item_num
)
...
@@ -61,7 +61,7 @@ void logsoftmax(hipStream_t stream, argument result, argument arg, int axis)
...
@@ -61,7 +61,7 @@ void logsoftmax(hipStream_t stream, argument result, argument arg, int axis)
__syncthreads
();
__syncthreads
();
lds_data
[
max_block_size
]
=
0
;
lds_data
[
max_block_size
]
=
0
;
remaining_item_num
=
batch_item_num
;
remaining_item_num
=
batch_item_num
;
for
(
size_t
i
=
thr_idx
;
i
<
round_item_num
;
i
+=
block_size
)
for
(
size_t
i
=
thr_idx
;
i
<
round_item_num
;
i
+=
block_size
)
{
{
if
(
i
<
batch_item_num
)
if
(
i
<
batch_item_num
)
...
...
src/targets/gpu/device/softmax.cpp
View file @
bfa455a1
...
@@ -41,7 +41,7 @@ void softmax(hipStream_t stream, argument result, argument arg, int axis)
...
@@ -41,7 +41,7 @@ void softmax(hipStream_t stream, argument result, argument arg, int axis)
// load data to lds and compute the batch max
// load data to lds and compute the batch max
size_t
remaining_item_num
=
batch_item_num
;
size_t
remaining_item_num
=
batch_item_num
;
size_t
round_item_num
=
(
batch_item_num
+
block_size
-
1
)
/
block_size
*
block_size
;
size_t
round_item_num
=
(
batch_item_num
+
block_size
-
1
)
/
block_size
*
block_size
;
lds_data
[
max_block_size
]
=
input
[
0
];
lds_data
[
max_block_size
]
=
input
[
0
];
for
(
size_t
i
=
thr_idx
;
i
<
round_item_num
;
i
+=
block_size
)
for
(
size_t
i
=
thr_idx
;
i
<
round_item_num
;
i
+=
block_size
)
{
{
if
(
i
<
batch_item_num
)
if
(
i
<
batch_item_num
)
...
@@ -62,7 +62,7 @@ void softmax(hipStream_t stream, argument result, argument arg, int axis)
...
@@ -62,7 +62,7 @@ void softmax(hipStream_t stream, argument result, argument arg, int axis)
__syncthreads
();
__syncthreads
();
lds_data
[
max_block_size
]
=
0
;
lds_data
[
max_block_size
]
=
0
;
remaining_item_num
=
batch_item_num
;
remaining_item_num
=
batch_item_num
;
for
(
size_t
i
=
thr_idx
;
i
<
round_item_num
;
i
+=
block_size
)
for
(
size_t
i
=
thr_idx
;
i
<
round_item_num
;
i
+=
block_size
)
{
{
if
(
i
<
batch_item_num
)
if
(
i
<
batch_item_num
)
...
...
src/targets/gpu/include/migraphx/gpu/device/reduce_opers.hpp
View file @
bfa455a1
...
@@ -11,7 +11,8 @@ namespace gpu {
...
@@ -11,7 +11,8 @@ namespace gpu {
namespace
device
{
namespace
device
{
template
<
class
T
>
template
<
class
T
>
inline
__device__
void
reduce_max
(
T
*
data_ptr
,
size_t
block_size
,
size_t
thr_idx
,
size_t
item_num
,
size_t
max_index
)
inline
__device__
void
reduce_max
(
T
*
data_ptr
,
size_t
block_size
,
size_t
thr_idx
,
size_t
item_num
,
size_t
max_index
)
{
{
while
(
true
)
while
(
true
)
{
{
...
@@ -38,7 +39,8 @@ inline __device__ void reduce_max(T* data_ptr, size_t block_size, size_t thr_idx
...
@@ -38,7 +39,8 @@ inline __device__ void reduce_max(T* data_ptr, size_t block_size, size_t thr_idx
}
}
template
<
class
T
>
template
<
class
T
>
inline
__device__
void
reduce_min
(
T
*
data_ptr
,
size_t
block_size
,
size_t
thr_idx
,
size_t
item_num
,
size_t
min_index
)
inline
__device__
void
reduce_min
(
T
*
data_ptr
,
size_t
block_size
,
size_t
thr_idx
,
size_t
item_num
,
size_t
min_index
)
{
{
while
(
true
)
while
(
true
)
{
{
...
@@ -65,7 +67,8 @@ inline __device__ void reduce_min(T* data_ptr, size_t block_size, size_t thr_idx
...
@@ -65,7 +67,8 @@ inline __device__ void reduce_min(T* data_ptr, size_t block_size, size_t thr_idx
}
}
template
<
class
T
>
template
<
class
T
>
inline
__device__
void
reduce_sum
(
T
*
data_ptr
,
size_t
block_size
,
size_t
thr_idx
,
size_t
item_num
,
size_t
sum_index
)
inline
__device__
void
reduce_sum
(
T
*
data_ptr
,
size_t
block_size
,
size_t
thr_idx
,
size_t
item_num
,
size_t
sum_index
)
{
{
while
(
true
)
while
(
true
)
{
{
...
...
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