Skip to content
GitLab
Menu
Projects
Groups
Snippets
Loading...
Help
Help
Support
Community forum
Keyboard shortcuts
?
Submit feedback
Contribute to GitLab
Sign in
Toggle navigation
Menu
Open sidebar
gaoqiong
MIGraphX
Commits
a83a717a
Commit
a83a717a
authored
Aug 17, 2018
by
Paul
Browse files
Dont create intermediate argument during eval
parent
f0861b1a
Changes
6
Hide whitespace changes
Inline
Side-by-side
Showing
6 changed files
with
64 additions
and
41 deletions
+64
-41
src/include/migraph/generate.hpp
src/include/migraph/generate.hpp
+3
-2
src/include/migraph/instruction.hpp
src/include/migraph/instruction.hpp
+13
-0
src/include/migraph/instruction_ref.hpp
src/include/migraph/instruction_ref.hpp
+1
-0
src/onnx/perf_onnx.cpp
src/onnx/perf_onnx.cpp
+1
-1
src/program.cpp
src/program.cpp
+45
-38
test/gpu/miopen.cpp
test/gpu/miopen.cpp
+1
-0
No files found.
src/include/migraph/generate.hpp
View file @
a83a717a
...
@@ -10,11 +10,12 @@ namespace migraph {
...
@@ -10,11 +10,12 @@ namespace migraph {
template
<
class
T
>
template
<
class
T
>
struct
xorshf96_generator
struct
xorshf96_generator
{
{
unsigned
long
max
=
31
;
unsigned
long
x
=
123456789
;
unsigned
long
x
=
123456789
;
unsigned
long
y
=
362436069
;
unsigned
long
y
=
362436069
;
unsigned
long
z
=
521288629
;
unsigned
long
z
=
521288629
;
constexpr
T
operator
()()
constexpr
T
operator
()()
noexcept
{
{
x
^=
x
<<
16U
;
x
^=
x
<<
16U
;
x
^=
x
>>
5U
;
x
^=
x
>>
5U
;
...
@@ -25,7 +26,7 @@ struct xorshf96_generator
...
@@ -25,7 +26,7 @@ struct xorshf96_generator
y
=
z
;
y
=
z
;
z
=
t
^
x
^
y
;
z
=
t
^
x
^
y
;
return
z
;
return
z
%
max
;
}
}
};
};
...
...
src/include/migraph/instruction.hpp
View file @
a83a717a
...
@@ -165,4 +165,17 @@ inline shape compute_shape(operation op, std::vector<instruction_ref> args)
...
@@ -165,4 +165,17 @@ inline shape compute_shape(operation op, std::vector<instruction_ref> args)
}
// namespace migraph
}
// namespace migraph
namespace
std
{
template
<
>
struct
hash
<
migraph
::
instruction_ref
>
{
using
argument_type
=
migraph
::
instruction_ref
;
using
result_type
=
std
::
size_t
;
result_type
operator
()(
const
argument_type
&
x
)
const
noexcept
{
return
std
::
hash
<
migraph
::
instruction
*>
{}(
&*
x
);
}
};
}
// namespace std
#endif
#endif
src/include/migraph/instruction_ref.hpp
View file @
a83a717a
...
@@ -2,6 +2,7 @@
...
@@ -2,6 +2,7 @@
#define MIGRAPH_GUARD_INSTRUCTION_REF_HPP
#define MIGRAPH_GUARD_INSTRUCTION_REF_HPP
#include <list>
#include <list>
#include <functional>
namespace
migraph
{
namespace
migraph
{
...
...
src/onnx/perf_onnx.cpp
View file @
a83a717a
...
@@ -30,6 +30,6 @@ int main(int argc, char const* argv[])
...
@@ -30,6 +30,6 @@ int main(int argc, char const* argv[])
std
::
cout
<<
"Allocating params ... "
<<
std
::
endl
;
std
::
cout
<<
"Allocating params ... "
<<
std
::
endl
;
auto
m
=
create_param_map
(
p
);
auto
m
=
create_param_map
(
p
);
std
::
cout
<<
"Running performance report ... "
<<
std
::
endl
;
std
::
cout
<<
"Running performance report ... "
<<
std
::
endl
;
p
.
perf_report
(
std
::
cout
,
1
0
,
m
);
p
.
perf_report
(
std
::
cout
,
5
0
,
m
);
}
}
}
}
src/program.cpp
View file @
a83a717a
...
@@ -3,6 +3,7 @@
...
@@ -3,6 +3,7 @@
#include <migraph/instruction.hpp>
#include <migraph/instruction.hpp>
#include <migraph/env.hpp>
#include <migraph/env.hpp>
#include <migraph/time.hpp>
#include <migraph/time.hpp>
#include <migraph/iterator_for.hpp>
#include <iostream>
#include <iostream>
#include <sstream>
#include <sstream>
#include <algorithm>
#include <algorithm>
...
@@ -23,48 +24,48 @@ const operation& get_operation(instruction_ref ins) { return ins->op; }
...
@@ -23,48 +24,48 @@ const operation& get_operation(instruction_ref ins) { return ins->op; }
template
<
class
F
>
template
<
class
F
>
static
void
print_program
(
std
::
ostream
&
os
,
const
program
&
p
,
F
annonate
)
static
void
print_program
(
std
::
ostream
&
os
,
const
program
&
p
,
F
annonate
)
{
{
std
::
unordered_map
<
const
instruction
*
,
std
::
string
>
names
;
std
::
unordered_map
<
instruction
_ref
,
std
::
string
>
names
;
int
count
=
0
;
int
count
=
0
;
for
(
auto
&
ins
:
p
)
for
(
auto
ins
:
iterator_for
(
p
)
)
{
{
std
::
string
var_name
=
"@"
+
std
::
to_string
(
count
);
std
::
string
var_name
=
"@"
+
std
::
to_string
(
count
);
if
(
ins
.
op
.
name
()
==
"@param"
)
if
(
ins
->
op
.
name
()
==
"@param"
)
{
{
var_name
=
any_cast
<
builtin
::
param
>
(
ins
.
op
).
parameter
;
var_name
=
any_cast
<
builtin
::
param
>
(
ins
->
op
).
parameter
;
}
}
os
<<
var_name
<<
" = "
;
os
<<
var_name
<<
" = "
;
os
<<
ins
.
op
;
os
<<
ins
->
op
;
if
(
ins
.
op
.
name
()
==
"@literal"
)
if
(
ins
->
op
.
name
()
==
"@literal"
)
{
{
if
(
ins
.
lit
.
get_shape
().
elements
()
>
10
)
if
(
ins
->
lit
.
get_shape
().
elements
()
>
10
)
os
<<
"{ ... }"
;
os
<<
"{ ... }"
;
else
else
os
<<
"{"
<<
ins
.
lit
<<
"}"
;
os
<<
"{"
<<
ins
->
lit
<<
"}"
;
}
}
if
(
!
ins
.
arguments
.
empty
())
if
(
!
ins
->
arguments
.
empty
())
{
{
char
delim
=
'('
;
char
delim
=
'('
;
for
(
auto
&&
arg
:
ins
.
arguments
)
for
(
auto
&&
arg
:
ins
->
arguments
)
{
{
assert
(
p
.
has_instruction
(
arg
)
&&
"Instruction not found"
);
assert
(
p
.
has_instruction
(
arg
)
&&
"Instruction not found"
);
os
<<
delim
<<
names
.
at
(
std
::
addressof
(
*
arg
)
)
;
os
<<
delim
<<
names
.
at
(
arg
);
delim
=
','
;
delim
=
','
;
}
}
os
<<
")"
;
os
<<
")"
;
}
}
os
<<
" -> "
<<
ins
.
result
;
os
<<
" -> "
<<
ins
->
result
;
annonate
(
ins
,
names
);
annonate
(
ins
,
names
);
os
<<
std
::
endl
;
os
<<
std
::
endl
;
names
.
emplace
(
std
::
addressof
(
ins
)
,
var_name
);
names
.
emplace
(
ins
,
var_name
);
count
++
;
count
++
;
}
}
}
}
...
@@ -276,42 +277,44 @@ argument generic_eval(const program& p,
...
@@ -276,42 +277,44 @@ argument generic_eval(const program& p,
F
trace
)
F
trace
)
{
{
assert
(
p
.
validate
()
==
p
.
end
());
assert
(
p
.
validate
()
==
p
.
end
());
std
::
unordered_map
<
const
instruction
*
,
argument
>
results
;
std
::
unordered_map
<
instruction_ref
,
argument
>
results
;
results
.
reserve
(
p
.
size
());
results
.
reserve
(
p
.
size
()
*
2
);
argument
result
;
std
::
vector
<
argument
>
values
;
std
::
vector
<
argument
>
values
;
values
.
reserve
(
16
);
values
.
reserve
(
16
);
for
(
auto
&
ins
:
p
)
for
(
auto
ins
:
iterator_for
(
p
)
)
{
{
if
(
ins
.
op
.
name
()
==
"@literal"
)
if
(
ins
->
op
.
name
()
==
"@literal"
)
{
{
trace
(
ins
,
[
&
]
{
re
sult
=
ins
.
lit
.
get_argument
();
});
results
.
emplace
(
ins
,
trace
(
ins
,
[
&
]
{
re
turn
ins
->
lit
.
get_argument
();
})
)
;
}
}
else
if
(
ins
.
op
.
name
()
==
"@param"
)
else
if
(
ins
->
op
.
name
()
==
"@param"
)
{
{
trace
(
ins
,
[
&
]
{
re
sult
=
params
.
at
(
any_cast
<
builtin
::
param
>
(
ins
.
op
).
parameter
);
});
results
.
emplace
(
ins
,
trace
(
ins
,
[
&
]
{
re
turn
params
.
at
(
any_cast
<
builtin
::
param
>
(
ins
->
op
).
parameter
);
})
)
;
}
}
else
if
(
ins
.
op
.
name
()
==
"@outline"
)
else
if
(
ins
->
op
.
name
()
==
"@outline"
)
{
{
trace
(
ins
,
[
&
]
{
re
sult
=
argument
{
ins
.
result
,
nullptr
};
});
results
.
emplace
(
ins
,
trace
(
ins
,
[
&
]
{
re
turn
argument
{
ins
->
result
,
nullptr
};
})
)
;
}
}
else
else
{
{
values
.
resize
(
ins
.
arguments
.
size
());
values
.
resize
(
ins
->
arguments
.
size
());
std
::
transform
(
ins
.
arguments
.
begin
(),
std
::
transform
(
ins
->
arguments
.
begin
(),
ins
.
arguments
.
end
(),
ins
->
arguments
.
end
(),
values
.
begin
(),
values
.
begin
(),
[
&
](
instruction_ref
i
)
{
return
results
.
at
(
std
::
addressof
(
*
i
));
});
[
&
](
instruction_ref
i
)
{
trace
(
ins
,
[
&
]
{
result
=
ins
.
op
.
compute
(
ctx
,
ins
.
result
,
values
);
});
assert
(
results
.
find
(
i
)
!=
results
.
end
());
return
results
[
i
];
});
results
.
emplace
(
ins
,
trace
(
ins
,
[
&
]
{
return
ins
->
op
.
compute
(
ctx
,
ins
->
result
,
values
);
}));
}
}
results
.
emplace
(
std
::
addressof
(
ins
)
,
result
);
assert
(
results
.
find
(
ins
)
!=
result
s
.
end
()
);
}
}
return
result
;
return
result
s
.
at
(
std
::
prev
(
p
.
end
()))
;
}
}
argument
program
::
eval
(
std
::
unordered_map
<
std
::
string
,
argument
>
params
)
const
argument
program
::
eval
(
std
::
unordered_map
<
std
::
string
,
argument
>
params
)
const
{
{
return
generic_eval
(
*
this
,
this
->
impl
->
ctx
,
params
,
[](
auto
&
,
auto
f
)
{
f
();
});
return
generic_eval
(
*
this
,
this
->
impl
->
ctx
,
params
,
[](
auto
&
,
auto
f
)
{
return
f
();
});
}
}
void
program
::
perf_report
(
std
::
ostream
&
os
,
std
::
size_t
n
,
parameter_map
params
)
const
void
program
::
perf_report
(
std
::
ostream
&
os
,
std
::
size_t
n
,
parameter_map
params
)
const
...
@@ -325,15 +328,19 @@ void program::perf_report(std::ostream& os, std::size_t n, parameter_map params)
...
@@ -325,15 +328,19 @@ void program::perf_report(std::ostream& os, std::size_t n, parameter_map params)
{
{
total_acc
+=
time
<
milliseconds
>
([
&
]
{
eval
(
params
);
});
total_acc
+=
time
<
milliseconds
>
([
&
]
{
eval
(
params
);
});
}
}
std
::
unordered_map
<
const
instruction
*
,
double
>
ins_acc
;
std
::
unordered_map
<
instruction
_ref
,
double
>
ins_acc
;
// Fill the map
// Fill the map
generic_eval
(
generic_eval
(
*
this
,
this
->
impl
->
ctx
,
params
,
[
&
](
auto
&
ins
,
auto
)
{
ins_acc
[
std
::
addressof
(
ins
)
]
=
0
;
});
*
this
,
this
->
impl
->
ctx
,
params
,
[
&
](
auto
ins
,
auto
)
{
ins_acc
[
ins
]
=
0
;
return
argument
{};
});
// Run and time each instruction
// Run and time each instruction
for
(
std
::
size_t
i
=
0
;
i
<
n
;
i
++
)
for
(
std
::
size_t
i
=
0
;
i
<
n
;
i
++
)
{
{
generic_eval
(
*
this
,
this
->
impl
->
ctx
,
params
,
[
&
](
auto
&
ins
,
auto
f
)
{
generic_eval
(
*
this
,
this
->
impl
->
ctx
,
params
,
[
&
](
auto
ins
,
auto
f
)
{
ins_acc
[
std
::
addressof
(
ins
)]
+=
time
<
milliseconds
>
(
f
);
argument
result
;
ins_acc
[
ins
]
+=
time
<
milliseconds
>
([
&
]{
result
=
f
();
});
return
result
;
});
});
}
}
// Run and time implicit overhead
// Run and time implicit overhead
...
@@ -341,7 +348,7 @@ void program::perf_report(std::ostream& os, std::size_t n, parameter_map params)
...
@@ -341,7 +348,7 @@ void program::perf_report(std::ostream& os, std::size_t n, parameter_map params)
for
(
std
::
size_t
i
=
0
;
i
<
n
;
i
++
)
for
(
std
::
size_t
i
=
0
;
i
<
n
;
i
++
)
{
{
overhead_acc
+=
time
<
milliseconds
>
(
overhead_acc
+=
time
<
milliseconds
>
(
[
&
]
{
generic_eval
(
*
this
,
this
->
impl
->
ctx
,
params
,
[](
auto
&&
...)
{});
});
[
&
]
{
generic_eval
(
*
this
,
this
->
impl
->
ctx
,
params
,
[](
auto
...)
{
return
argument
{};
});
});
}
}
double
total_time
=
total_acc
/
n
;
double
total_time
=
total_acc
/
n
;
...
@@ -353,8 +360,8 @@ void program::perf_report(std::ostream& os, std::size_t n, parameter_map params)
...
@@ -353,8 +360,8 @@ void program::perf_report(std::ostream& os, std::size_t n, parameter_map params)
double
calculate_overhead_time
=
total_time
-
total_instruction_time
;
double
calculate_overhead_time
=
total_time
-
total_instruction_time
;
double
calculate_overhead_percent
=
calculate_overhead_time
*
100.0
/
total_time
;
double
calculate_overhead_percent
=
calculate_overhead_time
*
100.0
/
total_time
;
print_program
(
os
,
*
this
,
[
&
](
auto
&
ins
,
auto
&&
)
{
print_program
(
os
,
*
this
,
[
&
](
auto
ins
,
auto
&&
)
{
os
<<
": "
<<
ins_acc
[
std
::
addressof
(
ins
)
]
/
n
<<
"ms"
;
os
<<
": "
<<
ins_acc
[
ins
]
/
n
<<
"ms"
;
});
});
os
<<
"Total time: "
<<
total_time
<<
"ms"
<<
std
::
endl
;
os
<<
"Total time: "
<<
total_time
<<
"ms"
<<
std
::
endl
;
...
...
test/gpu/miopen.cpp
View file @
a83a717a
...
@@ -104,6 +104,7 @@ void verify_program()
...
@@ -104,6 +104,7 @@ void verify_program()
visit_all
(
cpu_arg_f
.
get
(),
gpu_arg
)([](
auto
cpu
,
auto
gpu
)
{
visit_all
(
cpu_arg_f
.
get
(),
gpu_arg
)([](
auto
cpu
,
auto
gpu
)
{
if
(
not
migraph
::
verify_range
(
cpu
,
gpu
))
if
(
not
migraph
::
verify_range
(
cpu
,
gpu
))
{
{
// TODO: Check for nans
std
::
cout
<<
"FAILED: "
<<
migraph
::
get_type_name
<
V
>
()
<<
std
::
endl
;
std
::
cout
<<
"FAILED: "
<<
migraph
::
get_type_name
<
V
>
()
<<
std
::
endl
;
}
}
});
});
...
...
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