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
cfdce3eb
Unverified
Commit
cfdce3eb
authored
Sep 08, 2023
by
Po Yen Chen
Committed by
GitHub
Sep 07, 2023
Browse files
Replace variable by integral constant (#8)
parent
98109c8b
Changes
1
Hide whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
12 additions
and
8 deletions
+12
-8
include/ck/tile_program/block_tile_pipeline/block_gemm_pipeline_agmem_bgmem_creg_v1_default_policy.hpp
...lock_gemm_pipeline_agmem_bgmem_creg_v1_default_policy.hpp
+12
-8
No files found.
include/ck/tile_program/block_tile_pipeline/block_gemm_pipeline_agmem_bgmem_creg_v1_default_policy.hpp
View file @
cfdce3eb
...
...
@@ -115,20 +115,22 @@ struct BlockGemmPipelineAGmemBGmemCRegV1DefaultPolicy
constexpr
index_t
kKPerBlock
=
Problem
::
BlockGemmShape
::
kK
;
constexpr
auto
a_lds_block_desc_d1_d2_d3
=
make_naive_tensor_descriptor_packed
(
make_tuple
(
kMPerBlock
/
2
,
2
,
kKPerBlock
),
Number
<
kKPerBlock
>
{});
make_tuple
(
Number
<
kMPerBlock
/
2
>
{},
Number
<
2
>
{},
Number
<
kKPerBlock
>
{}),
Number
<
kKPerBlock
>
{});
constexpr
index_t
kK1
=
16
/
sizeof
(
ADataType
);
constexpr
auto
a_lds_block_desc_d4_d5_d6
=
transform_tensor_descriptor
(
a_lds_block_desc_d1_d2_d3
,
make_tuple
(
make_xor_transform
(
make_tuple
(
kMPerBlock
/
2
,
kKPerBlock
),
kK1
),
make_pass_through_transform
(
2
)),
make_tuple
(
make_xor_transform
(
make_tuple
(
Number
<
kMPerBlock
/
2
>
{},
Number
<
kKPerBlock
>
{}),
kK1
),
make_pass_through_transform
(
2
)),
make_tuple
(
Sequence
<
0
,
2
>
{},
Sequence
<
1
>
{}),
make_tuple
(
Sequence
<
0
,
2
>
{},
Sequence
<
1
>
{}));
constexpr
auto
a_lds_block_desc_m_k
=
transform_tensor_descriptor
(
a_lds_block_desc_d4_d5_d6
,
make_tuple
(
make_merge_transform
(
make_tuple
(
kMPerBlock
/
2
,
2
)),
make_tuple
(
make_merge_transform
(
make_tuple
(
Number
<
kMPerBlock
/
2
>
{},
Number
<
2
>
{}
)),
make_pass_through_transform
(
kKPerBlock
)),
make_tuple
(
Sequence
<
0
,
1
>
{},
Sequence
<
2
>
{}),
make_tuple
(
Sequence
<
0
>
{},
Sequence
<
1
>
{}));
...
...
@@ -148,20 +150,22 @@ struct BlockGemmPipelineAGmemBGmemCRegV1DefaultPolicy
constexpr
index_t
kKPerBlock
=
Problem
::
BlockGemmShape
::
kK
;
constexpr
auto
b_lds_block_desc_d1_d2_d3
=
make_naive_tensor_descriptor_packed
(
make_tuple
(
kNPerBlock
/
2
,
2
,
kKPerBlock
),
Number
<
kKPerBlock
>
{});
make_tuple
(
Number
<
kNPerBlock
/
2
>
{},
Number
<
2
>
{},
Number
<
kKPerBlock
>
{}),
Number
<
kKPerBlock
>
{});
constexpr
index_t
kK1
=
16
/
sizeof
(
BDataType
);
constexpr
auto
b_lds_block_desc_d4_d5_d6
=
transform_tensor_descriptor
(
b_lds_block_desc_d1_d2_d3
,
make_tuple
(
make_xor_transform
(
make_tuple
(
kNPerBlock
/
2
,
kKPerBlock
),
kK1
),
make_pass_through_transform
(
2
)),
make_tuple
(
make_xor_transform
(
make_tuple
(
Number
<
kNPerBlock
/
2
>
{},
Number
<
kKPerBlock
>
{}),
kK1
),
make_pass_through_transform
(
2
)),
make_tuple
(
Sequence
<
0
,
2
>
{},
Sequence
<
1
>
{}),
make_tuple
(
Sequence
<
0
,
2
>
{},
Sequence
<
1
>
{}));
constexpr
auto
b_lds_block_desc_n_k
=
transform_tensor_descriptor
(
b_lds_block_desc_d4_d5_d6
,
make_tuple
(
make_merge_transform
(
make_tuple
(
kNPerBlock
/
2
,
2
)),
make_tuple
(
make_merge_transform
(
make_tuple
(
Number
<
kNPerBlock
/
2
>
{},
Number
<
2
>
{}
)),
make_pass_through_transform
(
kKPerBlock
)),
make_tuple
(
Sequence
<
0
,
1
>
{},
Sequence
<
2
>
{}),
make_tuple
(
Sequence
<
0
>
{},
Sequence
<
1
>
{}));
...
...
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