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
902e7c62
Commit
902e7c62
authored
Jun 01, 2023
by
Po-Yen, Chen
Browse files
Rename type trait get_unsigned_int<> to get_carrier<>
parent
290791c8
Changes
1
Hide whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
7 additions
and
7 deletions
+7
-7
include/ck/utility/amd_wave_read_first_lane.hpp
include/ck/utility/amd_wave_read_first_lane.hpp
+7
-7
No files found.
include/ck/utility/amd_wave_read_first_lane.hpp
View file @
902e7c62
...
@@ -15,28 +15,28 @@ namespace ck {
...
@@ -15,28 +15,28 @@ namespace ck {
namespace
detail
{
namespace
detail
{
template
<
unsigned
Size
>
template
<
unsigned
Size
>
struct
get_
unsigned_int
;
struct
get_
carrier
;
template
<
>
template
<
>
struct
get_
unsigned_int
<
1
>
struct
get_
carrier
<
1
>
{
{
using
type
=
uint8_t
;
using
type
=
uint8_t
;
};
};
template
<
>
template
<
>
struct
get_
unsigned_int
<
2
>
struct
get_
carrier
<
2
>
{
{
using
type
=
uint16_t
;
using
type
=
uint16_t
;
};
};
template
<
>
template
<
>
struct
get_
unsigned_int
<
4
>
struct
get_
carrier
<
4
>
{
{
using
type
=
uint32_t
;
using
type
=
uint32_t
;
};
};
template
<
unsigned
Size
>
template
<
unsigned
Size
>
using
get_
unsigned_int
_t
=
typename
get_
unsigned_int
<
Size
>::
type
;
using
get_
carrier
_t
=
typename
get_
carrier
<
Size
>::
type
;
}
// namespace detail
}
// namespace detail
...
@@ -61,7 +61,7 @@ __device__ auto amd_wave_read_first_lane(const Object& obj)
...
@@ -61,7 +61,7 @@ __device__ auto amd_wave_read_first_lane(const Object& obj)
constexpr
Size
CompleteSgprCopyBoundary
=
ObjectSize
-
RemainedSize
;
constexpr
Size
CompleteSgprCopyBoundary
=
ObjectSize
-
RemainedSize
;
for
(
Size
offset
=
0
;
offset
<
CompleteSgprCopyBoundary
;
offset
+=
SgprSize
)
for
(
Size
offset
=
0
;
offset
<
CompleteSgprCopyBoundary
;
offset
+=
SgprSize
)
{
{
using
Sgpr
=
detail
::
get_
unsigned_int
_t
<
SgprSize
>
;
using
Sgpr
=
detail
::
get_
carrier
_t
<
SgprSize
>
;
*
reinterpret_cast
<
Sgpr
*>
(
to_obj
+
offset
)
=
*
reinterpret_cast
<
Sgpr
*>
(
to_obj
+
offset
)
=
amd_wave_read_first_lane
(
*
reinterpret_cast
<
const
Sgpr
*>
(
from_obj
+
offset
));
amd_wave_read_first_lane
(
*
reinterpret_cast
<
const
Sgpr
*>
(
from_obj
+
offset
));
...
@@ -69,7 +69,7 @@ __device__ auto amd_wave_read_first_lane(const Object& obj)
...
@@ -69,7 +69,7 @@ __device__ auto amd_wave_read_first_lane(const Object& obj)
if
constexpr
(
0
<
RemainedSize
)
if
constexpr
(
0
<
RemainedSize
)
{
{
using
Carrier
=
detail
::
get_
unsigned_int
_t
<
RemainedSize
>
;
using
Carrier
=
detail
::
get_
carrier
_t
<
RemainedSize
>
;
*
reinterpret_cast
<
Carrier
*>
(
to_obj
+
CompleteSgprCopyBoundary
)
=
amd_wave_read_first_lane
(
*
reinterpret_cast
<
Carrier
*>
(
to_obj
+
CompleteSgprCopyBoundary
)
=
amd_wave_read_first_lane
(
*
reinterpret_cast
<
const
Carrier
*>
(
from_obj
+
CompleteSgprCopyBoundary
));
*
reinterpret_cast
<
const
Carrier
*>
(
from_obj
+
CompleteSgprCopyBoundary
));
...
...
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