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
db108e89
Commit
db108e89
authored
Dec 17, 2024
by
Ville Pietilä
Browse files
Fix handling of mempool size when allocating/deallocating memory.
parent
d0a846c2
Changes
1
Hide whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
8 additions
and
0 deletions
+8
-0
include/ck/utility/host_memory_allocator.hpp
include/ck/utility/host_memory_allocator.hpp
+8
-0
No files found.
include/ck/utility/host_memory_allocator.hpp
View file @
db108e89
...
...
@@ -47,6 +47,7 @@ namespace memory {
{
void
*
p
=
memory_pool_
[
sizeInBytes
].
front
();
memory_pool_
[
sizeInBytes
].
pop
();
memPoolSizeInBytes_
-=
sizeInBytes
;
return
p
;
}
void
*
p
;
...
...
@@ -65,6 +66,11 @@ namespace memory {
// If the memory pool size exceeds the maximum size, free the memory.
if
(
memPoolSizeInBytes_
>
maxMemoryPoolSizeInBytes_
)
{
if
(
enableLogging_
)
{
std
::
cout
<<
"[ MemPool ] Clearing pool queue for size "
<<
sizeInBytes
<<
std
::
endl
;
}
memPoolSizeInBytes_
-=
sizeInBytes
*
q
.
size
();
clearMemoryPoolQueue
(
q
);
}
}
...
...
@@ -84,6 +90,8 @@ namespace memory {
{
void
*
p
=
q
.
front
();
q
.
pop
();
// This performs an implicit hipDeviceSynchronize().
// Does this create a deadlock situation when grouped GEMM is used in distributed training with NCCL?
hip_check_error
(
hipHostFree
(
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