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
d8214b28
Commit
d8214b28
authored
Jan 16, 2025
by
Rostyslav Geyyer
Browse files
Add tests
parent
3a64757f
Changes
2
Hide whitespace changes
Inline
Side-by-side
Showing
2 changed files
with
275 additions
and
1 deletion
+275
-1
test/data_type/test_bf6.cpp
test/data_type/test_bf6.cpp
+138
-1
test/data_type/test_fp6.cpp
test/data_type/test_fp6.cpp
+137
-0
No files found.
test/data_type/test_bf6.cpp
View file @
d8214b28
...
...
@@ -9,6 +9,8 @@
using
ck
::
bf6_convert_rne
;
using
ck
::
bf6_convert_sr
;
using
ck
::
bf6_t
;
using
ck
::
bf6x16_pk_t
;
using
ck
::
bf6x32_pk_t
;
using
ck
::
e8m0_bexp_t
;
using
ck
::
Number
;
using
ck
::
scaled_type_convert
;
...
...
@@ -40,7 +42,7 @@ TEST(BF6, ConvertFP32Nearest)
type_convert
<
float
>
(
bf6_convert_rne
(
std
::
numeric_limits
<
float
>::
infinity
())),
0.0
f
);
// convert float value less than bf6 subnorm to bf6 and back, check if equal to 0.0
float
less_than_subnorm
=
0.0
6
25
f
;
float
less_than_subnorm
=
0.0
31
25
f
;
ASSERT_NEAR
(
0.0
f
,
type_convert
<
float
>
(
bf6_convert_rne
(
less_than_subnorm
)),
0.0
f
);
// convert float NaN to bf6 and back, check if clipped to max_bf6
ASSERT_NEAR
(
max_bf6
,
...
...
@@ -216,3 +218,138 @@ TEST(BF6, ScaledConvertFP32Stochastic)
scaled_type_convert
<
float
>
(
e8m0_bexp_t
(
min_scale
),
bf6_convert_sr
(
neg_float
)),
abs_tol
);
}
TEST
(
FP6
,
TestSize
)
{
ASSERT_EQ
(
1
,
sizeof
(
bf6_t
));
ASSERT_EQ
(
12
,
sizeof
(
bf6x16_pk_t
));
ASSERT_EQ
(
24
,
sizeof
(
bf6x32_pk_t
));
ASSERT_EQ
(
16
,
sizeof
(
vector_type
<
bf6x16_pk_t
,
1
>
));
ASSERT_EQ
(
32
,
sizeof
(
vector_type
<
bf6x16_pk_t
,
2
>
));
ASSERT_EQ
(
32
,
sizeof
(
vector_type
<
bf6x32_pk_t
,
1
>
));
}
TEST
(
FP6
,
TestAlignment
)
{
ASSERT_EQ
(
1
,
alignof
(
bf6_t
));
ASSERT_EQ
(
4
,
alignof
(
bf6x16_pk_t
));
ASSERT_EQ
(
4
,
alignof
(
bf6x32_pk_t
));
ASSERT_EQ
(
16
,
alignof
(
vector_type
<
bf6x16_pk_t
,
1
>
));
ASSERT_EQ
(
32
,
alignof
(
vector_type
<
bf6x16_pk_t
,
2
>
));
ASSERT_EQ
(
32
,
alignof
(
vector_type
<
bf6x32_pk_t
,
1
>
));
}
// test vector of 1 bf6x16_pk_t, contains 16 bf6_t
TEST
(
FP6
,
TestAsType16x1
)
{
// test size
const
int
vector_size
=
1
;
const
int
packed_size
=
16
;
std
::
vector
<
bf6_t
>
test_vec
=
{
bf6_t
(
0b000000
),
bf6_t
(
0b100000
),
bf6_t
(
0b000001
),
bf6_t
(
0b100001
),
bf6_t
(
0b000010
),
bf6_t
(
0b100010
),
bf6_t
(
0b000011
),
bf6_t
(
0b100011
),
bf6_t
(
0b000100
),
bf6_t
(
0b100100
),
bf6_t
(
0b000101
),
bf6_t
(
0b100101
),
bf6_t
(
0b000110
),
bf6_t
(
0b100110
),
bf6_t
(
0b001011
),
bf6_t
(
0b101011
)};
// reference vector
vector_type
<
bf6x16_pk_t
,
vector_size
>
right_vec
;
// check default CTOR
ck
::
static_for
<
0
,
packed_size
,
1
>
{}([
&
](
auto
i
)
{
ASSERT_EQ
(
right_vec
.
template
AsType
<
bf6x16_pk_t
>()(
Number
<
0
>
{}).
template
unpack
<
i
>(),
0
);
});
// assign test values to the vector
ck
::
static_for
<
0
,
vector_size
,
1
>
{}([
&
](
auto
i
)
{
right_vec
.
template
AsType
<
bf6x16_pk_t
>()(
Number
<
i
>
{})
=
bf6x16_pk_t
{}.
pack
(
test_vec
.
data
());
});
// copy the vector
vector_type
<
bf6x16_pk_t
,
vector_size
>
left_vec
{
right_vec
};
// check if values were copied correctly
ck
::
static_for
<
0
,
packed_size
,
1
>
{}([
&
](
auto
i
)
{
ASSERT_EQ
(
left_vec
.
template
AsType
<
bf6x16_pk_t
>()(
Number
<
0
>
{}).
template
unpack
<
i
>(),
test_vec
.
at
(
i
));
});
}
// test vector of 2 bf6x16_pk_t, contains 32 bf6_t
TEST
(
FP6
,
TestAsType16x2
)
{
// test size
const
int
vector_size
=
2
;
const
int
packed_size
=
16
;
std
::
vector
<
bf6_t
>
test_vec
=
{
bf6_t
(
0b000000
),
bf6_t
(
0b100000
),
bf6_t
(
0b000001
),
bf6_t
(
0b100001
),
bf6_t
(
0b000010
),
bf6_t
(
0b100010
),
bf6_t
(
0b000011
),
bf6_t
(
0b100011
),
bf6_t
(
0b000100
),
bf6_t
(
0b100100
),
bf6_t
(
0b000101
),
bf6_t
(
0b100101
),
bf6_t
(
0b000110
),
bf6_t
(
0b100110
),
bf6_t
(
0b001011
),
bf6_t
(
0b101011
),
bf6_t
(
0b010000
),
bf6_t
(
0b110000
),
bf6_t
(
0b010001
),
bf6_t
(
0b110001
),
bf6_t
(
0b010010
),
bf6_t
(
0b110010
),
bf6_t
(
0b010011
),
bf6_t
(
0b110011
),
bf6_t
(
0b010100
),
bf6_t
(
0b110100
),
bf6_t
(
0b010101
),
bf6_t
(
0b110101
),
bf6_t
(
0b010110
),
bf6_t
(
0b110110
),
bf6_t
(
0b011011
),
bf6_t
(
0b111011
)};
// reference vector
vector_type
<
bf6x16_pk_t
,
vector_size
>
right_vec
;
// check default CTOR
ck
::
static_for
<
0
,
vector_size
,
1
>
{}([
&
](
auto
idx_vector
)
{
ck
::
static_for
<
0
,
packed_size
,
1
>
{}([
&
](
auto
idx_element
)
{
ASSERT_EQ
(
right_vec
.
template
AsType
<
bf6x16_pk_t
>()(
Number
<
idx_vector
>
{})
.
template
unpack
<
idx_element
>(),
0
);
});
});
// assign test values to the vector
ck
::
static_for
<
0
,
vector_size
,
1
>
{}([
&
](
auto
i
)
{
right_vec
.
template
AsType
<
bf6x16_pk_t
>()(
Number
<
i
>
{})
=
bf6x16_pk_t
{}.
pack
(
test_vec
.
data
()
+
i
*
packed_size
);
});
// copy the vector
vector_type
<
bf6x16_pk_t
,
vector_size
>
left_vec
{
right_vec
};
// check if values were copied correctly
ck
::
static_for
<
0
,
vector_size
,
1
>
{}([
&
](
auto
idx_vector
)
{
ck
::
static_for
<
0
,
packed_size
,
1
>
{}([
&
](
auto
idx_element
)
{
ASSERT_EQ
(
left_vec
.
template
AsType
<
bf6x16_pk_t
>()(
Number
<
idx_vector
>
{})
.
template
unpack
<
idx_element
>(),
test_vec
.
at
(
idx_vector
*
packed_size
+
idx_element
));
});
});
}
// test vector of 1 bf6x32_pk_t, contains 32 bf6_t
TEST
(
FP6
,
TestAsType32x1
)
{
// test size
const
int
vector_size
=
1
;
const
int
packed_size
=
32
;
std
::
vector
<
bf6_t
>
test_vec
=
{
bf6_t
(
0b000000
),
bf6_t
(
0b100000
),
bf6_t
(
0b000001
),
bf6_t
(
0b100001
),
bf6_t
(
0b000010
),
bf6_t
(
0b100010
),
bf6_t
(
0b000011
),
bf6_t
(
0b100011
),
bf6_t
(
0b000100
),
bf6_t
(
0b100100
),
bf6_t
(
0b000101
),
bf6_t
(
0b100101
),
bf6_t
(
0b000110
),
bf6_t
(
0b100110
),
bf6_t
(
0b001011
),
bf6_t
(
0b101011
),
bf6_t
(
0b010000
),
bf6_t
(
0b110000
),
bf6_t
(
0b010001
),
bf6_t
(
0b110001
),
bf6_t
(
0b010010
),
bf6_t
(
0b110010
),
bf6_t
(
0b010011
),
bf6_t
(
0b110011
),
bf6_t
(
0b010100
),
bf6_t
(
0b110100
),
bf6_t
(
0b010101
),
bf6_t
(
0b110101
),
bf6_t
(
0b010110
),
bf6_t
(
0b110110
),
bf6_t
(
0b011011
),
bf6_t
(
0b111011
)};
// reference vector
vector_type
<
bf6x32_pk_t
,
vector_size
>
right_vec
;
// check default CTOR
ck
::
static_for
<
0
,
packed_size
,
1
>
{}([
&
](
auto
i
)
{
ASSERT_EQ
(
right_vec
.
template
AsType
<
bf6x32_pk_t
>()(
Number
<
0
>
{}).
template
unpack
<
i
>(),
0
);
});
// assign test values to the vector
ck
::
static_for
<
0
,
vector_size
,
1
>
{}([
&
](
auto
i
)
{
right_vec
.
template
AsType
<
bf6x32_pk_t
>()(
Number
<
i
>
{})
=
bf6x32_pk_t
{}.
pack
(
test_vec
.
data
());
});
// copy the vector
vector_type
<
bf6x32_pk_t
,
vector_size
>
left_vec
{
right_vec
};
// check if values were copied correctly
ck
::
static_for
<
0
,
packed_size
,
1
>
{}([
&
](
auto
i
)
{
ASSERT_EQ
(
left_vec
.
template
AsType
<
bf6x32_pk_t
>()(
Number
<
0
>
{}).
template
unpack
<
i
>(),
test_vec
.
at
(
i
));
});
}
test/data_type/test_fp6.cpp
View file @
d8214b28
...
...
@@ -10,6 +10,8 @@ using ck::e8m0_bexp_t;
using
ck
::
f6_convert_rne
;
using
ck
::
f6_convert_sr
;
using
ck
::
f6_t
;
using
ck
::
f6x16_pk_t
;
using
ck
::
f6x32_pk_t
;
using
ck
::
Number
;
using
ck
::
scaled_type_convert
;
using
ck
::
type_convert
;
...
...
@@ -215,3 +217,138 @@ TEST(FP6, ScaledConvertFP32Stochastic)
scaled_type_convert
<
float
>
(
e8m0_bexp_t
(
min_scale
),
f6_convert_sr
(
neg_float
)),
abs_tol
);
}
TEST
(
FP6
,
TestSize
)
{
ASSERT_EQ
(
1
,
sizeof
(
f6_t
));
ASSERT_EQ
(
12
,
sizeof
(
f6x16_pk_t
));
ASSERT_EQ
(
24
,
sizeof
(
f6x32_pk_t
));
ASSERT_EQ
(
16
,
sizeof
(
vector_type
<
f6x16_pk_t
,
1
>
));
ASSERT_EQ
(
32
,
sizeof
(
vector_type
<
f6x16_pk_t
,
2
>
));
ASSERT_EQ
(
32
,
sizeof
(
vector_type
<
f6x32_pk_t
,
1
>
));
}
TEST
(
FP6
,
TestAlignment
)
{
ASSERT_EQ
(
1
,
alignof
(
f6_t
));
ASSERT_EQ
(
4
,
alignof
(
f6x16_pk_t
));
ASSERT_EQ
(
4
,
alignof
(
f6x32_pk_t
));
ASSERT_EQ
(
16
,
alignof
(
vector_type
<
f6x16_pk_t
,
1
>
));
ASSERT_EQ
(
32
,
alignof
(
vector_type
<
f6x16_pk_t
,
2
>
));
ASSERT_EQ
(
32
,
alignof
(
vector_type
<
f6x32_pk_t
,
1
>
));
}
// test vector of 1 f6x16_pk_t, contains 16 f6_t
TEST
(
FP6
,
TestAsType16x1
)
{
// test size
const
int
vector_size
=
1
;
const
int
packed_size
=
16
;
std
::
vector
<
f6_t
>
test_vec
=
{
f6_t
(
0b000000
),
f6_t
(
0b100000
),
f6_t
(
0b000001
),
f6_t
(
0b100001
),
f6_t
(
0b000010
),
f6_t
(
0b100010
),
f6_t
(
0b000011
),
f6_t
(
0b100011
),
f6_t
(
0b000100
),
f6_t
(
0b100100
),
f6_t
(
0b000101
),
f6_t
(
0b100101
),
f6_t
(
0b000110
),
f6_t
(
0b100110
),
f6_t
(
0b001011
),
f6_t
(
0b101011
)};
// reference vector
vector_type
<
f6x16_pk_t
,
vector_size
>
right_vec
;
// check default CTOR
ck
::
static_for
<
0
,
packed_size
,
1
>
{}([
&
](
auto
i
)
{
ASSERT_EQ
(
right_vec
.
template
AsType
<
f6x16_pk_t
>()(
Number
<
0
>
{}).
template
unpack
<
i
>(),
0
);
});
// assign test values to the vector
ck
::
static_for
<
0
,
vector_size
,
1
>
{}([
&
](
auto
i
)
{
right_vec
.
template
AsType
<
f6x16_pk_t
>()(
Number
<
i
>
{})
=
f6x16_pk_t
{}.
pack
(
test_vec
.
data
());
});
// copy the vector
vector_type
<
f6x16_pk_t
,
vector_size
>
left_vec
{
right_vec
};
// check if values were copied correctly
ck
::
static_for
<
0
,
packed_size
,
1
>
{}([
&
](
auto
i
)
{
ASSERT_EQ
(
left_vec
.
template
AsType
<
f6x16_pk_t
>()(
Number
<
0
>
{}).
template
unpack
<
i
>(),
test_vec
.
at
(
i
));
});
}
// test vector of 2 f6x16_pk_t, contains 32 f6_t
TEST
(
FP6
,
TestAsType16x2
)
{
// test size
const
int
vector_size
=
2
;
const
int
packed_size
=
16
;
std
::
vector
<
f6_t
>
test_vec
=
{
f6_t
(
0b000000
),
f6_t
(
0b100000
),
f6_t
(
0b000001
),
f6_t
(
0b100001
),
f6_t
(
0b000010
),
f6_t
(
0b100010
),
f6_t
(
0b000011
),
f6_t
(
0b100011
),
f6_t
(
0b000100
),
f6_t
(
0b100100
),
f6_t
(
0b000101
),
f6_t
(
0b100101
),
f6_t
(
0b000110
),
f6_t
(
0b100110
),
f6_t
(
0b001011
),
f6_t
(
0b101011
),
f6_t
(
0b010000
),
f6_t
(
0b110000
),
f6_t
(
0b010001
),
f6_t
(
0b110001
),
f6_t
(
0b010010
),
f6_t
(
0b110010
),
f6_t
(
0b010011
),
f6_t
(
0b110011
),
f6_t
(
0b010100
),
f6_t
(
0b110100
),
f6_t
(
0b010101
),
f6_t
(
0b110101
),
f6_t
(
0b010110
),
f6_t
(
0b110110
),
f6_t
(
0b011011
),
f6_t
(
0b111011
)};
// reference vector
vector_type
<
f6x16_pk_t
,
vector_size
>
right_vec
;
// check default CTOR
ck
::
static_for
<
0
,
vector_size
,
1
>
{}([
&
](
auto
idx_vector
)
{
ck
::
static_for
<
0
,
packed_size
,
1
>
{}([
&
](
auto
idx_element
)
{
ASSERT_EQ
(
right_vec
.
template
AsType
<
f6x16_pk_t
>()(
Number
<
idx_vector
>
{})
.
template
unpack
<
idx_element
>(),
0
);
});
});
// assign test values to the vector
ck
::
static_for
<
0
,
vector_size
,
1
>
{}([
&
](
auto
i
)
{
right_vec
.
template
AsType
<
f6x16_pk_t
>()(
Number
<
i
>
{})
=
f6x16_pk_t
{}.
pack
(
test_vec
.
data
()
+
i
*
packed_size
);
});
// copy the vector
vector_type
<
f6x16_pk_t
,
vector_size
>
left_vec
{
right_vec
};
// check if values were copied correctly
ck
::
static_for
<
0
,
vector_size
,
1
>
{}([
&
](
auto
idx_vector
)
{
ck
::
static_for
<
0
,
packed_size
,
1
>
{}([
&
](
auto
idx_element
)
{
ASSERT_EQ
(
left_vec
.
template
AsType
<
f6x16_pk_t
>()(
Number
<
idx_vector
>
{})
.
template
unpack
<
idx_element
>(),
test_vec
.
at
(
idx_vector
*
packed_size
+
idx_element
));
});
});
}
// test vector of 1 f6x32_pk_t, contains 32 f6_t
TEST
(
FP6
,
TestAsType32x1
)
{
// test size
const
int
vector_size
=
1
;
const
int
packed_size
=
32
;
std
::
vector
<
f6_t
>
test_vec
=
{
f6_t
(
0b000000
),
f6_t
(
0b100000
),
f6_t
(
0b000001
),
f6_t
(
0b100001
),
f6_t
(
0b000010
),
f6_t
(
0b100010
),
f6_t
(
0b000011
),
f6_t
(
0b100011
),
f6_t
(
0b000100
),
f6_t
(
0b100100
),
f6_t
(
0b000101
),
f6_t
(
0b100101
),
f6_t
(
0b000110
),
f6_t
(
0b100110
),
f6_t
(
0b001011
),
f6_t
(
0b101011
),
f6_t
(
0b010000
),
f6_t
(
0b110000
),
f6_t
(
0b010001
),
f6_t
(
0b110001
),
f6_t
(
0b010010
),
f6_t
(
0b110010
),
f6_t
(
0b010011
),
f6_t
(
0b110011
),
f6_t
(
0b010100
),
f6_t
(
0b110100
),
f6_t
(
0b010101
),
f6_t
(
0b110101
),
f6_t
(
0b010110
),
f6_t
(
0b110110
),
f6_t
(
0b011011
),
f6_t
(
0b111011
)};
// reference vector
vector_type
<
f6x32_pk_t
,
vector_size
>
right_vec
;
// check default CTOR
ck
::
static_for
<
0
,
packed_size
,
1
>
{}([
&
](
auto
i
)
{
ASSERT_EQ
(
right_vec
.
template
AsType
<
f6x32_pk_t
>()(
Number
<
0
>
{}).
template
unpack
<
i
>(),
0
);
});
// assign test values to the vector
ck
::
static_for
<
0
,
vector_size
,
1
>
{}([
&
](
auto
i
)
{
right_vec
.
template
AsType
<
f6x32_pk_t
>()(
Number
<
i
>
{})
=
f6x32_pk_t
{}.
pack
(
test_vec
.
data
());
});
// copy the vector
vector_type
<
f6x32_pk_t
,
vector_size
>
left_vec
{
right_vec
};
// check if values were copied correctly
ck
::
static_for
<
0
,
packed_size
,
1
>
{}([
&
](
auto
i
)
{
ASSERT_EQ
(
left_vec
.
template
AsType
<
f6x32_pk_t
>()(
Number
<
0
>
{}).
template
unpack
<
i
>(),
test_vec
.
at
(
i
));
});
}
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