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
b24c5f66
Commit
b24c5f66
authored
Sep 07, 2022
by
Po-Yen, Chen
Browse files
Rename 'GridwisePermute' to 'GridwiseCopy'
parent
18ba135f
Changes
2
Hide whitespace changes
Inline
Side-by-side
Showing
2 changed files
with
14 additions
and
14 deletions
+14
-14
include/ck/tensor_operation/gpu/device/device_permute.hpp
include/ck/tensor_operation/gpu/device/device_permute.hpp
+10
-10
include/ck/tensor_operation/gpu/grid/gridwise_copy.hpp
include/ck/tensor_operation/gpu/grid/gridwise_copy.hpp
+4
-4
No files found.
include/ck/tensor_operation/gpu/device/device_permute.hpp
View file @
b24c5f66
...
@@ -10,7 +10,7 @@
...
@@ -10,7 +10,7 @@
#include "ck/utility/math.hpp"
#include "ck/utility/math.hpp"
#include "ck/utility/sequence.hpp"
#include "ck/utility/sequence.hpp"
#include "ck/tensor_operation/gpu/device/device_base.hpp"
#include "ck/tensor_operation/gpu/device/device_base.hpp"
#include "ck/tensor_operation/gpu/grid/gridwise_
permute
.hpp"
#include "ck/tensor_operation/gpu/grid/gridwise_
copy
.hpp"
#include "ck/tensor_description/tensor_descriptor_helper.hpp"
#include "ck/tensor_description/tensor_descriptor_helper.hpp"
#include "ck/host_utility/kernel_launch.hpp"
#include "ck/host_utility/kernel_launch.hpp"
...
@@ -166,14 +166,14 @@ struct DevicePermute : detail::DevicePermuteBase<DevicePermute<InDataType,
...
@@ -166,14 +166,14 @@ struct DevicePermute : detail::DevicePermuteBase<DevicePermute<InDataType,
using
InGrid1dDesc
=
decltype
(
MakeDescriptor_N_H_W
({
1
,
1
},
{
1
,
1
},
1
,
1
));
using
InGrid1dDesc
=
decltype
(
MakeDescriptor_N_H_W
({
1
,
1
},
{
1
,
1
},
1
,
1
));
using
OutGrid1dDesc
=
decltype
(
MakeDescriptor_N_H_W
({
1
,
1
},
{
1
,
1
},
1
,
1
));
using
OutGrid1dDesc
=
decltype
(
MakeDescriptor_N_H_W
({
1
,
1
},
{
1
,
1
},
1
,
1
));
using
Gridwise
Permute
=
Gridwise
Permute
<
InGrid1dDesc
,
using
Gridwise
Copy
=
Gridwise
Copy
<
InGrid1dDesc
,
OutGrid1dDesc
,
OutGrid1dDesc
,
InDataTypePointer
,
InDataTypePointer
,
OutDataTypePointer
,
OutDataTypePointer
,
ElementwiseOperation
,
ElementwiseOperation
,
MPerThread
,
MPerThread
,
InScalarPerVector
,
InScalarPerVector
,
OutScalarPerVector
>
;
OutScalarPerVector
>
;
struct
Argument
:
public
BaseArgument
struct
Argument
:
public
BaseArgument
{
{
...
@@ -218,7 +218,7 @@ struct DevicePermute : detail::DevicePermuteBase<DevicePermute<InDataType,
...
@@ -218,7 +218,7 @@ struct DevicePermute : detail::DevicePermuteBase<DevicePermute<InDataType,
{
{
static
float
Run
(
const
Argument
&
arg
,
const
StreamConfig
&
stream_config
=
StreamConfig
{})
static
float
Run
(
const
Argument
&
arg
,
const
StreamConfig
&
stream_config
=
StreamConfig
{})
{
{
const
auto
kernel
=
kernel_
permute
<
Gridwise
Permute
,
const
auto
kernel
=
kernel_
nd_copy
<
Gridwise
Copy
,
InGrid1dDesc
,
InGrid1dDesc
,
OutGrid1dDesc
,
OutGrid1dDesc
,
InDataTypePointer
,
InDataTypePointer
,
...
...
include/ck/tensor_operation/gpu/grid/gridwise_
permute
.hpp
→
include/ck/tensor_operation/gpu/grid/gridwise_
copy
.hpp
View file @
b24c5f66
...
@@ -10,19 +10,19 @@
...
@@ -10,19 +10,19 @@
namespace
ck
{
namespace
ck
{
template
<
typename
Gridwise
Permute
Functor
,
template
<
typename
Gridwise
Copy
Functor
,
typename
InGrid1dDesc
,
typename
InGrid1dDesc
,
typename
OutGrid1dDesc
,
typename
OutGrid1dDesc
,
typename
InDataTypePointer
,
typename
InDataTypePointer
,
typename
OutDataTypePointer
,
typename
OutDataTypePointer
,
typename
ElementwiseOperation
>
typename
ElementwiseOperation
>
__global__
void
kernel_
permute
(
const
InGrid1dDesc
in_grid_1d_desc
,
__global__
void
kernel_
nd_copy
(
const
InGrid1dDesc
in_grid_1d_desc
,
const
OutGrid1dDesc
out_grid_1d_desc
,
const
OutGrid1dDesc
out_grid_1d_desc
,
const
InDataTypePointer
p_in_global
,
const
InDataTypePointer
p_in_global
,
const
OutDataTypePointer
p_out_global
,
const
OutDataTypePointer
p_out_global
,
const
ElementwiseOperation
elementwise_op
)
const
ElementwiseOperation
elementwise_op
)
{
{
Gridwise
Permute
Functor
::
Run
(
Gridwise
Copy
Functor
::
Run
(
in_grid_1d_desc
,
out_grid_1d_desc
,
p_in_global
,
p_out_global
,
elementwise_op
);
in_grid_1d_desc
,
out_grid_1d_desc
,
p_in_global
,
p_out_global
,
elementwise_op
);
}
}
...
@@ -34,7 +34,7 @@ template <typename InGrid1dDesc,
...
@@ -34,7 +34,7 @@ template <typename InGrid1dDesc,
index_t
MPerThread
,
index_t
MPerThread
,
index_t
InScalarPerVector
,
index_t
InScalarPerVector
,
index_t
OutScalarPerVector
>
index_t
OutScalarPerVector
>
struct
Gridwise
Permute
struct
Gridwise
Copy
{
{
static
constexpr
auto
I0
=
Number
<
0
>
{};
static
constexpr
auto
I0
=
Number
<
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