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
1b146563
Commit
1b146563
authored
Jun 01, 2023
by
Po-Yen, Chen
Browse files
Add 3-bytes carrier type
parent
902e7c62
Changes
1
Hide whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
40 additions
and
1 deletion
+40
-1
include/ck/utility/amd_wave_read_first_lane.hpp
include/ck/utility/amd_wave_read_first_lane.hpp
+40
-1
No files found.
include/ck/utility/amd_wave_read_first_lane.hpp
View file @
1b146563
...
...
@@ -7,6 +7,8 @@
#include "ck/utility/functional2.hpp"
#include "ck/utility/math.hpp"
#include <algorithm>
#include <array>
#include <cstddef>
#include <cstdint>
#include <type_traits>
...
...
@@ -14,7 +16,7 @@
namespace
ck
{
namespace
detail
{
template
<
unsigned
Size
>
template
<
unsigned
Size
InBytes
>
struct
get_carrier
;
template
<
>
...
...
@@ -29,6 +31,43 @@ struct get_carrier<2>
using
type
=
uint16_t
;
};
template
<
>
struct
get_carrier
<
3
>
{
using
type
=
class
carrier
{
using
value_type
=
uint32_t
;
std
::
array
<
std
::
byte
,
3
>
bytes
;
static_assert
(
sizeof
(
bytes
)
<=
sizeof
(
value_type
));
public:
inline
carrier
(
value_type
value
)
noexcept
{
auto
const
from
=
reinterpret_cast
<
const
std
::
byte
*>
(
&
value
);
std
::
copy_n
(
from
,
bytes
.
size
(),
bytes
.
data
());
}
// method to trigger template substitution failure
inline
carrier
&
operator
=
(
const
carrier
&
other
)
noexcept
{
std
::
copy_n
(
other
.
bytes
.
data
(),
bytes
.
size
(),
bytes
.
data
());
return
*
this
;
}
inline
operator
value_type
()
const
noexcept
{
value_type
value
{};
auto
const
to
=
reinterpret_cast
<
std
::
byte
*>
(
&
value
);
std
::
copy_n
(
bytes
.
data
(),
bytes
.
size
(),
to
);
return
value
;
}
};
};
static_assert
(
sizeof
(
get_carrier
<
3
>::
type
)
==
3
);
template
<
>
struct
get_carrier
<
4
>
{
...
...
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