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
ea23062c
Commit
ea23062c
authored
Sep 19, 2022
by
Po-Yen, Chen
Browse files
Remove template BaseInvokerCRTP<>
parent
1ca0b97c
Changes
2
Show whitespace changes
Inline
Side-by-side
Showing
2 changed files
with
13 additions
and
19 deletions
+13
-19
include/ck/tensor_operation/gpu/device/device_base.hpp
include/ck/tensor_operation/gpu/device/device_base.hpp
+0
-18
include/ck/tensor_operation/gpu/device/impl/device_permute_impl.hpp
.../tensor_operation/gpu/device/impl/device_permute_impl.hpp
+13
-1
No files found.
include/ck/tensor_operation/gpu/device/device_base.hpp
View file @
ea23062c
...
@@ -57,24 +57,6 @@ struct BaseOperator
...
@@ -57,24 +57,6 @@ struct BaseOperator
virtual
~
BaseOperator
()
{}
virtual
~
BaseOperator
()
{}
};
};
template
<
typename
DerivedInvoker
,
typename
Argument
>
struct
BaseInvokerCRTP
:
BaseInvoker
{
static_assert
(
std
::
is_class_v
<
Argument
>
&&
std
::
is_base_of_v
<
BaseArgument
,
Argument
>
);
float
Run
(
const
BaseArgument
*
arg
,
const
StreamConfig
&
stream_config
=
StreamConfig
{})
override
final
{
const
auto
*
const
argument
=
dynamic_cast
<
const
Argument
*>
(
arg
);
if
(
!
argument
)
{
return
NAN
;
}
return
DerivedInvoker
::
Run
(
*
argument
,
stream_config
);
}
};
}
// namespace device
}
// namespace device
}
// namespace tensor_operation
}
// namespace tensor_operation
}
// namespace ck
}
// namespace ck
include/ck/tensor_operation/gpu/device/impl/device_permute_impl.hpp
View file @
ea23062c
...
@@ -146,7 +146,7 @@ struct DevicePermuteImpl : DevicePermute<NumDim, InDataType, OutDataType, Elemen
...
@@ -146,7 +146,7 @@ struct DevicePermuteImpl : DevicePermute<NumDim, InDataType, OutDataType, Elemen
Block2TileMap
block_2_tile_map_
;
Block2TileMap
block_2_tile_map_
;
};
};
struct
Invoker
:
BaseInvoker
CRTP
<
Invoker
,
Argument
>
struct
Invoker
:
BaseInvoker
{
{
static
float
Run
(
const
Argument
&
arg
,
const
StreamConfig
&
stream_config
=
StreamConfig
{})
static
float
Run
(
const
Argument
&
arg
,
const
StreamConfig
&
stream_config
=
StreamConfig
{})
{
{
...
@@ -173,6 +173,18 @@ struct DevicePermuteImpl : DevicePermute<NumDim, InDataType, OutDataType, Elemen
...
@@ -173,6 +173,18 @@ struct DevicePermuteImpl : DevicePermute<NumDim, InDataType, OutDataType, Elemen
arg
.
block_2_tile_map_
);
arg
.
block_2_tile_map_
);
return
elapsed_time
;
return
elapsed_time
;
}
}
float
Run
(
const
BaseArgument
*
arg
,
const
StreamConfig
&
stream_config
=
StreamConfig
{})
override
final
{
const
auto
*
const
argument
=
dynamic_cast
<
const
Argument
*>
(
arg
);
if
(
!
argument
)
{
return
NAN
;
}
return
Run
(
*
argument
,
stream_config
);
}
};
};
static
bool
IsSupportedArgument
(
const
Argument
&
arg
)
static
bool
IsSupportedArgument
(
const
Argument
&
arg
)
...
...
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