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
0840e016
Commit
0840e016
authored
May 29, 2023
by
Po-Yen, Chen
Browse files
Use type alias to hint usage
parent
813d4067
Changes
1
Show whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
5 additions
and
4 deletions
+5
-4
include/ck/utility/readfirstlane.hpp
include/ck/utility/readfirstlane.hpp
+5
-4
No files found.
include/ck/utility/readfirstlane.hpp
View file @
0840e016
...
@@ -50,8 +50,9 @@ template <
...
@@ -50,8 +50,9 @@ template <
typename
=
std
::
enable_if_t
<
std
::
is_class_v
<
Object
>
&&
std
::
is_trivially_copyable_v
<
Object
>>>
typename
=
std
::
enable_if_t
<
std
::
is_class_v
<
Object
>
&&
std
::
is_trivially_copyable_v
<
Object
>>>
__device__
auto
readfirstlane
(
const
Object
&
obj
)
__device__
auto
readfirstlane
(
const
Object
&
obj
)
{
{
constexpr
unsigned
SgprSize
=
4
;
using
Size
=
unsigned
;
constexpr
unsigned
ObjectSize
=
sizeof
(
Object
);
constexpr
Size
SgprSize
=
4
;
constexpr
Size
ObjectSize
=
sizeof
(
Object
);
using
Sgpr
=
detail
::
get_unsigned_int_t
<
SgprSize
>
;
using
Sgpr
=
detail
::
get_unsigned_int_t
<
SgprSize
>
;
...
@@ -59,8 +60,8 @@ __device__ auto readfirstlane(const Object& obj)
...
@@ -59,8 +60,8 @@ __device__ auto readfirstlane(const Object& obj)
auto
*
const
from_obj
=
reinterpret_cast
<
const
std
::
byte
*>
(
&
obj
);
auto
*
const
from_obj
=
reinterpret_cast
<
const
std
::
byte
*>
(
&
obj
);
constexpr
unsigned
RemainedSize
=
ObjectSize
%
SgprSize
;
constexpr
Size
RemainedSize
=
ObjectSize
%
SgprSize
;
constexpr
unsigned
CompleteSgprCopyBoundary
=
ObjectSize
-
RemainedSize
;
constexpr
Size
CompleteSgprCopyBoundary
=
ObjectSize
-
RemainedSize
;
static_for
<
0
,
CompleteSgprCopyBoundary
,
SgprSize
>
{}([
&
](
auto
offset
)
{
static_for
<
0
,
CompleteSgprCopyBoundary
,
SgprSize
>
{}([
&
](
auto
offset
)
{
*
reinterpret_cast
<
Sgpr
*>
(
to_obj
+
offset
)
=
*
reinterpret_cast
<
Sgpr
*>
(
to_obj
+
offset
)
=
readfirstlane
(
*
reinterpret_cast
<
const
Sgpr
*>
(
from_obj
+
offset
));
readfirstlane
(
*
reinterpret_cast
<
const
Sgpr
*>
(
from_obj
+
offset
));
...
...
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