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
84af2e9e
Commit
84af2e9e
authored
Jun 25, 2018
by
Paul
Browse files
Formatting
parent
5f04935f
Changes
4
Hide whitespace changes
Inline
Side-by-side
Showing
4 changed files
with
43 additions
and
30 deletions
+43
-30
src/shape.cpp
src/shape.cpp
+4
-2
src/targets/miopen/miopen_target.cpp
src/targets/miopen/miopen_target.cpp
+25
-21
test/include/test.hpp
test/include/test.hpp
+12
-5
test/miopen/miopen.cpp
test/miopen/miopen.cpp
+2
-2
No files found.
src/shape.cpp
View file @
84af2e9e
...
@@ -84,8 +84,10 @@ bool shape::packed() const { return this->m_packed; }
...
@@ -84,8 +84,10 @@ bool shape::packed() const { return this->m_packed; }
bool
shape
::
broadcasted
()
const
bool
shape
::
broadcasted
()
const
{
{
assert
(
this
->
lens
().
size
()
==
this
->
strides
().
size
());
assert
(
this
->
lens
().
size
()
==
this
->
strides
().
size
());
return
std
::
accumulate
(
return
std
::
accumulate
(
this
->
strides
().
begin
(),
this
->
strides
().
begin
(),
this
->
strides
().
end
(),
std
::
size_t
{
1
},
std
::
multiplies
<
std
::
size_t
>
())
==
0
;
this
->
strides
().
end
(),
std
::
size_t
{
1
},
std
::
multiplies
<
std
::
size_t
>
())
==
0
;
}
}
std
::
size_t
shape
::
element_space
()
const
std
::
size_t
shape
::
element_space
()
const
...
...
src/targets/miopen/miopen_target.cpp
View file @
84af2e9e
...
@@ -105,33 +105,37 @@ struct miopen_add
...
@@ -105,33 +105,37 @@ struct miopen_add
argument
compute
(
shape
output_shape
,
std
::
vector
<
argument
>
args
)
const
argument
compute
(
shape
output_shape
,
std
::
vector
<
argument
>
args
)
const
{
{
if
(
args
[
2
].
get_shape
().
broadcasted
())
{
if
(
args
[
2
].
get_shape
().
broadcasted
())
{
argument
result
{
output_shape
};
argument
result
{
output_shape
};
visit_all
(
result
,
from_gpu
(
args
[
1
]),
from_gpu
(
args
[
2
]))([
&
](
auto
output
,
auto
input1
,
auto
input2
)
{
visit_all
(
result
,
from_gpu
(
args
[
1
]),
from_gpu
(
args
[
2
]))(
shape_for_each
(
output
.
get_shape
(),
[
&
](
const
auto
&
idx
)
{
[
&
](
auto
output
,
auto
input1
,
auto
input2
)
{
shape_for_each
(
output
.
get_shape
(),
[
&
](
const
auto
&
idx
)
{
output
(
idx
.
begin
(),
idx
.
end
())
=
output
(
idx
.
begin
(),
idx
.
end
())
=
input1
(
idx
.
begin
(),
idx
.
end
())
+
input2
(
idx
.
begin
(),
idx
.
end
());
input1
(
idx
.
begin
(),
idx
.
end
())
+
input2
(
idx
.
begin
(),
idx
.
end
());
});
});
});
});
return
to_gpu
(
result
);
return
to_gpu
(
result
);
}
else
{
}
float
alpha
=
1
,
beta
=
0
;
else
auto
a_desc
=
make_tensor
(
args
[
1
].
get_shape
());
{
auto
b_desc
=
make_tensor
(
args
[
2
].
get_shape
());
float
alpha
=
1
,
beta
=
0
;
auto
c_desc
=
make_tensor
(
output_shape
);
auto
a_desc
=
make_tensor
(
args
[
1
].
get_shape
());
miopenOpTensor
(
args
[
0
].
implicit
(),
auto
b_desc
=
make_tensor
(
args
[
2
].
get_shape
());
miopenTensorOpAdd
,
auto
c_desc
=
make_tensor
(
output_shape
);
&
alpha
,
miopenOpTensor
(
args
[
0
].
implicit
(),
a_desc
.
get
(),
miopenTensorOpAdd
,
args
[
1
].
implicit
(),
&
alpha
,
&
alpha
,
a_desc
.
get
(),
b_desc
.
get
(),
args
[
1
].
implicit
(),
args
[
2
].
implicit
(),
&
alpha
,
&
beta
,
b_desc
.
get
(),
c_desc
.
get
(),
args
[
2
].
implicit
(),
args
[
3
].
implicit
());
&
beta
,
return
args
[
3
];
c_desc
.
get
(),
args
[
3
].
implicit
());
return
args
[
3
];
}
}
}
}
};
};
...
...
test/include/test.hpp
View file @
84af2e9e
...
@@ -114,7 +114,7 @@ struct capture
...
@@ -114,7 +114,7 @@ struct capture
};
};
template
<
class
T
,
class
F
>
template
<
class
T
,
class
F
>
void
failed
(
T
x
,
const
char
*
msg
,
const
char
*
func
,
const
char
*
file
,
int
line
,
F
f
)
void
failed
(
T
x
,
const
char
*
msg
,
const
char
*
func
,
const
char
*
file
,
int
line
,
F
f
)
{
{
if
(
!
x
.
value
())
if
(
!
x
.
value
())
{
{
...
@@ -163,11 +163,18 @@ void run_test()
...
@@ -163,11 +163,18 @@ void run_test()
}
// namespace test
}
// namespace test
// NOLINTNEXTLINE
// NOLINTNEXTLINE
#define CHECK(...) \
#define CHECK(...) \
test::failed(test::capture{}->*__VA_ARGS__, #__VA_ARGS__, __PRETTY_FUNCTION__, __FILE__, __LINE__, [] {})
test::failed( \
test::capture{}->*__VA_ARGS__, #__VA_ARGS__, __PRETTY_FUNCTION__, __FILE__, __LINE__, [] { \
})
// NOLINTNEXTLINE
// NOLINTNEXTLINE
#define EXPECT(...) \
#define EXPECT(...) \
test::failed(test::capture{}->*__VA_ARGS__, #__VA_ARGS__, __PRETTY_FUNCTION__, __FILE__, __LINE__, &std::abort)
test::failed(test::capture{}->*__VA_ARGS__, \
#__VA_ARGS__, \
__PRETTY_FUNCTION__, \
__FILE__, \
__LINE__, \
&std::abort)
// NOLINTNEXTLINE
// NOLINTNEXTLINE
#define STATUS(...) EXPECT((__VA_ARGS__) == 0)
#define STATUS(...) EXPECT((__VA_ARGS__) == 0)
...
...
test/miopen/miopen.cpp
View file @
84af2e9e
...
@@ -77,8 +77,8 @@ struct test_add_broadcast
...
@@ -77,8 +77,8 @@ struct test_add_broadcast
{
{
rtg
::
program
p
;
rtg
::
program
p
;
rtg
::
shape
s
{
rtg
::
shape
::
float_type
,
{
3
}};
rtg
::
shape
s
{
rtg
::
shape
::
float_type
,
{
3
}};
auto
x
=
p
.
add_parameter
(
"x"
,
{
rtg
::
shape
::
float_type
,
{
2
,
2
,
3
}});
auto
x
=
p
.
add_parameter
(
"x"
,
{
rtg
::
shape
::
float_type
,
{
2
,
2
,
3
}});
auto
y
=
p
.
add_parameter
(
"y"
,
{
rtg
::
shape
::
float_type
,
{
2
,
2
}});
auto
y
=
p
.
add_parameter
(
"y"
,
{
rtg
::
shape
::
float_type
,
{
2
,
2
}});
auto
by
=
p
.
add_instruction
(
rtg
::
broadcast
{
0
},
x
,
y
);
auto
by
=
p
.
add_instruction
(
rtg
::
broadcast
{
0
},
x
,
y
);
p
.
add_instruction
(
rtg
::
add
{},
x
,
by
);
p
.
add_instruction
(
rtg
::
add
{},
x
,
by
);
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