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
206af368
Commit
206af368
authored
May 26, 2023
by
Paul
Browse files
Format
parent
ceee865f
Changes
3
Hide whitespace changes
Inline
Side-by-side
Showing
3 changed files
with
15 additions
and
13 deletions
+15
-13
src/targets/gpu/compile_ops.cpp
src/targets/gpu/compile_ops.cpp
+4
-3
src/targets/gpu/device/multinomial.cpp
src/targets/gpu/device/multinomial.cpp
+10
-9
src/targets/gpu/fuse_ck.cpp
src/targets/gpu/fuse_ck.cpp
+1
-1
No files found.
src/targets/gpu/compile_ops.cpp
View file @
206af368
...
...
@@ -88,8 +88,9 @@ struct problem_cache
assert
(
not
solution
.
is_null
());
cache
[
create_key
(
name
,
problem
)]
=
solution
;
}
void
mark
(
const
std
::
string
&
name
,
const
value
&
problem
)
{
cache
.
insert
(
std
::
make_pair
(
create_key
(
name
,
problem
),
value
{}));
void
mark
(
const
std
::
string
&
name
,
const
value
&
problem
)
{
cache
.
insert
(
std
::
make_pair
(
create_key
(
name
,
problem
),
value
{}));
}
optional
<
value
>
get
(
const
std
::
string
&
name
,
const
value
&
problem
)
const
{
...
...
@@ -242,7 +243,7 @@ void compile_ops::apply(module& m) const
cm
.
compile
(
m
);
// Compile already tuned configs
cm
.
compile
(
m
);
if
(
not
cm
.
cps
.
empty
())
if
(
not
cm
.
cps
.
empty
())
MIGRAPHX_THROW
(
"Untuned configs"
);
}
...
...
src/targets/gpu/device/multinomial.cpp
View file @
206af368
...
...
@@ -69,16 +69,17 @@ void multinomial(hipStream_t stream,
visit_all
(
arg0
,
arg1
)([
&
](
auto
cdf_host
,
auto
dist_host
)
{
result
.
visit
([
&
](
auto
output_host
)
{
hip_visit_views
(
cdf_host
,
dist_host
,
output_host
)([
&
](
auto
cdf
,
auto
dist
,
auto
output
)
{
gs_launch
(
stream
,
batch_size
*
sample_size
)([
=
](
auto
i
)
__device__
{
auto
idx
=
output
.
get_shape
().
multi
(
i
);
auto
cdf_begin
=
cdf
.
begin
()
+
(
idx
.
front
()
*
class_size
);
auto
cdf_end
=
cdf_begin
+
class_size
;
auto
sample_iter
=
upper_bound
(
cdf_begin
,
cdf_end
,
dist
[
i
]
*
*
(
std
::
prev
(
cdf_end
)));
output
[
i
]
=
std
::
distance
(
cdf_begin
,
sample_iter
);
hip_visit_views
(
cdf_host
,
dist_host
,
output_host
)(
[
&
](
auto
cdf
,
auto
dist
,
auto
output
)
{
gs_launch
(
stream
,
batch_size
*
sample_size
)([
=
](
auto
i
)
__device__
{
auto
idx
=
output
.
get_shape
().
multi
(
i
);
auto
cdf_begin
=
cdf
.
begin
()
+
(
idx
.
front
()
*
class_size
);
auto
cdf_end
=
cdf_begin
+
class_size
;
auto
sample_iter
=
upper_bound
(
cdf_begin
,
cdf_end
,
dist
[
i
]
*
*
(
std
::
prev
(
cdf_end
)));
output
[
i
]
=
std
::
distance
(
cdf_begin
,
sample_iter
);
});
});
});
});
});
}
...
...
src/targets/gpu/fuse_ck.cpp
View file @
206af368
...
...
@@ -84,7 +84,7 @@ struct find_ck_gemm_pointwise
auto
inputs
=
ins
->
inputs
();
auto
gemm_it
=
std
::
find
(
inputs
.
begin
(),
inputs
.
end
(),
x_ins
);
auto
gemm_idx
=
gemm_it
-
inputs
.
begin
();
if
(
ins
->
get_shape
().
type
()
!=
gemm_ins
->
get_shape
().
type
())
if
(
ins
->
get_shape
().
type
()
!=
gemm_ins
->
get_shape
().
type
())
return
;
assert
(
gemm_it
!=
inputs
.
end
());
if
(
gemm_idx
!=
0
)
...
...
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