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
a7ea008d
Commit
a7ea008d
authored
Jun 27, 2019
by
Paul
Browse files
Fix tidy warnings
parent
be7467a8
Changes
3
Hide whitespace changes
Inline
Side-by-side
Showing
3 changed files
with
15 additions
and
10 deletions
+15
-10
src/targets/gpu/device/include/migraphx/gpu/device/array.hpp
src/targets/gpu/device/include/migraphx/gpu/device/array.hpp
+1
-1
src/targets/gpu/device/include/migraphx/gpu/device/launch.hpp
...targets/gpu/device/include/migraphx/gpu/device/launch.hpp
+5
-5
src/targets/gpu/device/include/migraphx/gpu/device/reduce.hpp
...targets/gpu/device/include/migraphx/gpu/device/reduce.hpp
+9
-4
No files found.
src/targets/gpu/device/include/migraphx/gpu/device/array.hpp
View file @
a7ea008d
...
...
@@ -53,7 +53,7 @@ struct hip_array
friend
MIGRAPHX_DEVICE_CONSTEXPR
hip_array
operator
+
(
const
hip_array
&
x
,
const
hip_array
&
y
)
{
hip_array
result
;
hip_array
result
{}
;
for
(
std
::
size_t
i
=
0
;
i
<
N
;
i
++
)
result
[
i
]
=
x
[
i
]
+
y
[
i
];
return
result
;
...
...
src/targets/gpu/device/include/migraphx/gpu/device/launch.hpp
View file @
a7ea008d
...
...
@@ -11,13 +11,13 @@ namespace device {
struct
index
{
std
::
size_t
global
;
std
::
size_t
local
;
std
::
size_t
group
;
std
::
size_t
global
=
0
;
std
::
size_t
local
=
0
;
std
::
size_t
group
=
0
;
__device__
std
::
size_t
nglobal
()
const
{
return
blockDim
.
x
*
gridDim
.
x
;
}
__device__
std
::
size_t
nglobal
()
const
{
return
blockDim
.
x
*
gridDim
.
x
;
}
// NOLINT
__device__
std
::
size_t
nlocal
()
const
{
return
blockDim
.
x
;
}
__device__
std
::
size_t
nlocal
()
const
{
return
blockDim
.
x
;
}
// NOLINT
template
<
class
F
>
__device__
void
global_stride
(
std
::
size_t
n
,
F
f
)
const
...
...
src/targets/gpu/device/include/migraphx/gpu/device/reduce.hpp
View file @
a7ea008d
...
...
@@ -87,7 +87,7 @@ __device__ auto block_reduce(index idx, Op op, T init, std::size_t n, F f)
return
buffer
[
0
];
}
#else
constexpr
unsigned
int
dpp_row_shr
(
unsigned
int
x
)
{
return
0x110
|
x
;
}
constexpr
unsigned
int
dpp_row_shr
(
unsigned
int
x
)
{
return
0x110
u
|
x
;
}
constexpr
unsigned
int
dpp_row_bcast
(
unsigned
int
x
)
{
...
...
@@ -114,8 +114,9 @@ __device__ T dpp_mov(T& x)
uint32_t
reg
[
n
];
T
data
;
};
type
output
;
type
input
;
type
output
{};
type
input
{};
// cppcheck-suppress unreadVariable
input
.
data
=
x
;
for
(
std
::
size_t
i
=
0
;
i
<
n
;
i
++
)
{
...
...
@@ -142,8 +143,11 @@ __device__ void dpp_reduce(T& in, Op op)
in
=
op
(
in
,
out
);
}
__device__
void
dpp_reduce
(
float
&
x
,
sum
)
__device__
inline
void
dpp_reduce
(
float
&
x
,
sum
)
{
#ifdef MIGRAPHX_USE_CLANG_TIDY
(
void
)
x
;
#else
__asm__
volatile
(
"s_nop 4
\n
"
"v_add_f32 %0 %0 %0 row_shr:1
\n
"
"s_nop 1
\n
"
...
...
@@ -159,6 +163,7 @@ __device__ void dpp_reduce(float& x, sum)
"s_nop 1
\n
"
:
"=v"
(
x
)
:
"0"
(
x
));
#endif
}
template
<
std
::
size_t
N
,
class
Op
,
class
T
,
class
F
>
...
...
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