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
9f6153b2
Commit
9f6153b2
authored
Aug 02, 2019
by
Shucai Xiao
Browse files
fix code review comments
parent
5645b165
Changes
7
Hide whitespace changes
Inline
Side-by-side
Showing
7 changed files
with
12 additions
and
57 deletions
+12
-57
src/include/migraphx/generate.hpp
src/include/migraphx/generate.hpp
+1
-2
src/include/migraphx/op/convert.hpp
src/include/migraphx/op/convert.hpp
+2
-17
src/quantization.cpp
src/quantization.cpp
+2
-9
src/targets/gpu/convert.cpp
src/targets/gpu/convert.cpp
+1
-1
src/targets/gpu/device/convert.cpp
src/targets/gpu/device/convert.cpp
+3
-20
src/targets/gpu/include/migraphx/gpu/device/convert.hpp
src/targets/gpu/include/migraphx/gpu/device/convert.hpp
+1
-6
test/gpu/miopen.cpp
test/gpu/miopen.cpp
+2
-2
No files found.
src/include/migraphx/generate.hpp
View file @
9f6153b2
...
@@ -82,8 +82,7 @@ std::vector<T> generate_tensor_data(const migraphx::shape& s, unsigned long seed
...
@@ -82,8 +82,7 @@ std::vector<T> generate_tensor_data(const migraphx::shape& s, unsigned long seed
{
{
std
::
vector
<
T
>
result
(
s
.
elements
());
std
::
vector
<
T
>
result
(
s
.
elements
());
std
::
generate
(
result
.
begin
(),
result
.
end
(),
xorshf96_generator
<
T
>
{
seed
});
std
::
generate
(
result
.
begin
(),
result
.
end
(),
xorshf96_generator
<
T
>
{
seed
});
// divide a value to avoid integer overflow
std
::
transform
(
result
.
begin
(),
result
.
end
(),
result
.
begin
(),
[](
auto
i
)
{
return
i
;
});
std
::
transform
(
result
.
begin
(),
result
.
end
(),
result
.
begin
(),
[](
auto
i
)
{
return
i
/
32
;
});
// std::generate(result.begin(), result.end(), [&]{ return seed % 7; });
// std::generate(result.begin(), result.end(), [&]{ return seed % 7; });
// std::generate(result.begin(), result.end(), []{ return 1; });
// std::generate(result.begin(), result.end(), []{ return 1; });
return
result
;
return
result
;
...
...
src/include/migraphx/op/convert.hpp
View file @
9f6153b2
...
@@ -20,14 +20,11 @@ namespace op {
...
@@ -20,14 +20,11 @@ namespace op {
struct
convert
:
unary
<
convert
>
struct
convert
:
unary
<
convert
>
{
{
shape
::
type_t
target_type
=
shape
::
half_type
;
shape
::
type_t
target_type
=
shape
::
half_type
;
float
scale
=
1.0
f
;
float
shift
=
0.0
f
;
template
<
class
Self
,
class
F
>
template
<
class
Self
,
class
F
>
static
auto
reflect
(
Self
&
self
,
F
f
)
static
auto
reflect
(
Self
&
self
,
F
f
)
{
{
return
pack
(
return
pack
(
f
(
self
.
target_type
,
"target_type"
));
f
(
self
.
target_type
,
"target_type"
),
f
(
self
.
scale
,
"scale"
),
f
(
self
.
shift
,
"shift"
));
}
}
shape
compute_shape
(
std
::
vector
<
shape
>
inputs
)
const
shape
compute_shape
(
std
::
vector
<
shape
>
inputs
)
const
...
@@ -38,22 +35,10 @@ struct convert : unary<convert>
...
@@ -38,22 +35,10 @@ struct convert : unary<convert>
auto
apply
()
const
auto
apply
()
const
{
{
return
[
&
](
auto
x
)
{
return
[](
auto
x
)
{
return
x
;
};
float
res
=
scale
*
x
+
shift
;
if
(
target_type
==
shape
::
int8_type
)
{
int
factor
=
(
res
>=
0.0
f
)
?
1
:
-
1
;
res
=
res
+
factor
*
0.5
f
;
res
=
res
>
127.0
f
?
127.0
f
:
res
;
res
=
res
<
-
128.0
f
?
-
128.0
f
:
res
;
}
return
res
;
};
}
}
convert
(
shape
::
type_t
t
)
:
target_type
{
t
}
{}
convert
(
shape
::
type_t
t
)
:
target_type
{
t
}
{}
convert
(
shape
::
type_t
t
,
float
sle
,
float
sft
)
:
target_type
{
t
},
scale
{
sle
},
shift
{
sft
}
{}
convert
()
{}
convert
()
{}
};
};
...
...
src/quantization.cpp
View file @
9f6153b2
...
@@ -23,9 +23,7 @@ inline namespace MIGRAPHX_INLINE_NS {
...
@@ -23,9 +23,7 @@ inline namespace MIGRAPHX_INLINE_NS {
instruction_ref
insert_quant_ins
(
program
&
prog
,
instruction_ref
insert_quant_ins
(
program
&
prog
,
instruction_ref
&
ins
,
instruction_ref
&
ins
,
shape
::
type_t
type
,
shape
::
type_t
type
,
std
::
unordered_map
<
instruction_ref
,
instruction_ref
>&
map_ins
,
std
::
unordered_map
<
instruction_ref
,
instruction_ref
>&
map_ins
)
float
scale
=
1.0
f
,
float
shift
=
0.0
f
)
{
{
if
(
map_ins
.
count
(
ins
)
>
0
)
if
(
map_ins
.
count
(
ins
)
>
0
)
{
{
...
@@ -37,16 +35,11 @@ instruction_ref insert_quant_ins(program& prog,
...
@@ -37,16 +35,11 @@ instruction_ref insert_quant_ins(program& prog,
return
ins
;
return
ins
;
}
}
if
(
scale
<
0.0
f
)
{
MIGRAPHX_THROW
(
"INSERT_QUANT_INS: scale less than 0"
);
}
assert
(
ins
->
get_shape
().
type
()
==
shape
::
float_type
||
assert
(
ins
->
get_shape
().
type
()
==
shape
::
float_type
||
ins
->
get_shape
().
type
()
==
shape
::
double_type
||
ins
->
get_shape
().
type
()
==
shape
::
double_type
||
ins
->
get_shape
().
type
()
==
shape
::
int32_type
);
ins
->
get_shape
().
type
()
==
shape
::
int32_type
);
instruction_ref
quant_ins
{};
instruction_ref
quant_ins
{};
quant_ins
=
prog
.
insert_instruction
(
std
::
next
(
ins
),
op
::
convert
{
type
,
scale
,
shift
},
ins
);
quant_ins
=
prog
.
insert_instruction
(
std
::
next
(
ins
),
op
::
convert
{
type
},
ins
);
map_ins
[
ins
]
=
quant_ins
;
map_ins
[
ins
]
=
quant_ins
;
return
quant_ins
;
return
quant_ins
;
...
...
src/targets/gpu/convert.cpp
View file @
9f6153b2
...
@@ -15,7 +15,7 @@ shape hip_convert::compute_shape(std::vector<shape> inputs) const
...
@@ -15,7 +15,7 @@ shape hip_convert::compute_shape(std::vector<shape> inputs) const
argument
hip_convert
::
compute
(
context
&
ctx
,
const
shape
&
,
const
std
::
vector
<
argument
>&
args
)
const
argument
hip_convert
::
compute
(
context
&
ctx
,
const
shape
&
,
const
std
::
vector
<
argument
>&
args
)
const
{
{
device
::
convert
(
ctx
.
get_stream
().
get
(),
args
[
1
],
args
[
0
]
,
op
.
scale
,
op
.
shift
,
op
.
target_type
);
device
::
convert
(
ctx
.
get_stream
().
get
(),
args
[
1
],
args
[
0
]);
return
args
[
1
];
return
args
[
1
];
}
}
...
...
src/targets/gpu/device/convert.cpp
View file @
9f6153b2
...
@@ -6,31 +6,14 @@ inline namespace MIGRAPHX_INLINE_NS {
...
@@ -6,31 +6,14 @@ inline namespace MIGRAPHX_INLINE_NS {
namespace
gpu
{
namespace
gpu
{
namespace
device
{
namespace
device
{
void
convert
(
hipStream_t
stream
,
void
convert
(
hipStream_t
stream
,
const
argument
&
result
,
const
argument
&
arg
)
const
argument
&
result
,
const
argument
&
arg
,
float
scale
,
float
shift
,
shape
::
type_t
target_type
)
{
{
result
.
visit
([
&
](
auto
output
)
{
result
.
visit
([
&
](
auto
output
)
{
arg
.
visit
([
&
](
auto
input
)
{
arg
.
visit
([
&
](
auto
input
)
{
const
auto
*
input_ptr
=
device_cast
(
input
.
data
());
const
auto
*
input_ptr
=
device_cast
(
input
.
data
());
auto
*
output_ptr
=
device_cast
(
output
.
data
());
auto
*
output_ptr
=
device_cast
(
output
.
data
());
if
(
target_type
==
shape
::
int8_type
)
gs_launch
(
stream
,
{
result
.
get_shape
().
elements
())([
=
](
auto
i
)
{
output_ptr
[
i
]
=
input_ptr
[
i
];
});
gs_launch
(
stream
,
result
.
get_shape
().
elements
())([
=
](
auto
i
)
{
float
res
=
input_ptr
[
i
]
*
scale
+
shift
;
int
factor
=
(
res
>=
0.0
f
)
?
1
:
-
1
;
output_ptr
[
i
]
=
static_cast
<
int8_t
>
(
std
::
min
<
float
>
(
std
::
max
<
float
>
(
-
128.0
f
,
res
+
factor
*
0.5
),
127.0
f
));
});
}
else
{
gs_launch
(
stream
,
result
.
get_shape
().
elements
())(
[
=
](
auto
i
)
{
output_ptr
[
i
]
=
input_ptr
[
i
]
*
scale
+
shift
;
});
}
});
});
});
});
}
}
...
...
src/targets/gpu/include/migraphx/gpu/device/convert.hpp
View file @
9f6153b2
...
@@ -11,12 +11,7 @@ inline namespace MIGRAPHX_INLINE_NS {
...
@@ -11,12 +11,7 @@ inline namespace MIGRAPHX_INLINE_NS {
namespace
gpu
{
namespace
gpu
{
namespace
device
{
namespace
device
{
void
convert
(
hipStream_t
stream
,
void
convert
(
hipStream_t
stream
,
const
argument
&
result
,
const
argument
&
arg
);
const
argument
&
result
,
const
argument
&
arg
,
float
scale
,
float
shift
,
shape
::
type_t
target_type
);
}
// namespace device
}
// namespace device
}
// namespace gpu
}
// namespace gpu
...
...
test/gpu/miopen.cpp
View file @
9f6153b2
...
@@ -3797,9 +3797,9 @@ struct test_convert : verify_program<test_convert>
...
@@ -3797,9 +3797,9 @@ struct test_convert : verify_program<test_convert>
auto
pa
=
p
.
add_parameter
(
"a"
,
sa
);
auto
pa
=
p
.
add_parameter
(
"a"
,
sa
);
auto
pb
=
p
.
add_parameter
(
"b"
,
sb
);
auto
pb
=
p
.
add_parameter
(
"b"
,
sb
);
auto
ia
=
auto
ia
=
p
.
add_instruction
(
migraphx
::
op
::
convert
{
migraphx
::
shape
::
int8_type
,
16.0
f
,
1.0
f
},
pa
);
p
.
add_instruction
(
migraphx
::
op
::
convert
{
migraphx
::
shape
::
int8_type
},
pa
);
auto
ib
=
auto
ib
=
p
.
add_instruction
(
migraphx
::
op
::
convert
{
migraphx
::
shape
::
int8_type
,
16.0
f
,
2.0
f
},
pb
);
p
.
add_instruction
(
migraphx
::
op
::
convert
{
migraphx
::
shape
::
int8_type
},
pb
);
p
.
add_instruction
(
migraphx
::
op
::
quant_dot
{},
ia
,
ib
);
p
.
add_instruction
(
migraphx
::
op
::
quant_dot
{},
ia
,
ib
);
return
p
;
return
p
;
...
...
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