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
263579c2
Commit
263579c2
authored
Jun 03, 2018
by
Paul
Browse files
Formatting
parent
d007b98f
Changes
8
Show whitespace changes
Inline
Side-by-side
Showing
8 changed files
with
67 additions
and
62 deletions
+67
-62
src/include/rtg/argument.hpp
src/include/rtg/argument.hpp
+4
-2
src/include/rtg/manage_ptr.hpp
src/include/rtg/manage_ptr.hpp
+1
-1
src/include/rtg/operators.hpp
src/include/rtg/operators.hpp
+1
-4
src/include/rtg/program.hpp
src/include/rtg/program.hpp
+1
-1
src/program.cpp
src/program.cpp
+14
-12
src/targets/cpu/cpu_target.cpp
src/targets/cpu/cpu_target.cpp
+4
-1
src/targets/miopen/miopen_target.cpp
src/targets/miopen/miopen_target.cpp
+27
-16
test/miopen.cpp
test/miopen.cpp
+15
-25
No files found.
src/include/rtg/argument.hpp
View file @
263579c2
...
@@ -26,8 +26,10 @@ struct argument : raw_data<argument>
...
@@ -26,8 +26,10 @@ struct argument : raw_data<argument>
}
}
argument
(
shape
s
,
std
::
function
<
char
*
()
>
d
)
:
data
(
d
),
m_shape
(
s
)
{}
argument
(
shape
s
,
std
::
function
<
char
*
()
>
d
)
:
data
(
d
),
m_shape
(
s
)
{}
template
<
class
T
>
template
<
class
T
>
argument
(
shape
s
,
T
*
d
)
:
data
([
d
]
{
return
reinterpret_cast
<
char
*>
(
d
);
}),
m_shape
(
s
)
{}
argument
(
shape
s
,
T
*
d
)
:
data
([
d
]
{
return
reinterpret_cast
<
char
*>
(
d
);
}),
m_shape
(
s
)
{
}
/// Provides a raw pointer to the data
/// Provides a raw pointer to the data
std
::
function
<
char
*
()
>
data
;
std
::
function
<
char
*
()
>
data
;
...
...
src/include/rtg/manage_ptr.hpp
View file @
263579c2
...
@@ -43,7 +43,7 @@ using remove_ptr = typename std::
...
@@ -43,7 +43,7 @@ using remove_ptr = typename std::
template
<
class
T
>
template
<
class
T
>
using
shared
=
std
::
shared_ptr
<
remove_ptr
<
T
>>
;
using
shared
=
std
::
shared_ptr
<
remove_ptr
<
T
>>
;
template
<
class
T
>
template
<
class
T
>
shared
<
T
>
share
(
T
p
)
shared
<
T
>
share
(
T
p
)
{
{
return
shared
<
T
>
{
std
::
move
(
p
)};
return
shared
<
T
>
{
std
::
move
(
p
)};
...
...
src/include/rtg/operators.hpp
View file @
263579c2
...
@@ -227,10 +227,7 @@ struct outline
...
@@ -227,10 +227,7 @@ struct outline
check_shapes
{
inputs
}.
has
(
0
);
check_shapes
{
inputs
}.
has
(
0
);
return
s
;
return
s
;
}
}
argument
compute
(
shape
,
std
::
vector
<
argument
>
)
const
argument
compute
(
shape
,
std
::
vector
<
argument
>
)
const
{
return
{
s
,
nullptr
};
}
{
return
{
s
,
nullptr
};
}
};
};
}
// namespace rtg
}
// namespace rtg
...
...
src/include/rtg/program.hpp
View file @
263579c2
src/program.cpp
View file @
263579c2
...
@@ -80,8 +80,10 @@ shape program::get_parameter_shape(std::string name)
...
@@ -80,8 +80,10 @@ shape program::get_parameter_shape(std::string name)
return
false
;
return
false
;
}
}
});
});
if
(
ins
!=
this
->
end
())
return
ins
->
result
;
if
(
ins
!=
this
->
end
())
else
return
{};
return
ins
->
result
;
else
return
{};
}
}
bool
program
::
has_instruction
(
instruction_ref
ins
)
const
bool
program
::
has_instruction
(
instruction_ref
ins
)
const
...
...
src/targets/cpu/cpu_target.cpp
View file @
263579c2
...
@@ -26,7 +26,10 @@ struct cpu_convolution
...
@@ -26,7 +26,10 @@ struct cpu_convolution
auto
wei_h
=
weights
.
get_shape
().
lens
()[
2
];
auto
wei_h
=
weights
.
get_shape
().
lens
()[
2
];
auto
wei_w
=
weights
.
get_shape
().
lens
()[
3
];
auto
wei_w
=
weights
.
get_shape
().
lens
()[
3
];
dfor
(
output_shape
.
lens
()[
0
],
output_shape
.
lens
()[
1
],
output_shape
.
lens
()[
2
],
output_shape
.
lens
()[
3
])(
dfor
(
output_shape
.
lens
()[
0
],
output_shape
.
lens
()[
1
],
output_shape
.
lens
()[
2
],
output_shape
.
lens
()[
3
])(
[
&
](
std
::
size_t
o
,
std
::
size_t
w
,
std
::
size_t
i
,
std
::
size_t
j
)
{
[
&
](
std
::
size_t
o
,
std
::
size_t
w
,
std
::
size_t
i
,
std
::
size_t
j
)
{
const
int
start_x
=
i
*
op
.
stride
[
0
]
-
op
.
padding
[
0
];
const
int
start_x
=
i
*
op
.
stride
[
0
]
-
op
.
padding
[
0
];
const
int
start_y
=
j
*
op
.
stride
[
1
]
-
op
.
padding
[
1
];
const
int
start_y
=
j
*
op
.
stride
[
1
]
-
op
.
padding
[
1
];
...
...
src/targets/miopen/miopen_target.cpp
View file @
263579c2
...
@@ -8,7 +8,6 @@
...
@@ -8,7 +8,6 @@
namespace
rtg
{
namespace
rtg
{
namespace
miopen
{
namespace
miopen
{
struct
hip_allocate
struct
hip_allocate
{
{
std
::
string
name
()
const
{
return
"hip::allocate"
;
}
std
::
string
name
()
const
{
return
"hip::allocate"
;
}
...
@@ -19,7 +18,7 @@ struct hip_allocate
...
@@ -19,7 +18,7 @@ struct hip_allocate
}
}
argument
compute
(
shape
output_shape
,
std
::
vector
<
argument
>
)
const
argument
compute
(
shape
output_shape
,
std
::
vector
<
argument
>
)
const
{
{
char
*
data
=
nullptr
;
char
*
data
=
nullptr
;
// TODO: Check return status
// TODO: Check return status
hipMalloc
(
&
data
,
output_shape
.
bytes
());
hipMalloc
(
&
data
,
output_shape
.
bytes
());
return
{
output_shape
,
data
};
return
{
output_shape
,
data
};
...
@@ -42,7 +41,6 @@ struct hip_free
...
@@ -42,7 +41,6 @@ struct hip_free
}
}
};
};
using
miopen_handle
=
RTG_MANAGE_PTR
(
miopenHandle_t
,
miopenDestroy
);
using
miopen_handle
=
RTG_MANAGE_PTR
(
miopenHandle_t
,
miopenDestroy
);
using
tensor_descriptor
=
RTG_MANAGE_PTR
(
miopenTensorDescriptor_t
,
miopenDestroyTensorDescriptor
);
using
tensor_descriptor
=
RTG_MANAGE_PTR
(
miopenTensorDescriptor_t
,
miopenDestroyTensorDescriptor
);
using
convolution_descriptor
=
RTG_MANAGE_PTR
(
miopenConvolutionDescriptor_t
,
using
convolution_descriptor
=
RTG_MANAGE_PTR
(
miopenConvolutionDescriptor_t
,
...
@@ -163,7 +161,14 @@ struct miopen_relu
...
@@ -163,7 +161,14 @@ struct miopen_relu
float
alpha
=
1
,
beta
=
0
;
float
alpha
=
1
,
beta
=
0
;
auto
x_desc
=
make_tensor
(
args
[
1
].
get_shape
());
auto
x_desc
=
make_tensor
(
args
[
1
].
get_shape
());
auto
y_desc
=
make_tensor
(
output_shape
);
auto
y_desc
=
make_tensor
(
output_shape
);
miopenActivationForward
(
args
[
0
].
get
(),
ad
.
get
(),
&
alpha
,
x_desc
.
get
(),
args
[
1
].
get
(),
&
beta
,
y_desc
.
get
(),
args
[
2
].
get
());
miopenActivationForward
(
args
[
0
].
get
(),
ad
.
get
(),
&
alpha
,
x_desc
.
get
(),
args
[
1
].
get
(),
&
beta
,
y_desc
.
get
(),
args
[
2
].
get
());
return
args
[
2
];
return
args
[
2
];
}
}
...
@@ -192,7 +197,7 @@ struct miopen_apply
...
@@ -192,7 +197,7 @@ struct miopen_apply
instruction_ref
insert_allocation
(
instruction_ref
ins
,
const
shape
&
s
)
instruction_ref
insert_allocation
(
instruction_ref
ins
,
const
shape
&
s
)
{
{
if
(
ins
==
--
prog
->
end
())
if
(
ins
==
--
prog
->
end
())
{
{
return
prog
->
add_parameter
(
"output"
,
s
);
return
prog
->
add_parameter
(
"output"
,
s
);
}
}
...
@@ -211,7 +216,12 @@ struct miopen_apply
...
@@ -211,7 +216,12 @@ struct miopen_apply
auto
cd
=
make_conv
(
op
);
auto
cd
=
make_conv
(
op
);
auto
output
=
insert_allocation
(
ins
,
ins
->
result
);
auto
output
=
insert_allocation
(
ins
,
ins
->
result
);
prog
->
replace_instruction
(
ins
,
miopen_convolution
{
op
,
std
::
move
(
cd
)},
handle
,
ins
->
arguments
.
at
(
0
),
ins
->
arguments
.
at
(
1
),
output
);
prog
->
replace_instruction
(
ins
,
miopen_convolution
{
op
,
std
::
move
(
cd
)},
handle
,
ins
->
arguments
.
at
(
0
),
ins
->
arguments
.
at
(
1
),
output
);
}
}
void
apply_activation
(
instruction_ref
ins
)
void
apply_activation
(
instruction_ref
ins
)
...
@@ -221,7 +231,8 @@ struct miopen_apply
...
@@ -221,7 +231,8 @@ struct miopen_apply
if
(
op
.
mode
==
"relu"
)
if
(
op
.
mode
==
"relu"
)
{
{
auto
output
=
insert_allocation
(
ins
,
ins
->
result
);
auto
output
=
insert_allocation
(
ins
,
ins
->
result
);
prog
->
replace_instruction
(
ins
,
miopen_relu
{
std
::
move
(
ad
)},
handle
,
ins
->
arguments
.
at
(
0
),
output
);
prog
->
replace_instruction
(
ins
,
miopen_relu
{
std
::
move
(
ad
)},
handle
,
ins
->
arguments
.
at
(
0
),
output
);
}
}
}
}
};
};
...
...
test/miopen.cpp
View file @
263579c2
...
@@ -33,7 +33,7 @@ hip_ptr hip_allocate(std::size_t sz)
...
@@ -33,7 +33,7 @@ hip_ptr hip_allocate(std::size_t sz)
return
hip_ptr
{
result
};
return
hip_ptr
{
result
};
}
}
template
<
class
T
>
template
<
class
T
>
hip_ptr
write
(
const
T
&
x
)
hip_ptr
write
(
const
T
&
x
)
{
{
using
type
=
typename
T
::
value_type
;
using
type
=
typename
T
::
value_type
;
...
@@ -44,7 +44,7 @@ hip_ptr write(const T& x)
...
@@ -44,7 +44,7 @@ hip_ptr write(const T& x)
return
result
;
return
result
;
}
}
template
<
class
T
>
template
<
class
T
>
std
::
vector
<
T
>
read
(
const
hip_ptr
&
x
,
std
::
size_t
sz
)
std
::
vector
<
T
>
read
(
const
hip_ptr
&
x
,
std
::
size_t
sz
)
{
{
std
::
vector
<
T
>
result
(
sz
);
std
::
vector
<
T
>
result
(
sz
);
...
@@ -92,10 +92,7 @@ std::vector<float> cpu()
...
@@ -92,10 +92,7 @@ std::vector<float> cpu()
auto
x
=
get_tensor_argument_cpu
({
rtg
::
shape
::
float_type
,
{
4
,
3
,
3
,
3
}});
auto
x
=
get_tensor_argument_cpu
({
rtg
::
shape
::
float_type
,
{
4
,
3
,
3
,
3
}});
auto
w
=
get_tensor_argument_cpu
({
rtg
::
shape
::
float_type
,
{
4
,
3
,
3
,
3
}});
auto
w
=
get_tensor_argument_cpu
({
rtg
::
shape
::
float_type
,
{
4
,
3
,
3
,
3
}});
p
.
compile
(
rtg
::
cpu
::
cpu_target
{});
p
.
compile
(
rtg
::
cpu
::
cpu_target
{});
auto
r
=
p
.
eval
({
auto
r
=
p
.
eval
({{
"x"
,
x
},
{
"w"
,
w
}});
{
"x"
,
x
},
{
"w"
,
w
}
});
r
.
visit
([
&
](
auto
output
)
{
result
.
assign
(
output
.
begin
(),
output
.
end
());
});
r
.
visit
([
&
](
auto
output
)
{
result
.
assign
(
output
.
begin
(),
output
.
end
());
});
return
result
;
return
result
;
}
}
...
@@ -109,12 +106,8 @@ std::vector<float> gpu()
...
@@ -109,12 +106,8 @@ std::vector<float> gpu()
p
.
compile
(
rtg
::
miopen
::
miopen_target
{});
p
.
compile
(
rtg
::
miopen
::
miopen_target
{});
auto
y
=
get_tensor_argument_gpu
(
p
.
get_parameter_shape
(
"output"
));
auto
y
=
get_tensor_argument_gpu
(
p
.
get_parameter_shape
(
"output"
));
auto
handle
=
make_obj
<
miopen_handle
>
(
&
miopenCreate
);
auto
handle
=
make_obj
<
miopen_handle
>
(
&
miopenCreate
);
auto
r
=
p
.
eval
({
auto
r
=
p
.
eval
(
{
"x"
,
x
},
{{
"x"
,
x
},
{
"w"
,
w
},
{
"output"
,
y
},
{
"handle"
,
{
rtg
::
shape
::
any_type
,
handle
.
get
()}}});
{
"w"
,
w
},
{
"output"
,
y
},
{
"handle"
,
{
rtg
::
shape
::
any_type
,
handle
.
get
()}}
});
r
.
visit
([
&
](
auto
output
)
{
result
.
assign
(
output
.
begin
(),
output
.
end
());
});
r
.
visit
([
&
](
auto
output
)
{
result
.
assign
(
output
.
begin
(),
output
.
end
());
});
return
result
;
return
result
;
}
}
...
@@ -123,11 +116,8 @@ void test1()
...
@@ -123,11 +116,8 @@ void test1()
{
{
auto
x
=
cpu
();
auto
x
=
cpu
();
auto
y
=
gpu
();
auto
y
=
gpu
();
if
(
x
==
y
)
if
(
x
==
y
)
printf
(
"FAILED
\n
"
);
printf
(
"FAILED
\n
"
);
}
}
int
main
()
int
main
()
{
test1
();
}
{
test1
();
}
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