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
c1d244d9
"git@developer.sourcefind.cn:modelzoo/resnet50_tensorflow.git" did not exist on "59c7683e997a97363777a0600740c4a982421ab6"
Commit
c1d244d9
authored
Jun 17, 2019
by
Paul
Browse files
Formatting
parent
cf8cc835
Changes
2
Hide whitespace changes
Inline
Side-by-side
Showing
2 changed files
with
16 additions
and
21 deletions
+16
-21
src/targets/gpu/device/include/migraphx/gpu/device/nary.hpp
src/targets/gpu/device/include/migraphx/gpu/device/nary.hpp
+6
-8
src/targets/gpu/device/include/migraphx/gpu/device/tensor.hpp
...targets/gpu/device/include/migraphx/gpu/device/tensor.hpp
+10
-13
No files found.
src/targets/gpu/device/include/migraphx/gpu/device/nary.hpp
View file @
c1d244d9
...
@@ -43,7 +43,7 @@ auto nary_nonstandard_impl(hipStream_t stream, F f, argument result, Arguments..
...
@@ -43,7 +43,7 @@ auto nary_nonstandard_impl(hipStream_t stream, F f, argument result, Arguments..
std
::
size_t
nelements
=
result
.
get_shape
().
elements
();
std
::
size_t
nelements
=
result
.
get_shape
().
elements
();
hip_visit_all
(
result
,
args
...)([
&
](
auto
output
,
auto
...
inputs
)
{
hip_visit_all
(
result
,
args
...)([
&
](
auto
output
,
auto
...
inputs
)
{
gs_launch
(
stream
,
nelements
)([
=
](
auto
i
)
{
gs_launch
(
stream
,
nelements
)([
=
](
auto
i
)
{
auto
idx
=
output
.
get_shape
().
multi
(
i
);
auto
idx
=
output
.
get_shape
().
multi
(
i
);
output
[
i
]
=
f
(
inputs
[
idx
]...);
output
[
i
]
=
f
(
inputs
[
idx
]...);
});
});
});
});
...
@@ -93,8 +93,8 @@ void trinary_broadcast_vec_impl(hipStream_t stream,
...
@@ -93,8 +93,8 @@ void trinary_broadcast_vec_impl(hipStream_t stream,
// Process the data
// Process the data
for
(
size_t
i
=
idx
.
global
;
i
<
n
;
i
+=
nglobal
)
for
(
size_t
i
=
idx
.
global
;
i
<
n
;
i
+=
nglobal
)
{
{
auto
bidx
=
((
i
*
vec_size
)
%
bdim_next_stride
)
/
bdim_stride
;
auto
bidx
=
((
i
*
vec_size
)
%
bdim_next_stride
)
/
bdim_stride
;
auto
b
=
bp
[
bidx
];
auto
b
=
bp
[
bidx
];
vec
<
type
,
4
>
x
=
xp
[
i
];
vec
<
type
,
4
>
x
=
xp
[
i
];
vec
<
type
,
4
>
y
=
yp
[
i
];
vec
<
type
,
4
>
y
=
yp
[
i
];
vec
<
type
,
4
>
out
=
outp
[
i
];
vec
<
type
,
4
>
out
=
outp
[
i
];
...
@@ -198,8 +198,8 @@ void binary_broadcast_vec_impl(
...
@@ -198,8 +198,8 @@ void binary_broadcast_vec_impl(
// Process the data
// Process the data
for
(
size_t
i
=
idx
.
global
;
i
<
n
;
i
+=
nglobal
)
for
(
size_t
i
=
idx
.
global
;
i
<
n
;
i
+=
nglobal
)
{
{
auto
bidx
=
((
i
*
vec_size
)
%
bdim_next_stride
)
/
bdim_stride
;
auto
bidx
=
((
i
*
vec_size
)
%
bdim_next_stride
)
/
bdim_stride
;
auto
b
=
bp
[
bidx
];
auto
b
=
bp
[
bidx
];
vec
<
type
,
4
>
x
=
xp
[
i
];
vec
<
type
,
4
>
x
=
xp
[
i
];
vec
<
type
,
4
>
out
=
outp
[
i
];
vec
<
type
,
4
>
out
=
outp
[
i
];
for
(
std
::
size_t
j
=
0
;
j
<
vec_size
;
j
++
)
for
(
std
::
size_t
j
=
0
;
j
<
vec_size
;
j
++
)
...
@@ -287,9 +287,7 @@ void nary_standard_impl(hipStream_t stream, F f, argument result, Arguments... a
...
@@ -287,9 +287,7 @@ void nary_standard_impl(hipStream_t stream, F f, argument result, Arguments... a
{
{
std
::
size_t
nelements
=
result
.
get_shape
().
elements
();
std
::
size_t
nelements
=
result
.
get_shape
().
elements
();
hip_visit_all
(
result
,
args
...)([
&
](
auto
output
,
auto
...
inputs
)
{
hip_visit_all
(
result
,
args
...)([
&
](
auto
output
,
auto
...
inputs
)
{
gs_launch
(
stream
,
nelements
)([
=
](
auto
i
)
{
gs_launch
(
stream
,
nelements
)([
=
](
auto
i
)
{
output
.
data
()[
i
]
=
f
(
inputs
.
data
()[
i
]...);
});
output
.
data
()[
i
]
=
f
(
inputs
.
data
()[
i
]...);
});
});
});
}
}
...
...
src/targets/gpu/device/include/migraphx/gpu/device/tensor.hpp
View file @
c1d244d9
...
@@ -85,7 +85,7 @@ struct hip_array
...
@@ -85,7 +85,7 @@ struct hip_array
friend
MIGRAPHX_DEVICE_CONSTEXPR
hip_array
operator
*
(
const
hip_array
&
x
,
const
hip_array
&
y
)
friend
MIGRAPHX_DEVICE_CONSTEXPR
hip_array
operator
*
(
const
hip_array
&
x
,
const
hip_array
&
y
)
{
{
hip_array
result
;
hip_array
result
;
for
(
std
::
size_t
i
=
0
;
i
<
N
;
i
++
)
for
(
std
::
size_t
i
=
0
;
i
<
N
;
i
++
)
result
[
i
]
=
x
[
i
]
*
y
[
i
];
result
[
i
]
=
x
[
i
]
*
y
[
i
];
return
result
;
return
result
;
}
}
...
@@ -167,15 +167,9 @@ struct hip_shape
...
@@ -167,15 +167,9 @@ struct hip_shape
std
::
copy
(
s
.
strides
().
begin
(),
s
.
strides
().
end
(),
strides
.
begin
());
std
::
copy
(
s
.
strides
().
begin
(),
s
.
strides
().
end
(),
strides
.
begin
());
}
}
MIGRAPHX_DEVICE_CONSTEXPR
std
::
size_t
elements
()
const
MIGRAPHX_DEVICE_CONSTEXPR
std
::
size_t
elements
()
const
{
return
lens
.
product
();
}
{
return
lens
.
product
();
}
MIGRAPHX_DEVICE_CONSTEXPR
std
::
size_t
index
(
hip_index
x
)
const
MIGRAPHX_DEVICE_CONSTEXPR
std
::
size_t
index
(
hip_index
x
)
const
{
return
x
.
dot
(
strides
);
}
{
return
x
.
dot
(
strides
);
}
MIGRAPHX_DEVICE_CONSTEXPR
std
::
size_t
index
(
std
::
initializer_list
<
std
::
size_t
>
x
)
const
MIGRAPHX_DEVICE_CONSTEXPR
std
::
size_t
index
(
std
::
initializer_list
<
std
::
size_t
>
x
)
const
{
{
...
@@ -199,7 +193,7 @@ struct hip_shape
...
@@ -199,7 +193,7 @@ struct hip_shape
const
std
::
size_t
k
=
rank
-
j
-
1
;
const
std
::
size_t
k
=
rank
-
j
-
1
;
const
std
::
size_t
stride
=
this
->
strides
[
k
];
const
std
::
size_t
stride
=
this
->
strides
[
k
];
const
std
::
size_t
len
=
this
->
lens
[
k
];
const
std
::
size_t
len
=
this
->
lens
[
k
];
const
std
::
size_t
slen
=
s
*
len
;
const
std
::
size_t
slen
=
s
*
len
;
const
std
::
size_t
idx
=
(
i
%
slen
)
/
s
;
const
std
::
size_t
idx
=
(
i
%
slen
)
/
s
;
result
+=
stride
*
idx
;
result
+=
stride
*
idx
;
s
=
slen
;
s
=
slen
;
...
@@ -224,7 +218,7 @@ struct hip_shape
...
@@ -224,7 +218,7 @@ struct hip_shape
template
<
class
T
,
std
::
size_t
N
>
template
<
class
T
,
std
::
size_t
N
>
struct
hip_tensor_view
struct
hip_tensor_view
{
{
using
value_type
=
device_type
<
T
>
;
using
value_type
=
device_type
<
T
>
;
__device__
__host__
hip_tensor_view
()
=
default
;
__device__
__host__
hip_tensor_view
()
=
default
;
__host__
hip_tensor_view
(
tensor_view
<
T
>
x
)
:
d
(
device_cast
(
x
.
data
())),
s
(
x
.
get_shape
())
{}
__host__
hip_tensor_view
(
tensor_view
<
T
>
x
)
:
d
(
device_cast
(
x
.
data
())),
s
(
x
.
get_shape
())
{}
...
@@ -234,8 +228,11 @@ struct hip_tensor_view
...
@@ -234,8 +228,11 @@ struct hip_tensor_view
MIGRAPHX_DEVICE_CONSTEXPR
value_type
*
data
()
const
{
return
d
;
}
MIGRAPHX_DEVICE_CONSTEXPR
value_type
*
data
()
const
{
return
d
;
}
template
<
class
U
>
template
<
class
U
>
MIGRAPHX_DEVICE_CONSTEXPR
value_type
&
operator
[](
U
i
)
const
{
return
d
[
s
.
index
(
i
)];
}
MIGRAPHX_DEVICE_CONSTEXPR
value_type
&
operator
[](
U
i
)
const
{
return
d
[
s
.
index
(
i
)];
}
MIGRAPHX_DEVICE_CONSTEXPR
value_type
*
begin
()
const
{
return
d
;
}
MIGRAPHX_DEVICE_CONSTEXPR
value_type
*
begin
()
const
{
return
d
;
}
...
...
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