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
composable_kernel_ROCM
Commits
557a794d
"git@developer.sourcefind.cn:modelzoo/resnet50_tensorflow.git" did not exist on "8bc9fe94016cf24e9284b9835e55b1a452174540"
Commit
557a794d
authored
Apr 24, 2024
by
Harisankar Sadasivan
Browse files
modified calculating total iterations to find ceil
parent
5898ba83
Changes
1
Hide whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
4 additions
and
8 deletions
+4
-8
include/ck/host_utility/kernel_launch.hpp
include/ck/host_utility/kernel_launch.hpp
+4
-8
No files found.
include/ck/host_utility/kernel_launch.hpp
View file @
557a794d
...
@@ -21,7 +21,6 @@ float launch_and_time_kernel(const StreamConfig& stream_config,
...
@@ -21,7 +21,6 @@ float launch_and_time_kernel(const StreamConfig& stream_config,
#if CK_TIME_KERNEL
#if CK_TIME_KERNEL
if
(
stream_config
.
time_kernel_
)
if
(
stream_config
.
time_kernel_
)
{
{
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
if
(
ck
::
get_device_name
()
==
"gfx940"
||
ck
::
get_device_name
()
==
"gfx941"
||
if
(
ck
::
get_device_name
()
==
"gfx940"
||
ck
::
get_device_name
()
==
"gfx941"
||
ck
::
get_device_name
()
==
"gfx942"
)
ck
::
get_device_name
()
==
"gfx942"
)
{
{
...
@@ -47,11 +46,10 @@ float launch_and_time_kernel(const StreamConfig& stream_config,
...
@@ -47,11 +46,10 @@ float launch_and_time_kernel(const StreamConfig& stream_config,
hip_check_error
(
hipEventElapsedTime
(
&
total_time
,
start
,
stop
));
hip_check_error
(
hipEventElapsedTime
(
&
total_time
,
start
,
stop
));
total_time
/=
stream_config
.
nrepeat_
;
total_time
/=
stream_config
.
nrepeat_
;
stream_config
.
cold_niters_
=
stream_config
.
cold_niters_
=
(
stream_config
.
time_limit_ms
/
(
stream_config
.
time_limit_ms
+
total_time
-
1
)
/
total_time
)
;
// we need longer runtime to ramp up the clk on MI300s
total_time
;
// we need longer runtime to ramp up the clk on MI300s
stream_config
.
nrepeat_
=
stream_config
.
cold_niters_
;
stream_config
.
nrepeat_
=
stream_config
.
cold_niters_
;
}
}
#endif
#if DEBUG_LOG
#if DEBUG_LOG
printf
(
"%s: grid_dim {%d, %d, %d}, block_dim {%d, %d, %d}
\n
"
,
printf
(
"%s: grid_dim {%d, %d, %d}, block_dim {%d, %d, %d}
\n
"
,
__func__
,
__func__
,
...
@@ -125,7 +123,6 @@ float launch_and_time_kernel_with_preprocess(const StreamConfig& stream_config,
...
@@ -125,7 +123,6 @@ float launch_and_time_kernel_with_preprocess(const StreamConfig& stream_config,
#if CK_TIME_KERNEL
#if CK_TIME_KERNEL
if
(
stream_config
.
time_kernel_
)
if
(
stream_config
.
time_kernel_
)
{
{
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
if
(
ck
::
get_device_name
()
==
"gfx940"
||
ck
::
get_device_name
()
==
"gfx941"
||
if
(
ck
::
get_device_name
()
==
"gfx940"
||
ck
::
get_device_name
()
==
"gfx941"
||
ck
::
get_device_name
()
==
"gfx942"
)
ck
::
get_device_name
()
==
"gfx942"
)
{
{
...
@@ -151,11 +148,10 @@ float launch_and_time_kernel_with_preprocess(const StreamConfig& stream_config,
...
@@ -151,11 +148,10 @@ float launch_and_time_kernel_with_preprocess(const StreamConfig& stream_config,
hip_check_error
(
hipEventElapsedTime
(
&
total_time
,
start
,
stop
));
hip_check_error
(
hipEventElapsedTime
(
&
total_time
,
start
,
stop
));
total_time
/=
stream_config
.
nrepeat_
;
total_time
/=
stream_config
.
nrepeat_
;
stream_config
.
cold_niters_
=
stream_config
.
cold_niters_
=
(
stream_config
.
nrepeat_
/
(
stream_config
.
time_limit_ms
+
total_time
-
1
)
/
total_time
)
;
// we need longer runtime to ramp up the clk on MI300s
total_time
;
// we need longer runtime to ramp up the clk on MI300s
stream_config
.
nrepeat_
=
stream_config
.
cold_niters_
;
stream_config
.
nrepeat_
=
stream_config
.
cold_niters_
;
}
}
#endif
#if DEBUG_LOG
#if DEBUG_LOG
printf
(
"%s: grid_dim {%d, %d, %d}, block_dim {%d, %d, %d}
\n
"
,
printf
(
"%s: grid_dim {%d, %d, %d}, block_dim {%d, %d, %d}
\n
"
,
__func__
,
__func__
,
...
...
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