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
d1f7a3cf
Commit
d1f7a3cf
authored
Feb 26, 2024
by
Adam Osewski
Browse files
Fix B2C Linear KSplit test.
parent
d8dc850e
Changes
1
Hide whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
14 additions
and
10 deletions
+14
-10
test/block_to_ctile_map/test_block_to_ctile_map.cpp
test/block_to_ctile_map/test_block_to_ctile_map.cpp
+14
-10
No files found.
test/block_to_ctile_map/test_block_to_ctile_map.cpp
View file @
d1f7a3cf
// SPDX-License-Identifier: MIT
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-202
3
, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2018-202
4
, Advanced Micro Devices, Inc. All rights reserved.
#include <iostream>
#include <iostream>
#include <vector>
#include <vector>
...
@@ -364,12 +364,11 @@ TEST(BlockToCTileMap, BlockToCTileMap_LinearKSplit_BottomIndex)
...
@@ -364,12 +364,11 @@ TEST(BlockToCTileMap, BlockToCTileMap_LinearKSplit_BottomIndex)
TEST
(
BlockToCTileMap
,
BlockToCTileMap_LinearKSplit_NextKTile
)
TEST
(
BlockToCTileMap
,
BlockToCTileMap_LinearKSplit_NextKTile
)
{
{
const
index_t
M
=
768
;
const
index_t
M
=
768
;
const
index_t
N
=
384
;
const
index_t
N
=
384
;
const
index_t
MPerBlock
=
128
;
const
index_t
MPerBlock
=
128
;
const
index_t
NPerBlock
=
64
;
const
index_t
NPerBlock
=
64
;
const
index_t
KSplit
=
3
;
const
index_t
KSplit
=
3
;
const
index_t
tiles_per_block
=
1
;
auto
c_grid_desc_m_n
=
make_naive_tensor_descriptor_packed
(
make_tuple
(
M
,
N
));
auto
c_grid_desc_m_n
=
make_naive_tensor_descriptor_packed
(
make_tuple
(
M
,
N
));
...
@@ -378,7 +377,7 @@ TEST(BlockToCTileMap, BlockToCTileMap_LinearKSplit_NextKTile)
...
@@ -378,7 +377,7 @@ TEST(BlockToCTileMap, BlockToCTileMap_LinearKSplit_NextKTile)
auto
m0n0k0_idx
=
tile_map
.
CalculateBottomIndex
(
3
);
auto
m0n0k0_idx
=
tile_map
.
CalculateBottomIndex
(
3
);
EXPECT_EQ
((
std
::
vector
<
int
>
{
m0n0k0_idx
[
I0
],
m0n0k0_idx
[
I1
],
m0n0k0_idx
[
I2
]}),
EXPECT_EQ
((
std
::
vector
<
int
>
{
m0n0k0_idx
[
I0
],
m0n0k0_idx
[
I1
],
m0n0k0_idx
[
I2
]}),
(
std
::
vector
<
int
>
{
0
,
1
,
0
}));
(
std
::
vector
<
int
>
{
0
,
1
,
0
}));
EXPECT_TRUE
(
tile_map
.
IsFirstKSplitBlock
(
tiles_per_block
));
EXPECT_TRUE
(
tile_map
.
IsFirstKSplitBlock
());
for
(
index_t
i
=
0
;
i
<
KSplit
-
1
;
i
++
)
for
(
index_t
i
=
0
;
i
<
KSplit
-
1
;
i
++
)
{
{
...
@@ -386,11 +385,16 @@ TEST(BlockToCTileMap, BlockToCTileMap_LinearKSplit_NextKTile)
...
@@ -386,11 +385,16 @@ TEST(BlockToCTileMap, BlockToCTileMap_LinearKSplit_NextKTile)
m0n0k0_idx
=
tile_map
.
GetBottomIndex
();
m0n0k0_idx
=
tile_map
.
GetBottomIndex
();
EXPECT_EQ
((
std
::
vector
<
int
>
{
m0n0k0_idx
[
I0
],
m0n0k0_idx
[
I1
],
m0n0k0_idx
[
I2
]}),
EXPECT_EQ
((
std
::
vector
<
int
>
{
m0n0k0_idx
[
I0
],
m0n0k0_idx
[
I1
],
m0n0k0_idx
[
I2
]}),
(
std
::
vector
<
int
>
{
0
,
1
,
i
+
1
}));
(
std
::
vector
<
int
>
{
0
,
1
,
i
+
1
}));
EXPECT_
FALS
E
(
tile_map
.
IsFirstKSplitBlock
(
tiles_per_block
));
EXPECT_
TRU
E
(
tile_map
.
IsFirstKSplitBlock
());
}
}
EXPECT_FALSE
(
tile_map
.
GetNextKTileIdx
());
EXPECT_FALSE
(
tile_map
.
GetNextKTileIdx
());
m0n0k0_idx
=
tile_map
.
GetBottomIndex
();
m0n0k0_idx
=
tile_map
.
GetBottomIndex
();
EXPECT_EQ
((
std
::
vector
<
int
>
{
m0n0k0_idx
[
I0
],
m0n0k0_idx
[
I1
],
m0n0k0_idx
[
I2
]}),
EXPECT_EQ
((
std
::
vector
<
int
>
{
m0n0k0_idx
[
I0
],
m0n0k0_idx
[
I1
],
m0n0k0_idx
[
I2
]}),
(
std
::
vector
<
int
>
{
0
,
1
,
2
}));
(
std
::
vector
<
int
>
{
0
,
1
,
2
}));
EXPECT_FALSE
(
tile_map
.
IsFirstKSplitBlock
(
tiles_per_block
));
EXPECT_TRUE
(
tile_map
.
IsFirstKSplitBlock
());
m0n0k0_idx
=
tile_map
.
CalculateBottomIndex
(
7
);
EXPECT_EQ
((
std
::
vector
<
int
>
{
m0n0k0_idx
[
I0
],
m0n0k0_idx
[
I1
],
m0n0k0_idx
[
I2
]}),
(
std
::
vector
<
int
>
{
0
,
2
,
1
}));
EXPECT_FALSE
(
tile_map
.
IsFirstKSplitBlock
());
}
}
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