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
0b217041
Commit
0b217041
authored
Nov 27, 2018
by
Paul
Browse files
Rename more things to migraphx
parent
53fd3f74
Changes
155
Hide whitespace changes
Inline
Side-by-side
Showing
20 changed files
with
94 additions
and
94 deletions
+94
-94
src/targets/gpu/device/contiguous.cpp
src/targets/gpu/device/contiguous.cpp
+2
-2
src/targets/gpu/device/include/migraphx/gpu/device/launch.hpp
...targets/gpu/device/include/migraphx/gpu/device/launch.hpp
+6
-6
src/targets/gpu/device/include/migraphx/gpu/device/nary.hpp
src/targets/gpu/device/include/migraphx/gpu/device/nary.hpp
+8
-8
src/targets/gpu/device/include/migraphx/gpu/device/tensor.hpp
...targets/gpu/device/include/migraphx/gpu/device/tensor.hpp
+4
-4
src/targets/gpu/device/include/migraphx/gpu/device/types.hpp
src/targets/gpu/device/include/migraphx/gpu/device/types.hpp
+4
-4
src/targets/gpu/device/mul.cpp
src/targets/gpu/device/mul.cpp
+2
-2
src/targets/gpu/device/sin.cpp
src/targets/gpu/device/sin.cpp
+2
-2
src/targets/gpu/eliminate_workspace.cpp
src/targets/gpu/eliminate_workspace.cpp
+3
-3
src/targets/gpu/fuse_ops.cpp
src/targets/gpu/fuse_ops.cpp
+10
-10
src/targets/gpu/gemm.cpp
src/targets/gpu/gemm.cpp
+3
-3
src/targets/gpu/hip.cpp
src/targets/gpu/hip.cpp
+11
-11
src/targets/gpu/include/migraphx/gpu/add.hpp
src/targets/gpu/include/migraphx/gpu/add.hpp
+4
-4
src/targets/gpu/include/migraphx/gpu/batchnorm.hpp
src/targets/gpu/include/migraphx/gpu/batchnorm.hpp
+4
-4
src/targets/gpu/include/migraphx/gpu/concat.hpp
src/targets/gpu/include/migraphx/gpu/concat.hpp
+4
-4
src/targets/gpu/include/migraphx/gpu/concat_gpu_opt.hpp
src/targets/gpu/include/migraphx/gpu/concat_gpu_opt.hpp
+2
-2
src/targets/gpu/include/migraphx/gpu/context.hpp
src/targets/gpu/include/migraphx/gpu/context.hpp
+9
-9
src/targets/gpu/include/migraphx/gpu/contiguous.hpp
src/targets/gpu/include/migraphx/gpu/contiguous.hpp
+4
-4
src/targets/gpu/include/migraphx/gpu/convolution.hpp
src/targets/gpu/include/migraphx/gpu/convolution.hpp
+4
-4
src/targets/gpu/include/migraphx/gpu/device/add.hpp
src/targets/gpu/include/migraphx/gpu/device/add.hpp
+4
-4
src/targets/gpu/include/migraphx/gpu/device/add_relu.hpp
src/targets/gpu/include/migraphx/gpu/device/add_relu.hpp
+4
-4
No files found.
src/targets/gpu/device/contiguous.cpp
View file @
0b217041
...
@@ -3,7 +3,7 @@
...
@@ -3,7 +3,7 @@
#include <migraphx/gpu/device/nary.hpp>
#include <migraphx/gpu/device/nary.hpp>
namespace
migraphx
{
namespace
migraphx
{
inline
namespace
MIGRAPH_INLINE_NS
{
inline
namespace
MIGRAPH
X
_INLINE_NS
{
namespace
gpu
{
namespace
gpu
{
namespace
device
{
namespace
device
{
...
@@ -14,5 +14,5 @@ void contiguous(hipStream_t stream, argument result, argument arg)
...
@@ -14,5 +14,5 @@ void contiguous(hipStream_t stream, argument result, argument arg)
}
// namespace device
}
// namespace device
}
// namespace gpu
}
// namespace gpu
}
// namespace MIGRAPH_INLINE_NS
}
// namespace MIGRAPH
X
_INLINE_NS
}
// namespace migraphx
}
// namespace migraphx
src/targets/gpu/device/include/migraphx/gpu/device/launch.hpp
View file @
0b217041
#ifndef MIGRAPH_GUARD_RTGLIB_DEVICE_LAUNCH_HPP
#ifndef MIGRAPH
X
_GUARD_RTGLIB_DEVICE_LAUNCH_HPP
#define MIGRAPH_GUARD_RTGLIB_DEVICE_LAUNCH_HPP
#define MIGRAPH
X
_GUARD_RTGLIB_DEVICE_LAUNCH_HPP
#include <hip/hip_runtime.h>
#include <hip/hip_runtime.h>
#include <migraphx/config.hpp>
#include <migraphx/config.hpp>
namespace
migraphx
{
namespace
migraphx
{
inline
namespace
MIGRAPH_INLINE_NS
{
inline
namespace
MIGRAPH
X
_INLINE_NS
{
namespace
gpu
{
namespace
gpu
{
namespace
device
{
namespace
device
{
...
@@ -53,14 +53,14 @@ inline auto gs_launch(hipStream_t stream, std::size_t n, std::size_t local = 102
...
@@ -53,14 +53,14 @@ inline auto gs_launch(hipStream_t stream, std::size_t n, std::size_t local = 102
// Workaround hcc's broken tile_static macro
// Workaround hcc's broken tile_static macro
#ifdef tile_static
#ifdef tile_static
#undef tile_static
#undef tile_static
#define MIGRAPH_DEVICE_SHARED __attribute__((tile_static))
#define MIGRAPH
X
_DEVICE_SHARED __attribute__((tile_static))
#else
#else
#define MIGRAPH_DEVICE_SHARED __shared__
#define MIGRAPH
X
_DEVICE_SHARED __shared__
#endif
#endif
}
// namespace device
}
// namespace device
}
// namespace gpu
}
// namespace gpu
}
// namespace MIGRAPH_INLINE_NS
}
// namespace MIGRAPH
X
_INLINE_NS
}
// namespace migraphx
}
// namespace migraphx
#endif
#endif
src/targets/gpu/device/include/migraphx/gpu/device/nary.hpp
View file @
0b217041
#ifndef MIGRAPH_GUARD_RTGLIB_DEVICE_NARY_HPP
#ifndef MIGRAPH
X
_GUARD_RTGLIB_DEVICE_NARY_HPP
#define MIGRAPH_GUARD_RTGLIB_DEVICE_NARY_HPP
#define MIGRAPH
X
_GUARD_RTGLIB_DEVICE_NARY_HPP
#include <migraphx/gpu/device/tensor.hpp>
#include <migraphx/gpu/device/tensor.hpp>
#include <migraphx/gpu/device/launch.hpp>
#include <migraphx/gpu/device/launch.hpp>
...
@@ -9,7 +9,7 @@
...
@@ -9,7 +9,7 @@
#include <migraphx/config.hpp>
#include <migraphx/config.hpp>
namespace
migraphx
{
namespace
migraphx
{
inline
namespace
MIGRAPH_INLINE_NS
{
inline
namespace
MIGRAPH
X
_INLINE_NS
{
namespace
gpu
{
namespace
gpu
{
namespace
device
{
namespace
device
{
...
@@ -87,7 +87,7 @@ void trinary_broadcast_vec_impl(hipStream_t stream,
...
@@ -87,7 +87,7 @@ void trinary_broadcast_vec_impl(hipStream_t stream,
const
std
::
size_t
bdim_vec_len
=
bdim_len
/
vec_size
;
const
std
::
size_t
bdim_vec_len
=
bdim_len
/
vec_size
;
launch
(
stream
,
nglobal
,
nlocal
)([
=
](
auto
idx
)
__device__
{
launch
(
stream
,
nglobal
,
nlocal
)([
=
](
auto
idx
)
__device__
{
MIGRAPH_DEVICE_SHARED
vec4
<
type
>
buffer
[
2048
/
vec_size
];
MIGRAPH
X
_DEVICE_SHARED
vec4
<
type
>
buffer
[
2048
/
vec_size
];
// Load bias into LDS
// Load bias into LDS
for
(
size_t
i
=
idx
.
local
;
i
<
bdim_vec_len
;
i
+=
nlocal
)
for
(
size_t
i
=
idx
.
local
;
i
<
bdim_vec_len
;
i
+=
nlocal
)
{
{
...
@@ -144,7 +144,7 @@ void trinary_broadcast_impl(hipStream_t stream,
...
@@ -144,7 +144,7 @@ void trinary_broadcast_impl(hipStream_t stream,
const
std
::
size_t
n
=
output
.
size
();
const
std
::
size_t
n
=
output
.
size
();
launch
(
stream
,
nglobal
,
nlocal
)([
=
](
auto
idx
)
__device__
{
launch
(
stream
,
nglobal
,
nlocal
)([
=
](
auto
idx
)
__device__
{
MIGRAPH_DEVICE_SHARED
type
buffer
[
2048
];
MIGRAPH
X
_DEVICE_SHARED
type
buffer
[
2048
];
// Load bias into LDS
// Load bias into LDS
for
(
size_t
i
=
idx
.
local
;
i
<
bdim_len
;
i
+=
nlocal
)
for
(
size_t
i
=
idx
.
local
;
i
<
bdim_len
;
i
+=
nlocal
)
{
{
...
@@ -192,7 +192,7 @@ void binary_broadcast_vec_impl(
...
@@ -192,7 +192,7 @@ void binary_broadcast_vec_impl(
const
std
::
size_t
bdim_vec_len
=
bdim_len
/
vec_size
;
const
std
::
size_t
bdim_vec_len
=
bdim_len
/
vec_size
;
launch
(
stream
,
nglobal
,
nlocal
)([
=
](
auto
idx
)
__device__
{
launch
(
stream
,
nglobal
,
nlocal
)([
=
](
auto
idx
)
__device__
{
MIGRAPH_DEVICE_SHARED
vec4
<
type
>
buffer
[
2048
/
vec_size
];
MIGRAPH
X
_DEVICE_SHARED
vec4
<
type
>
buffer
[
2048
/
vec_size
];
// Load bias into LDS
// Load bias into LDS
for
(
size_t
i
=
idx
.
local
;
i
<
bdim_vec_len
;
i
+=
nlocal
)
for
(
size_t
i
=
idx
.
local
;
i
<
bdim_vec_len
;
i
+=
nlocal
)
{
{
...
@@ -243,7 +243,7 @@ void binary_broadcast_impl(
...
@@ -243,7 +243,7 @@ void binary_broadcast_impl(
const
std
::
size_t
n
=
output
.
size
();
const
std
::
size_t
n
=
output
.
size
();
launch
(
stream
,
nglobal
,
nlocal
)([
=
](
auto
idx
)
__device__
{
launch
(
stream
,
nglobal
,
nlocal
)([
=
](
auto
idx
)
__device__
{
MIGRAPH_DEVICE_SHARED
type
buffer
[
2048
];
MIGRAPH
X
_DEVICE_SHARED
type
buffer
[
2048
];
// Load bias into LDS
// Load bias into LDS
for
(
size_t
i
=
idx
.
local
;
i
<
bdim_len
;
i
+=
nlocal
)
for
(
size_t
i
=
idx
.
local
;
i
<
bdim_len
;
i
+=
nlocal
)
{
{
...
@@ -396,7 +396,7 @@ inline auto nary(hipStream_t stream,
...
@@ -396,7 +396,7 @@ inline auto nary(hipStream_t stream,
}
// namespace device
}
// namespace device
}
// namespace gpu
}
// namespace gpu
}
// namespace MIGRAPH_INLINE_NS
}
// namespace MIGRAPH
X
_INLINE_NS
}
// namespace migraphx
}
// namespace migraphx
#endif
#endif
src/targets/gpu/device/include/migraphx/gpu/device/tensor.hpp
View file @
0b217041
#ifndef MIGRAPH_GUARD_RTGLIB_DEAVICE_TENSOR_HPP
#ifndef MIGRAPH
X
_GUARD_RTGLIB_DEAVICE_TENSOR_HPP
#define MIGRAPH_GUARD_RTGLIB_DEAVICE_TENSOR_HPP
#define MIGRAPH
X
_GUARD_RTGLIB_DEAVICE_TENSOR_HPP
#include <hip/hip_runtime.h>
#include <hip/hip_runtime.h>
#include <migraphx/functional.hpp>
#include <migraphx/functional.hpp>
#include <migraphx/config.hpp>
#include <migraphx/config.hpp>
namespace
migraphx
{
namespace
migraphx
{
inline
namespace
MIGRAPH_INLINE_NS
{
inline
namespace
MIGRAPH
X
_INLINE_NS
{
namespace
gpu
{
namespace
gpu
{
namespace
device
{
namespace
device
{
...
@@ -87,7 +87,7 @@ struct hip_tensor_descriptor
...
@@ -87,7 +87,7 @@ struct hip_tensor_descriptor
}
// namespace device
}
// namespace device
}
// namespace gpu
}
// namespace gpu
}
// namespace MIGRAPH_INLINE_NS
}
// namespace MIGRAPH
X
_INLINE_NS
}
// namespace migraphx
}
// namespace migraphx
#endif
#endif
src/targets/gpu/device/include/migraphx/gpu/device/types.hpp
View file @
0b217041
...
@@ -5,14 +5,14 @@
...
@@ -5,14 +5,14 @@
file LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt)
file LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt)
==============================================================================*/
==============================================================================*/
#ifndef MIGRAPH_GUARD_RTGLIB_GPU_DEVICE_TYPES_HPP
#ifndef MIGRAPH
X
_GUARD_RTGLIB_GPU_DEVICE_TYPES_HPP
#define MIGRAPH_GUARD_RTGLIB_GPU_DEVICE_TYPES_HPP
#define MIGRAPH
X
_GUARD_RTGLIB_GPU_DEVICE_TYPES_HPP
#include <migraphx/half.hpp>
#include <migraphx/half.hpp>
#include <migraphx/config.hpp>
#include <migraphx/config.hpp>
namespace
migraphx
{
namespace
migraphx
{
inline
namespace
MIGRAPH_INLINE_NS
{
inline
namespace
MIGRAPH
X
_INLINE_NS
{
namespace
gpu
{
namespace
gpu
{
namespace
device
{
namespace
device
{
...
@@ -86,7 +86,7 @@ inline float to_hip_type(gpu_half x) { return x; }
...
@@ -86,7 +86,7 @@ inline float to_hip_type(gpu_half x) { return x; }
}
// namespace device
}
// namespace device
}
// namespace gpu
}
// namespace gpu
}
// namespace MIGRAPH_INLINE_NS
}
// namespace MIGRAPH
X
_INLINE_NS
}
// namespace migraphx
}
// namespace migraphx
#endif
#endif
src/targets/gpu/device/mul.cpp
View file @
0b217041
...
@@ -2,7 +2,7 @@
...
@@ -2,7 +2,7 @@
#include <migraphx/gpu/device/nary.hpp>
#include <migraphx/gpu/device/nary.hpp>
namespace
migraphx
{
namespace
migraphx
{
inline
namespace
MIGRAPH_INLINE_NS
{
inline
namespace
MIGRAPH
X
_INLINE_NS
{
namespace
gpu
{
namespace
gpu
{
namespace
device
{
namespace
device
{
...
@@ -22,5 +22,5 @@ void mul(hipStream_t stream,
...
@@ -22,5 +22,5 @@ void mul(hipStream_t stream,
}
// namespace device
}
// namespace device
}
// namespace gpu
}
// namespace gpu
}
// namespace MIGRAPH_INLINE_NS
}
// namespace MIGRAPH
X
_INLINE_NS
}
// namespace migraphx
}
// namespace migraphx
src/targets/gpu/device/sin.cpp
View file @
0b217041
...
@@ -3,7 +3,7 @@
...
@@ -3,7 +3,7 @@
#include <migraphx/gpu/device/types.hpp>
#include <migraphx/gpu/device/types.hpp>
namespace
migraphx
{
namespace
migraphx
{
inline
namespace
MIGRAPH_INLINE_NS
{
inline
namespace
MIGRAPH
X
_INLINE_NS
{
namespace
gpu
{
namespace
gpu
{
namespace
device
{
namespace
device
{
...
@@ -14,5 +14,5 @@ void sin(hipStream_t stream, const argument& result, const argument& arg)
...
@@ -14,5 +14,5 @@ void sin(hipStream_t stream, const argument& result, const argument& arg)
}
// namespace device
}
// namespace device
}
// namespace gpu
}
// namespace gpu
}
// namespace MIGRAPH_INLINE_NS
}
// namespace MIGRAPH
X
_INLINE_NS
}
// namespace migraphx
}
// namespace migraphx
src/targets/gpu/eliminate_workspace.cpp
View file @
0b217041
...
@@ -9,12 +9,12 @@
...
@@ -9,12 +9,12 @@
#include <migraphx/pass_config.hpp>
#include <migraphx/pass_config.hpp>
namespace
migraphx
{
namespace
migraphx
{
inline
namespace
MIGRAPH_INLINE_NS
{
inline
namespace
MIGRAPH
X
_INLINE_NS
{
namespace
gpu
{
namespace
gpu
{
void
eliminate_workspace
::
apply
(
program
&
p
)
const
void
eliminate_workspace
::
apply
(
program
&
p
)
const
{
{
if
(
!
enabled
(
MIGRAPH_DISABLE_MEMORY_COLORING
{}))
if
(
!
enabled
(
MIGRAPH
X
_DISABLE_MEMORY_COLORING
{}))
return
;
return
;
std
::
size_t
n
=
0
;
std
::
size_t
n
=
0
;
...
@@ -41,5 +41,5 @@ void eliminate_workspace::apply(program& p) const
...
@@ -41,5 +41,5 @@ void eliminate_workspace::apply(program& p) const
}
}
}
// namespace gpu
}
// namespace gpu
}
// namespace MIGRAPH_INLINE_NS
}
// namespace MIGRAPH
X
_INLINE_NS
}
// namespace migraphx
}
// namespace migraphx
src/targets/gpu/fuse_ops.cpp
View file @
0b217041
...
@@ -6,7 +6,7 @@
...
@@ -6,7 +6,7 @@
#include <migraphx/instruction.hpp>
#include <migraphx/instruction.hpp>
namespace
migraphx
{
namespace
migraphx
{
inline
namespace
MIGRAPH_INLINE_NS
{
inline
namespace
MIGRAPH
X
_INLINE_NS
{
namespace
gpu
{
namespace
gpu
{
struct
fusion
struct
fusion
...
@@ -38,7 +38,7 @@ struct fusion
...
@@ -38,7 +38,7 @@ struct fusion
op_t
result
;
op_t
result
;
auto
status
=
miopenFusionPlanGetOp
(
fp
.
get
(),
i
,
&
result
);
auto
status
=
miopenFusionPlanGetOp
(
fp
.
get
(),
i
,
&
result
);
if
(
status
!=
miopenStatusSuccess
)
if
(
status
!=
miopenStatusSuccess
)
MIGRAPH_THROW
(
"Failed retrieving operator at "
+
std
::
to_string
(
i
));
MIGRAPH
X
_THROW
(
"Failed retrieving operator at "
+
std
::
to_string
(
i
));
return
result
;
return
result
;
}
}
...
@@ -51,7 +51,7 @@ struct fusion
...
@@ -51,7 +51,7 @@ struct fusion
auto
t
=
keep_alive
(
make_tensor
(
b
));
auto
t
=
keep_alive
(
make_tensor
(
b
));
auto
status
=
miopenCreateOpBiasForward
(
fp
.
get
(),
&
result
,
t
.
get
());
auto
status
=
miopenCreateOpBiasForward
(
fp
.
get
(),
&
result
,
t
.
get
());
if
(
status
!=
miopenStatusSuccess
)
if
(
status
!=
miopenStatusSuccess
)
MIGRAPH_THROW
(
"Creating operator failed"
);
MIGRAPH
X
_THROW
(
"Creating operator failed"
);
return
result
;
return
result
;
}
}
...
@@ -60,7 +60,7 @@ struct fusion
...
@@ -60,7 +60,7 @@ struct fusion
op_t
result
;
op_t
result
;
auto
status
=
miopenCreateOpActivationForward
(
fp
.
get
(),
&
result
,
miopenActivationRELU
);
auto
status
=
miopenCreateOpActivationForward
(
fp
.
get
(),
&
result
,
miopenActivationRELU
);
if
(
status
!=
miopenStatusSuccess
)
if
(
status
!=
miopenStatusSuccess
)
MIGRAPH_THROW
(
"Creating operator failed"
);
MIGRAPH
X
_THROW
(
"Creating operator failed"
);
return
result
;
return
result
;
}
}
...
@@ -71,7 +71,7 @@ struct fusion
...
@@ -71,7 +71,7 @@ struct fusion
auto
t
=
keep_alive
(
make_tensor
(
weights
));
auto
t
=
keep_alive
(
make_tensor
(
weights
));
auto
status
=
miopenCreateOpConvForward
(
fp
.
get
(),
&
result
,
cd
.
get
(),
t
.
get
());
auto
status
=
miopenCreateOpConvForward
(
fp
.
get
(),
&
result
,
cd
.
get
(),
t
.
get
());
if
(
status
!=
miopenStatusSuccess
)
if
(
status
!=
miopenStatusSuccess
)
MIGRAPH_THROW
(
"Creating operator failed"
);
MIGRAPH
X
_THROW
(
"Creating operator failed"
);
return
result
;
return
result
;
}
}
...
@@ -91,7 +91,7 @@ struct fusion
...
@@ -91,7 +91,7 @@ struct fusion
{
{
auto
status
=
miopenCompileFusionPlan
(
ctx
.
get_stream
().
get_miopen
(),
fp
.
get
());
auto
status
=
miopenCompileFusionPlan
(
ctx
.
get_stream
().
get_miopen
(),
fp
.
get
());
if
(
status
!=
miopenStatusSuccess
)
if
(
status
!=
miopenStatusSuccess
)
MIGRAPH_THROW
(
"Compiling fusion plan failed"
);
MIGRAPH
X
_THROW
(
"Compiling fusion plan failed"
);
}
}
argument
execute
(
context
&
ctx
,
argument
execute
(
context
&
ctx
,
...
@@ -109,12 +109,12 @@ struct fusion
...
@@ -109,12 +109,12 @@ struct fusion
y
.
implicit
(),
y
.
implicit
(),
fargs
.
get
());
fargs
.
get
());
if
(
status
!=
miopenStatusSuccess
)
if
(
status
!=
miopenStatusSuccess
)
MIGRAPH_THROW
(
"Failed to execute fusion plan"
);
MIGRAPH
X
_THROW
(
"Failed to execute fusion plan"
);
return
y
;
return
y
;
}
}
};
};
MIGRAPH_PRED_MATCHER
(
bias_shape
,
instruction_ref
ins
)
MIGRAPH
X
_PRED_MATCHER
(
bias_shape
,
instruction_ref
ins
)
{
{
auto
&&
s
=
ins
->
get_shape
();
auto
&&
s
=
ins
->
get_shape
();
return
s
.
broadcasted
()
and
s
.
strides
().
size
()
==
4
and
s
.
strides
()[
0
]
==
0
and
return
s
.
broadcasted
()
and
s
.
strides
().
size
()
==
4
and
s
.
strides
()[
0
]
==
0
and
...
@@ -128,7 +128,7 @@ std::array<T, sizeof...(Ts) + 1> make_array(T x, Ts... xs)
...
@@ -128,7 +128,7 @@ std::array<T, sizeof...(Ts) + 1> make_array(T x, Ts... xs)
return
{
std
::
move
(
x
),
std
::
move
(
static_cast
<
T
>
(
xs
))...};
return
{
std
::
move
(
x
),
std
::
move
(
static_cast
<
T
>
(
xs
))...};
}
}
MIGRAPH_PRED_MATCHER
(
fusable_conv
,
instruction_ref
ins
)
MIGRAPH
X
_PRED_MATCHER
(
fusable_conv
,
instruction_ref
ins
)
{
{
if
(
ins
->
name
()
!=
"gpu::convolution"
)
if
(
ins
->
name
()
!=
"gpu::convolution"
)
return
false
;
return
false
;
...
@@ -389,5 +389,5 @@ void fuse_ops::apply(program& p) const
...
@@ -389,5 +389,5 @@ void fuse_ops::apply(program& p) const
}
}
}
// namespace gpu
}
// namespace gpu
}
// namespace MIGRAPH_INLINE_NS
}
// namespace MIGRAPH
X
_INLINE_NS
}
// namespace migraphx
}
// namespace migraphx
src/targets/gpu/gemm.cpp
View file @
0b217041
...
@@ -5,7 +5,7 @@
...
@@ -5,7 +5,7 @@
#include <utility>
#include <utility>
namespace
migraphx
{
namespace
migraphx
{
inline
namespace
MIGRAPH_INLINE_NS
{
inline
namespace
MIGRAPH
X
_INLINE_NS
{
namespace
gpu
{
namespace
gpu
{
template
<
class
...
Ts
>
template
<
class
...
Ts
>
...
@@ -29,7 +29,7 @@ void generic_rocblas_gemm(shape::as<half>, Ts&&... xs)
...
@@ -29,7 +29,7 @@ void generic_rocblas_gemm(shape::as<half>, Ts&&... xs)
template
<
class
T
,
class
...
Ts
>
template
<
class
T
,
class
...
Ts
>
void
generic_rocblas_gemm
(
shape
::
as
<
T
>
,
Ts
&&
...)
void
generic_rocblas_gemm
(
shape
::
as
<
T
>
,
Ts
&&
...)
{
{
MIGRAPH_THROW
(
"Type unsupported by rocblas"
);
MIGRAPH
X
_THROW
(
"Type unsupported by rocblas"
);
}
}
template
<
class
T
>
template
<
class
T
>
...
@@ -111,5 +111,5 @@ argument miopen_gemm::compute(context& ctx,
...
@@ -111,5 +111,5 @@ argument miopen_gemm::compute(context& ctx,
}
}
}
// namespace gpu
}
// namespace gpu
}
// namespace MIGRAPH_INLINE_NS
}
// namespace MIGRAPH
X
_INLINE_NS
}
// namespace migraphx
}
// namespace migraphx
src/targets/gpu/hip.cpp
View file @
0b217041
...
@@ -7,10 +7,10 @@
...
@@ -7,10 +7,10 @@
#include <vector>
#include <vector>
namespace
migraphx
{
namespace
migraphx
{
inline
namespace
MIGRAPH_INLINE_NS
{
inline
namespace
MIGRAPH
X
_INLINE_NS
{
namespace
gpu
{
namespace
gpu
{
using
hip_ptr
=
MIGRAPH_MANAGE_PTR
(
void
,
hipFree
);
using
hip_ptr
=
MIGRAPH
X
_MANAGE_PTR
(
void
,
hipFree
);
std
::
string
hip_error
(
int
error
)
{
return
hipGetErrorString
(
static_cast
<
hipError_t
>
(
error
));
}
std
::
string
hip_error
(
int
error
)
{
return
hipGetErrorString
(
static_cast
<
hipError_t
>
(
error
));
}
...
@@ -19,20 +19,20 @@ std::size_t get_available_gpu_memory()
...
@@ -19,20 +19,20 @@ std::size_t get_available_gpu_memory()
size_t
free
,
total
;
size_t
free
,
total
;
auto
status
=
hipMemGetInfo
(
&
free
,
&
total
);
auto
status
=
hipMemGetInfo
(
&
free
,
&
total
);
if
(
status
!=
hipSuccess
)
if
(
status
!=
hipSuccess
)
MIGRAPH_THROW
(
"Failed getting available memory: "
+
hip_error
(
status
));
MIGRAPH
X
_THROW
(
"Failed getting available memory: "
+
hip_error
(
status
));
return
free
;
return
free
;
}
}
hip_ptr
allocate_gpu
(
std
::
size_t
sz
,
bool
host
=
false
)
hip_ptr
allocate_gpu
(
std
::
size_t
sz
,
bool
host
=
false
)
{
{
if
(
sz
>
get_available_gpu_memory
())
if
(
sz
>
get_available_gpu_memory
())
MIGRAPH_THROW
(
"Memory not available to allocate buffer: "
+
std
::
to_string
(
sz
));
MIGRAPH
X
_THROW
(
"Memory not available to allocate buffer: "
+
std
::
to_string
(
sz
));
void
*
result
;
void
*
result
;
auto
status
=
host
?
hipHostMalloc
(
&
result
,
sz
)
:
hipMalloc
(
&
result
,
sz
);
auto
status
=
host
?
hipHostMalloc
(
&
result
,
sz
)
:
hipMalloc
(
&
result
,
sz
);
if
(
status
!=
hipSuccess
)
if
(
status
!=
hipSuccess
)
{
{
if
(
host
)
if
(
host
)
MIGRAPH_THROW
(
"Gpu allocation failed: "
+
hip_error
(
status
));
MIGRAPH
X
_THROW
(
"Gpu allocation failed: "
+
hip_error
(
status
));
else
else
allocate_gpu
(
sz
,
true
);
allocate_gpu
(
sz
,
true
);
}
}
...
@@ -45,7 +45,7 @@ std::vector<T> read_from_gpu(const void* x, std::size_t sz)
...
@@ -45,7 +45,7 @@ std::vector<T> read_from_gpu(const void* x, std::size_t sz)
std
::
vector
<
T
>
result
(
sz
);
std
::
vector
<
T
>
result
(
sz
);
auto
status
=
hipMemcpy
(
result
.
data
(),
x
,
sz
*
sizeof
(
T
),
hipMemcpyDeviceToHost
);
auto
status
=
hipMemcpy
(
result
.
data
(),
x
,
sz
*
sizeof
(
T
),
hipMemcpyDeviceToHost
);
if
(
status
!=
hipSuccess
)
if
(
status
!=
hipSuccess
)
MIGRAPH_THROW
(
"Copy from gpu failed: "
+
hip_error
(
status
));
// NOLINT
MIGRAPH
X
_THROW
(
"Copy from gpu failed: "
+
hip_error
(
status
));
// NOLINT
return
result
;
return
result
;
}
}
...
@@ -54,7 +54,7 @@ hip_ptr write_to_gpu(const void* x, std::size_t sz, bool host = false)
...
@@ -54,7 +54,7 @@ hip_ptr write_to_gpu(const void* x, std::size_t sz, bool host = false)
auto
result
=
allocate_gpu
(
sz
,
host
);
auto
result
=
allocate_gpu
(
sz
,
host
);
auto
status
=
hipMemcpy
(
result
.
get
(),
x
,
sz
,
hipMemcpyHostToDevice
);
auto
status
=
hipMemcpy
(
result
.
get
(),
x
,
sz
,
hipMemcpyHostToDevice
);
if
(
status
!=
hipSuccess
)
if
(
status
!=
hipSuccess
)
MIGRAPH_THROW
(
"Copy to gpu failed: "
+
hip_error
(
status
));
MIGRAPH
X
_THROW
(
"Copy to gpu failed: "
+
hip_error
(
status
));
return
result
;
return
result
;
}
}
...
@@ -93,7 +93,7 @@ void set_device(std::size_t id)
...
@@ -93,7 +93,7 @@ void set_device(std::size_t id)
{
{
auto
status
=
hipSetDevice
(
id
);
auto
status
=
hipSetDevice
(
id
);
if
(
status
!=
hipSuccess
)
if
(
status
!=
hipSuccess
)
MIGRAPH_THROW
(
"Error setting device"
);
MIGRAPH
X
_THROW
(
"Error setting device"
);
}
}
void
gpu_sync
()
{
hipDeviceSynchronize
();
}
void
gpu_sync
()
{
hipDeviceSynchronize
();
}
...
@@ -103,12 +103,12 @@ void copy_to_gpu(argument src, argument dst)
...
@@ -103,12 +103,12 @@ void copy_to_gpu(argument src, argument dst)
std
::
size_t
src_size
=
src
.
get_shape
().
bytes
();
std
::
size_t
src_size
=
src
.
get_shape
().
bytes
();
std
::
size_t
dst_size
=
dst
.
get_shape
().
bytes
();
std
::
size_t
dst_size
=
dst
.
get_shape
().
bytes
();
if
(
src_size
>
dst_size
)
if
(
src_size
>
dst_size
)
MIGRAPH_THROW
(
"Not enough memory available in destination to do copy"
);
MIGRAPH
X
_THROW
(
"Not enough memory available in destination to do copy"
);
auto
status
=
hipMemcpy
(
dst
.
data
(),
src
.
data
(),
src_size
,
hipMemcpyHostToDevice
);
auto
status
=
hipMemcpy
(
dst
.
data
(),
src
.
data
(),
src_size
,
hipMemcpyHostToDevice
);
if
(
status
!=
hipSuccess
)
if
(
status
!=
hipSuccess
)
MIGRAPH_THROW
(
"Copy to gpu failed: "
+
hip_error
(
status
));
MIGRAPH
X
_THROW
(
"Copy to gpu failed: "
+
hip_error
(
status
));
}
}
}
// namespace gpu
}
// namespace gpu
}
// namespace MIGRAPH_INLINE_NS
}
// namespace MIGRAPH
X
_INLINE_NS
}
// namespace migraphx
}
// namespace migraphx
src/targets/gpu/include/migraphx/gpu/add.hpp
View file @
0b217041
#ifndef MIGRAPH_GUARD_RTGLIB_ADD_HPP
#ifndef MIGRAPH
X
_GUARD_RTGLIB_ADD_HPP
#define MIGRAPH_GUARD_RTGLIB_ADD_HPP
#define MIGRAPH
X
_GUARD_RTGLIB_ADD_HPP
#include <migraphx/gpu/lowering.hpp>
#include <migraphx/gpu/lowering.hpp>
#include <migraphx/manage_ptr.hpp>
#include <migraphx/manage_ptr.hpp>
...
@@ -19,7 +19,7 @@
...
@@ -19,7 +19,7 @@
#include <utility>
#include <utility>
namespace
migraphx
{
namespace
migraphx
{
inline
namespace
MIGRAPH_INLINE_NS
{
inline
namespace
MIGRAPH
X
_INLINE_NS
{
namespace
gpu
{
namespace
gpu
{
struct
hip_add
struct
hip_add
...
@@ -40,7 +40,7 @@ struct miopen_add
...
@@ -40,7 +40,7 @@ struct miopen_add
};
};
}
// namespace gpu
}
// namespace gpu
}
// namespace MIGRAPH_INLINE_NS
}
// namespace MIGRAPH
X
_INLINE_NS
}
// namespace migraphx
}
// namespace migraphx
#endif
#endif
src/targets/gpu/include/migraphx/gpu/batchnorm.hpp
View file @
0b217041
#ifndef MIGRAPH_GUARD_RTGLIB_BATCHNORM_HPP
#ifndef MIGRAPH
X
_GUARD_RTGLIB_BATCHNORM_HPP
#define MIGRAPH_GUARD_RTGLIB_BATCHNORM_HPP
#define MIGRAPH
X
_GUARD_RTGLIB_BATCHNORM_HPP
#include <migraphx/gpu/lowering.hpp>
#include <migraphx/gpu/lowering.hpp>
#include <migraphx/manage_ptr.hpp>
#include <migraphx/manage_ptr.hpp>
...
@@ -19,7 +19,7 @@
...
@@ -19,7 +19,7 @@
#include <utility>
#include <utility>
namespace
migraphx
{
namespace
migraphx
{
inline
namespace
MIGRAPH_INLINE_NS
{
inline
namespace
MIGRAPH
X
_INLINE_NS
{
namespace
gpu
{
namespace
gpu
{
struct
miopen_batch_norm_inference
struct
miopen_batch_norm_inference
...
@@ -33,7 +33,7 @@ struct miopen_batch_norm_inference
...
@@ -33,7 +33,7 @@ struct miopen_batch_norm_inference
};
};
}
// namespace gpu
}
// namespace gpu
}
// namespace MIGRAPH_INLINE_NS
}
// namespace MIGRAPH
X
_INLINE_NS
}
// namespace migraphx
}
// namespace migraphx
#endif
#endif
src/targets/gpu/include/migraphx/gpu/concat.hpp
View file @
0b217041
#ifndef MIGRAPH_GUARD_RTGLIB_CONCAT_HPP
#ifndef MIGRAPH
X
_GUARD_RTGLIB_CONCAT_HPP
#define MIGRAPH_GUARD_RTGLIB_CONCAT_HPP
#define MIGRAPH
X
_GUARD_RTGLIB_CONCAT_HPP
#include <migraphx/gpu/lowering.hpp>
#include <migraphx/gpu/lowering.hpp>
#include <migraphx/manage_ptr.hpp>
#include <migraphx/manage_ptr.hpp>
...
@@ -19,7 +19,7 @@
...
@@ -19,7 +19,7 @@
#include <utility>
#include <utility>
namespace
migraphx
{
namespace
migraphx
{
inline
namespace
MIGRAPH_INLINE_NS
{
inline
namespace
MIGRAPH
X
_INLINE_NS
{
namespace
gpu
{
namespace
gpu
{
struct
hip_concat
struct
hip_concat
...
@@ -34,7 +34,7 @@ struct hip_concat
...
@@ -34,7 +34,7 @@ struct hip_concat
};
};
}
// namespace gpu
}
// namespace gpu
}
// namespace MIGRAPH_INLINE_NS
}
// namespace MIGRAPH
X
_INLINE_NS
}
// namespace migraphx
}
// namespace migraphx
#endif
#endif
src/targets/gpu/include/migraphx/gpu/concat_gpu_opt.hpp
View file @
0b217041
#ifndef MIGRAPH_GUARD_RTGLIB_CONCAT_GPU_OPT_HPP
#ifndef MIGRAPH
X
_GUARD_RTGLIB_CONCAT_GPU_OPT_HPP
#define MIGRAPH_GUARD_RTGLIB_CONCAT_GPU_OPT_HPP
#define MIGRAPH
X
_GUARD_RTGLIB_CONCAT_GPU_OPT_HPP
#include <migraphx/gpu/concat.hpp>
#include <migraphx/gpu/concat.hpp>
...
...
src/targets/gpu/include/migraphx/gpu/context.hpp
View file @
0b217041
#ifndef MIGRAPH_GUARD_RTGLIB_CONTEXT_HPP
#ifndef MIGRAPH
X
_GUARD_RTGLIB_CONTEXT_HPP
#define MIGRAPH_GUARD_RTGLIB_CONTEXT_HPP
#define MIGRAPH
X
_GUARD_RTGLIB_CONTEXT_HPP
#include <migraphx/gpu/miopen.hpp>
#include <migraphx/gpu/miopen.hpp>
#include <migraphx/gpu/rocblas.hpp>
#include <migraphx/gpu/rocblas.hpp>
...
@@ -8,10 +8,10 @@
...
@@ -8,10 +8,10 @@
#include <migraphx/config.hpp>
#include <migraphx/config.hpp>
namespace
migraphx
{
namespace
migraphx
{
inline
namespace
MIGRAPH_INLINE_NS
{
inline
namespace
MIGRAPH
X
_INLINE_NS
{
namespace
gpu
{
namespace
gpu
{
MIGRAPH_DECLARE_ENV_VAR
(
MIGRAPH_DISABLE_NULL_STREAM
)
MIGRAPH
X
_DECLARE_ENV_VAR
(
MIGRAPH
X
_DISABLE_NULL_STREAM
)
struct
hip_device
struct
hip_device
{
{
...
@@ -21,7 +21,7 @@ struct hip_device
...
@@ -21,7 +21,7 @@ struct hip_device
struct
stream
struct
stream
{
{
using
hip_stream_ptr
=
MIGRAPH_MANAGE_PTR
(
hipStream_t
,
hipStreamDestroy
);
using
hip_stream_ptr
=
MIGRAPH
X
_MANAGE_PTR
(
hipStream_t
,
hipStreamDestroy
);
stream
()
{}
stream
()
{}
...
@@ -34,13 +34,13 @@ struct hip_device
...
@@ -34,13 +34,13 @@ struct hip_device
hipStream_t
result
=
nullptr
;
hipStream_t
result
=
nullptr
;
auto
status
=
hipStreamCreate
(
&
result
);
auto
status
=
hipStreamCreate
(
&
result
);
if
(
status
!=
hipSuccess
)
if
(
status
!=
hipSuccess
)
MIGRAPH_THROW
(
"Failed to allocate stream"
);
MIGRAPH
X
_THROW
(
"Failed to allocate stream"
);
return
hip_stream_ptr
{
result
};
return
hip_stream_ptr
{
result
};
}
}
hipStream_t
get
()
hipStream_t
get
()
{
{
if
(
enabled
(
MIGRAPH_DISABLE_NULL_STREAM
{}))
if
(
enabled
(
MIGRAPH
X
_DISABLE_NULL_STREAM
{}))
{
{
setup
();
setup
();
if
(
s
==
nullptr
)
if
(
s
==
nullptr
)
...
@@ -53,7 +53,7 @@ struct hip_device
...
@@ -53,7 +53,7 @@ struct hip_device
auto
create_miopen_handle
()
auto
create_miopen_handle
()
{
{
if
(
enabled
(
MIGRAPH_DISABLE_NULL_STREAM
{}))
if
(
enabled
(
MIGRAPH
X
_DISABLE_NULL_STREAM
{}))
return
make_obj
<
miopen_handle
>
(
&
miopenCreateWithStream
,
get
());
return
make_obj
<
miopen_handle
>
(
&
miopenCreateWithStream
,
get
());
else
else
return
make_obj
<
miopen_handle
>
(
&
miopenCreate
);
return
make_obj
<
miopen_handle
>
(
&
miopenCreate
);
...
@@ -116,7 +116,7 @@ struct context
...
@@ -116,7 +116,7 @@ struct context
std
::
shared_ptr
<
hip_device
>
current_device
;
std
::
shared_ptr
<
hip_device
>
current_device
;
};
};
}
// namespace gpu
}
// namespace gpu
}
// namespace MIGRAPH_INLINE_NS
}
// namespace MIGRAPH
X
_INLINE_NS
}
// namespace migraphx
}
// namespace migraphx
#endif
#endif
src/targets/gpu/include/migraphx/gpu/contiguous.hpp
View file @
0b217041
#ifndef MIGRAPH_GUARD_RTGLIB_CONTIGUOUS_HPP
#ifndef MIGRAPH
X
_GUARD_RTGLIB_CONTIGUOUS_HPP
#define MIGRAPH_GUARD_RTGLIB_CONTIGUOUS_HPP
#define MIGRAPH
X
_GUARD_RTGLIB_CONTIGUOUS_HPP
#include <migraphx/gpu/lowering.hpp>
#include <migraphx/gpu/lowering.hpp>
#include <migraphx/manage_ptr.hpp>
#include <migraphx/manage_ptr.hpp>
...
@@ -19,7 +19,7 @@
...
@@ -19,7 +19,7 @@
#include <utility>
#include <utility>
namespace
migraphx
{
namespace
migraphx
{
inline
namespace
MIGRAPH_INLINE_NS
{
inline
namespace
MIGRAPH
X
_INLINE_NS
{
namespace
gpu
{
namespace
gpu
{
struct
miopen_contiguous
struct
miopen_contiguous
...
@@ -32,7 +32,7 @@ struct miopen_contiguous
...
@@ -32,7 +32,7 @@ struct miopen_contiguous
};
};
}
// namespace gpu
}
// namespace gpu
}
// namespace MIGRAPH_INLINE_NS
}
// namespace MIGRAPH
X
_INLINE_NS
}
// namespace migraphx
}
// namespace migraphx
#endif
#endif
src/targets/gpu/include/migraphx/gpu/convolution.hpp
View file @
0b217041
#ifndef MIGRAPH_GUARD_RTGLIB_CONVOLUTION_HPP
#ifndef MIGRAPH
X
_GUARD_RTGLIB_CONVOLUTION_HPP
#define MIGRAPH_GUARD_RTGLIB_CONVOLUTION_HPP
#define MIGRAPH
X
_GUARD_RTGLIB_CONVOLUTION_HPP
#include <migraphx/gpu/lowering.hpp>
#include <migraphx/gpu/lowering.hpp>
#include <migraphx/manage_ptr.hpp>
#include <migraphx/manage_ptr.hpp>
...
@@ -19,7 +19,7 @@
...
@@ -19,7 +19,7 @@
#include <utility>
#include <utility>
namespace
migraphx
{
namespace
migraphx
{
inline
namespace
MIGRAPH_INLINE_NS
{
inline
namespace
MIGRAPH
X
_INLINE_NS
{
namespace
gpu
{
namespace
gpu
{
struct
miopen_convolution
struct
miopen_convolution
...
@@ -44,7 +44,7 @@ struct miopen_convolution
...
@@ -44,7 +44,7 @@ struct miopen_convolution
};
};
}
// namespace gpu
}
// namespace gpu
}
// namespace MIGRAPH_INLINE_NS
}
// namespace MIGRAPH
X
_INLINE_NS
}
// namespace migraphx
}
// namespace migraphx
#endif
#endif
src/targets/gpu/include/migraphx/gpu/device/add.hpp
View file @
0b217041
#ifndef MIGRAPH_GUARD_RTGLIB_DEVICE_ADD_HPP
#ifndef MIGRAPH
X
_GUARD_RTGLIB_DEVICE_ADD_HPP
#define MIGRAPH_GUARD_RTGLIB_DEVICE_ADD_HPP
#define MIGRAPH
X
_GUARD_RTGLIB_DEVICE_ADD_HPP
#include <migraphx/argument.hpp>
#include <migraphx/argument.hpp>
#include <migraphx/config.hpp>
#include <migraphx/config.hpp>
#include <hip/hip_runtime_api.h>
#include <hip/hip_runtime_api.h>
namespace
migraphx
{
namespace
migraphx
{
inline
namespace
MIGRAPH_INLINE_NS
{
inline
namespace
MIGRAPH
X
_INLINE_NS
{
namespace
gpu
{
namespace
gpu
{
namespace
device
{
namespace
device
{
...
@@ -21,7 +21,7 @@ void add(hipStream_t stream,
...
@@ -21,7 +21,7 @@ void add(hipStream_t stream,
}
// namespace device
}
// namespace device
}
// namespace gpu
}
// namespace gpu
}
// namespace MIGRAPH_INLINE_NS
}
// namespace MIGRAPH
X
_INLINE_NS
}
// namespace migraphx
}
// namespace migraphx
#endif
#endif
src/targets/gpu/include/migraphx/gpu/device/add_relu.hpp
View file @
0b217041
#ifndef MIGRAPH_GUARD_RTGLIB_DEVICE_ADD_RELU_HPP
#ifndef MIGRAPH
X
_GUARD_RTGLIB_DEVICE_ADD_RELU_HPP
#define MIGRAPH_GUARD_RTGLIB_DEVICE_ADD_RELU_HPP
#define MIGRAPH
X
_GUARD_RTGLIB_DEVICE_ADD_RELU_HPP
#include <migraphx/argument.hpp>
#include <migraphx/argument.hpp>
#include <migraphx/config.hpp>
#include <migraphx/config.hpp>
#include <hip/hip_runtime_api.h>
#include <hip/hip_runtime_api.h>
namespace
migraphx
{
namespace
migraphx
{
inline
namespace
MIGRAPH_INLINE_NS
{
inline
namespace
MIGRAPH
X
_INLINE_NS
{
namespace
gpu
{
namespace
gpu
{
namespace
device
{
namespace
device
{
...
@@ -24,7 +24,7 @@ void add_relu(hipStream_t stream,
...
@@ -24,7 +24,7 @@ void add_relu(hipStream_t stream,
}
// namespace device
}
// namespace device
}
// namespace gpu
}
// namespace gpu
}
// namespace MIGRAPH_INLINE_NS
}
// namespace MIGRAPH
X
_INLINE_NS
}
// namespace migraphx
}
// namespace migraphx
#endif
#endif
Prev
1
2
3
4
5
6
7
8
Next
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