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_ROCM
Commits
91fa13b0
Commit
91fa13b0
authored
Jan 31, 2025
by
Rostyslav Geyyer
Browse files
Use pointers instead of array indices
parent
9f58449c
Changes
1
Show whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
6 additions
and
66 deletions
+6
-66
include/ck/utility/type_convert.hpp
include/ck/utility/type_convert.hpp
+6
-66
No files found.
include/ck/utility/type_convert.hpp
View file @
91fa13b0
...
@@ -1431,39 +1431,9 @@ inline __host__ __device__ f6_t f6_convert_rne(float x, float scale = 1.0f)
...
@@ -1431,39 +1431,9 @@ inline __host__ __device__ f6_t f6_convert_rne(float x, float scale = 1.0f)
inline
__host__
__device__
f6x32_t
f6_convert_rne
(
float32_t
x
,
float
scale
=
1.0
f
)
inline
__host__
__device__
f6x32_t
f6_convert_rne
(
float32_t
x
,
float
scale
=
1.0
f
)
{
{
#if defined(__gfx950__)
#if defined(__gfx950__)
float16_t
in1
{
x
[
0
],
float16_t
*
in1
=
reinterpret_cast
<
float16_t
*>
(
&
x
);
x
[
1
],
float16_t
*
in2
=
reinterpret_cast
<
float16_t
*>
(
&
x
+
16
);
x
[
2
],
return
__builtin_amdgcn_cvt_scalef32_2xpk16_bf6_f32
(
*
in1
,
*
in2
,
scale
);
x
[
3
],
x
[
4
],
x
[
5
],
x
[
6
],
x
[
7
],
x
[
8
],
x
[
9
],
x
[
10
],
x
[
11
],
x
[
12
],
x
[
13
],
x
[
14
],
x
[
15
]};
float16_t
in2
=
{
x
[
16
],
x
[
17
],
x
[
18
],
x
[
19
],
x
[
20
],
x
[
21
],
x
[
22
],
x
[
23
],
x
[
24
],
x
[
25
],
x
[
26
],
x
[
27
],
x
[
28
],
x
[
29
],
x
[
30
],
x
[
31
]};
return
__builtin_amdgcn_cvt_scalef32_2xpk16_fp6_f32
(
in1
,
in2
,
scale
);
#else
#else
union
union
{
{
...
@@ -1719,39 +1689,9 @@ inline __host__ __device__ bf6_t bf6_convert_rne(float x, float scale = 1.0f)
...
@@ -1719,39 +1689,9 @@ inline __host__ __device__ bf6_t bf6_convert_rne(float x, float scale = 1.0f)
inline
__host__
__device__
bf6x32_t
bf6_convert_rne
(
float32_t
x
,
float
scale
=
1.0
f
)
inline
__host__
__device__
bf6x32_t
bf6_convert_rne
(
float32_t
x
,
float
scale
=
1.0
f
)
{
{
#if defined(__gfx950__)
#if defined(__gfx950__)
float16_t
in1
{
x
[
0
],
float16_t
*
in1
=
reinterpret_cast
<
float16_t
*>
(
&
x
);
x
[
1
],
float16_t
*
in2
=
reinterpret_cast
<
float16_t
*>
(
&
x
+
16
);
x
[
2
],
return
__builtin_amdgcn_cvt_scalef32_2xpk16_bf6_f32
(
*
in1
,
*
in2
,
scale
);
x
[
3
],
x
[
4
],
x
[
5
],
x
[
6
],
x
[
7
],
x
[
8
],
x
[
9
],
x
[
10
],
x
[
11
],
x
[
12
],
x
[
13
],
x
[
14
],
x
[
15
]};
float16_t
in2
=
{
x
[
16
],
x
[
17
],
x
[
18
],
x
[
19
],
x
[
20
],
x
[
21
],
x
[
22
],
x
[
23
],
x
[
24
],
x
[
25
],
x
[
26
],
x
[
27
],
x
[
28
],
x
[
29
],
x
[
30
],
x
[
31
]};
return
__builtin_amdgcn_cvt_scalef32_2xpk16_bf6_f32
(
in1
,
in2
,
scale
);
#else
#else
union
union
{
{
...
...
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