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
Commits
a2598b8a
Commit
a2598b8a
authored
Jun 15, 2023
by
rocking
Browse files
Move initialize of workspace to the run
parent
4f1dbdf5
Changes
2
Hide whitespace changes
Inline
Side-by-side
Showing
2 changed files
with
8 additions
and
4 deletions
+8
-4
example/49_maxpool2d_bwd/maxpool2d_bwd_common.hpp
example/49_maxpool2d_bwd/maxpool2d_bwd_common.hpp
+0
-3
include/ck/tensor_operation/gpu/device/impl/device_index_pool_bwd_impl.hpp
..._operation/gpu/device/impl/device_index_pool_bwd_impl.hpp
+8
-1
No files found.
example/49_maxpool2d_bwd/maxpool2d_bwd_common.hpp
View file @
a2598b8a
...
@@ -163,9 +163,6 @@ bool maxpool_bwd_test(bool do_verification,
...
@@ -163,9 +163,6 @@ bool maxpool_bwd_test(bool do_verification,
size_t
pool_bwd_workspace_sz
=
pool_bwd
.
GetWorkSpaceSize
(
pool_bwd_argument_ptr
.
get
());
size_t
pool_bwd_workspace_sz
=
pool_bwd
.
GetWorkSpaceSize
(
pool_bwd_argument_ptr
.
get
());
DeviceMem
pool_bwd_workspace_device_buf
(
pool_bwd_workspace_sz
);
DeviceMem
pool_bwd_workspace_device_buf
(
pool_bwd_workspace_sz
);
// similar to din_device_buf.SetZero()
// we need to set workspace to be zero
pool_bwd_workspace_device_buf
.
SetZero
();
pool_bwd
.
SetWorkSpacePointer
(
pool_bwd_argument_ptr
.
get
(),
pool_bwd
.
SetWorkSpacePointer
(
pool_bwd_argument_ptr
.
get
(),
pool_bwd_workspace_device_buf
.
GetDeviceBuffer
());
pool_bwd_workspace_device_buf
.
GetDeviceBuffer
());
...
...
include/ck/tensor_operation/gpu/device/impl/device_index_pool_bwd_impl.hpp
View file @
a2598b8a
...
@@ -173,6 +173,14 @@ struct DeviceIndexPoolBwdImpl : public DeviceIndexPoolBwd<DOutDataType, IndexDat
...
@@ -173,6 +173,14 @@ struct DeviceIndexPoolBwdImpl : public DeviceIndexPoolBwd<DOutDataType, IndexDat
if
(
arg
.
p_workspace_
==
nullptr
)
if
(
arg
.
p_workspace_
==
nullptr
)
throw
std
::
runtime_error
(
"wrong! WorkSpace pointer has not been set"
);
throw
std
::
runtime_error
(
"wrong! WorkSpace pointer has not been set"
);
index_t
din_length_raw
=
arg
.
din_grid_desc_
.
GetTransforms
()[
I0
].
GetUpperLengths
()[
I0
];
hip_check_error
(
hipMemset
(
arg
.
p_workspace_
,
0
,
din_length_raw
*
sizeof
(
DInDataType_AutomicAddPreCast
)));
const
auto
put_kernel
=
kernel_put_element_1d
<
GridwisePutElementAtomicAdd
,
const
auto
put_kernel
=
kernel_put_element_1d
<
GridwisePutElementAtomicAdd
,
InOutGrid1dDesc
,
InOutGrid1dDesc
,
DOutDataType
,
DOutDataType
,
...
@@ -244,7 +252,6 @@ struct DeviceIndexPoolBwdImpl : public DeviceIndexPoolBwd<DOutDataType, IndexDat
...
@@ -244,7 +252,6 @@ struct DeviceIndexPoolBwdImpl : public DeviceIndexPoolBwd<DOutDataType, IndexDat
}
}
};
};
// User need to set the value of workspace to zero
size_t
GetWorkSpaceSize
(
const
BaseArgument
*
pArg
)
const
override
size_t
GetWorkSpaceSize
(
const
BaseArgument
*
pArg
)
const
override
{
{
const
Argument
*
pArg_
=
dynamic_cast
<
const
Argument
*>
(
pArg
);
const
Argument
*
pArg_
=
dynamic_cast
<
const
Argument
*>
(
pArg
);
...
...
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