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
beed1068
Commit
beed1068
authored
Sep 06, 2022
by
Po-Yen, Chen
Browse files
Use CRTP to generate overridden virtual method
parent
c7aa455f
Changes
1
Show whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
19 additions
and
10 deletions
+19
-10
include/ck/tensor_operation/gpu/device/device_permute.hpp
include/ck/tensor_operation/gpu/device/device_permute.hpp
+19
-10
No files found.
include/ck/tensor_operation/gpu/device/device_permute.hpp
View file @
beed1068
...
@@ -49,6 +49,22 @@ struct DevicePermuteBase : BaseOperator
...
@@ -49,6 +49,22 @@ struct DevicePermuteBase : BaseOperator
static
auto
MakeInvokerPointer
()
{
return
std
::
make_unique
<
typename
Derived
::
Invoker
>
();
};
static
auto
MakeInvokerPointer
()
{
return
std
::
make_unique
<
typename
Derived
::
Invoker
>
();
};
};
};
template
<
typename
Derived
,
typename
Argument
>
struct
InvokerBase
:
BaseInvoker
{
float
Run
(
const
BaseArgument
*
arg
,
const
StreamConfig
&
stream_config
=
StreamConfig
{})
override
final
{
const
auto
*
argument
=
dynamic_cast
<
const
Argument
*>
(
arg
);
if
(
!
argument
)
{
return
false
;
}
return
Derived
::
Run
(
*
argument
,
stream_config
);
}
};
}
// namespace detail
}
// namespace detail
template
<
typename
InDataType
,
template
<
typename
InDataType
,
...
@@ -195,9 +211,9 @@ struct DevicePermute : detail::DevicePermuteBase<DevicePermute<InDataType,
...
@@ -195,9 +211,9 @@ struct DevicePermute : detail::DevicePermuteBase<DevicePermute<InDataType,
ElementwiseOperation
elementwise_op_
;
ElementwiseOperation
elementwise_op_
;
};
};
struct
Invoker
:
public
BaseInvoker
struct
Invoker
:
detail
::
Invoker
Base
<
Invoker
,
Argument
>
{
{
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_elementwise_1d
<
GridwiseElementwise
,
const
auto
kernel
=
kernel_elementwise_1d
<
GridwiseElementwise
,
InGrid1dDescTuple
,
InGrid1dDescTuple
,
...
@@ -218,13 +234,6 @@ struct DevicePermute : detail::DevicePermuteBase<DevicePermute<InDataType,
...
@@ -218,13 +234,6 @@ struct DevicePermute : detail::DevicePermuteBase<DevicePermute<InDataType,
arg
.
elementwise_op_
);
arg
.
elementwise_op_
);
return
elapsed_time
;
return
elapsed_time
;
}
}
// polymorphic
float
Run
(
const
BaseArgument
*
p_arg
,
const
StreamConfig
&
stream_config
=
StreamConfig
{})
override
{
return
Run
(
*
dynamic_cast
<
const
Argument
*>
(
p_arg
),
stream_config
);
}
};
};
static
bool
IsSupportedArgument
(
const
Argument
&
arg
)
static
bool
IsSupportedArgument
(
const
Argument
&
arg
)
...
@@ -259,7 +268,7 @@ struct DevicePermute : detail::DevicePermuteBase<DevicePermute<InDataType,
...
@@ -259,7 +268,7 @@ struct DevicePermute : detail::DevicePermuteBase<DevicePermute<InDataType,
return
valid
;
return
valid
;
};
};
};
// namespace device
};
}
// namespace device
}
// namespace device
}
// namespace tensor_operation
}
// namespace tensor_operation
...
...
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