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
48f35aa0
Commit
48f35aa0
authored
Aug 13, 2018
by
Paul
Browse files
Formatting
parent
fa52b516
Changes
2
Hide whitespace changes
Inline
Side-by-side
Showing
2 changed files
with
11 additions
and
12 deletions
+11
-12
src/targets/gpu/hip.cpp
src/targets/gpu/hip.cpp
+2
-5
src/targets/gpu/lowering.cpp
src/targets/gpu/lowering.cpp
+9
-7
No files found.
src/targets/gpu/hip.cpp
View file @
48f35aa0
...
@@ -11,10 +11,7 @@ namespace gpu {
...
@@ -11,10 +11,7 @@ namespace gpu {
using
hip_ptr
=
MIGRAPH_MANAGE_PTR
(
void
,
hipFree
);
using
hip_ptr
=
MIGRAPH_MANAGE_PTR
(
void
,
hipFree
);
std
::
string
hip_error
(
int
error
)
std
::
string
hip_error
(
int
error
)
{
return
hipGetErrorString
(
static_cast
<
hipError_t
>
(
error
));
}
{
return
hipGetErrorString
(
static_cast
<
hipError_t
>
(
error
));
}
hip_ptr
allocate_gpu
(
std
::
size_t
sz
)
hip_ptr
allocate_gpu
(
std
::
size_t
sz
)
{
{
...
@@ -54,7 +51,7 @@ hip_ptr write_to_gpu(const void* x, std::size_t sz)
...
@@ -54,7 +51,7 @@ hip_ptr write_to_gpu(const void* x, std::size_t sz)
argument
allocate_gpu
(
shape
s
)
argument
allocate_gpu
(
shape
s
)
{
{
auto
p
=
share
(
allocate_gpu
(
s
.
bytes
()
+
1
));
auto
p
=
share
(
allocate_gpu
(
s
.
bytes
()
+
1
));
return
{
s
,
[
p
]()
mutable
{
return
reinterpret_cast
<
char
*>
(
p
.
get
());
}};
return
{
s
,
[
p
]()
mutable
{
return
reinterpret_cast
<
char
*>
(
p
.
get
());
}};
}
}
...
...
src/targets/gpu/lowering.cpp
View file @
48f35aa0
...
@@ -102,12 +102,13 @@ struct miopen_convolution
...
@@ -102,12 +102,13 @@ struct miopen_convolution
auto
y_desc
=
make_tensor
(
output_shape
);
auto
y_desc
=
make_tensor
(
output_shape
);
std
::
size_t
workspace_size
=
0
;
std
::
size_t
workspace_size
=
0
;
miopenConvolutionForwardGetWorkSpaceSize
(
ctx
.
handle
.
get
(),
x_desc
.
get
(),
w_desc
.
get
(),
cd
.
get
(),
y_desc
.
get
(),
&
workspace_size
);
miopenConvolutionForwardGetWorkSpaceSize
(
ctx
.
handle
.
get
(),
x_desc
.
get
(),
w_desc
.
get
(),
cd
.
get
(),
y_desc
.
get
(),
&
workspace_size
);
workspace_shape
=
shape
{
shape
::
int8_type
,
{
workspace_size
}};
workspace_shape
=
shape
{
shape
::
int8_type
,
{
workspace_size
}};
auto
x
=
to_gpu
(
generate_argument
(
inputs
[
0
]
->
get_shape
()));
auto
x
=
to_gpu
(
generate_argument
(
inputs
[
0
]
->
get_shape
()));
auto
w
=
to_gpu
(
generate_argument
(
inputs
[
1
]
->
get_shape
()));
auto
w
=
to_gpu
(
generate_argument
(
inputs
[
1
]
->
get_shape
()));
auto
y
=
to_gpu
(
generate_argument
(
output_shape
));
auto
y
=
to_gpu
(
generate_argument
(
output_shape
));
auto
workspace
=
allocate_gpu
(
workspace_shape
);
auto
workspace
=
allocate_gpu
(
workspace_shape
);
int
algo_count
;
int
algo_count
;
...
@@ -355,12 +356,13 @@ struct miopen_apply
...
@@ -355,12 +356,13 @@ struct miopen_apply
auto
&&
op
=
any_cast
<
convolution
>
(
ins
->
op
);
auto
&&
op
=
any_cast
<
convolution
>
(
ins
->
op
);
auto
conv
=
miopen_convolution
{
op
,
make_conv
(
op
)};
auto
conv
=
miopen_convolution
{
op
,
make_conv
(
op
)};
auto
ws
=
conv
.
compile
(
ctx
,
ins
->
result
,
ins
->
arguments
);
auto
ws
=
conv
.
compile
(
ctx
,
ins
->
result
,
ins
->
arguments
);
auto
workspace
=
insert_allocation
(
ins
,
ws
);
auto
workspace
=
insert_allocation
(
ins
,
ws
);
auto
output
=
insert_allocation
(
ins
,
ins
->
result
);
auto
output
=
insert_allocation
(
ins
,
ins
->
result
);
prog
->
replace_instruction
(
ins
,
conv
,
ins
->
arguments
.
at
(
0
),
ins
->
arguments
.
at
(
1
),
workspace
,
output
);
prog
->
replace_instruction
(
ins
,
conv
,
ins
->
arguments
.
at
(
0
),
ins
->
arguments
.
at
(
1
),
workspace
,
output
);
}
}
void
apply_pooling
(
instruction_ref
ins
)
void
apply_pooling
(
instruction_ref
ins
)
...
...
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