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
b0376100
Commit
b0376100
authored
Oct 15, 2024
by
Rostyslav Geyyer
Browse files
Update default stride value to -1
parent
d02a92cc
Changes
4
Show whitespace changes
Inline
Side-by-side
Showing
4 changed files
with
26 additions
and
26 deletions
+26
-26
example/01_gemm/common.hpp
example/01_gemm/common.hpp
+12
-12
example/01_gemm/run_gemm_example.inc
example/01_gemm/run_gemm_example.inc
+6
-6
example/01_gemm/run_gemm_example_streamk_v2.inc
example/01_gemm/run_gemm_example_streamk_v2.inc
+2
-2
example/01_gemm/run_gemm_example_v2.inc
example/01_gemm/run_gemm_example_v2.inc
+6
-6
No files found.
example/01_gemm/common.hpp
View file @
b0376100
...
...
@@ -29,9 +29,9 @@ struct ProblemSize final
ck
::
index_t
N
=
4096
;
ck
::
index_t
K
=
4096
;
ck
::
index_t
StrideA
=
0
;
ck
::
index_t
StrideB
=
0
;
ck
::
index_t
StrideC
=
0
;
ck
::
index_t
StrideA
=
-
1
;
ck
::
index_t
StrideB
=
-
1
;
ck
::
index_t
StrideC
=
-
1
;
};
struct
ProblemSizeStreamK
final
...
...
@@ -40,9 +40,9 @@ struct ProblemSizeStreamK final
ck
::
index_t
N
=
4096
;
ck
::
index_t
K
=
4096
;
ck
::
index_t
StrideA
=
0
;
ck
::
index_t
StrideB
=
0
;
ck
::
index_t
StrideC
=
0
;
ck
::
index_t
StrideA
=
-
1
;
ck
::
index_t
StrideB
=
-
1
;
ck
::
index_t
StrideC
=
-
1
;
ck
::
index_t
NumSKBlocks
=
-
1
;
};
...
...
@@ -52,9 +52,9 @@ struct ProblemSizeStreamK_universal final
ck
::
index_t
N
=
4096
;
ck
::
index_t
K
=
4096
;
ck
::
index_t
StrideA
=
0
;
ck
::
index_t
StrideB
=
0
;
ck
::
index_t
StrideC
=
0
;
ck
::
index_t
StrideA
=
-
1
;
ck
::
index_t
StrideB
=
-
1
;
ck
::
index_t
StrideC
=
-
1
;
ck
::
index_t
Grid_size
=
-
1
;
// defaults to max occupancy
ck
::
index_t
Streamk_sel
=
1
;
// defaults to 1-tile SK
...
...
@@ -66,9 +66,9 @@ struct ProblemSizeSplitK final
ck
::
index_t
N
=
4096
;
ck
::
index_t
K
=
4096
;
ck
::
index_t
StrideA
=
0
;
ck
::
index_t
StrideB
=
0
;
ck
::
index_t
StrideC
=
0
;
ck
::
index_t
StrideA
=
-
1
;
ck
::
index_t
StrideB
=
-
1
;
ck
::
index_t
StrideC
=
-
1
;
ck
::
index_t
KBatch
=
1
;
};
...
...
example/01_gemm/run_gemm_example.inc
View file @
b0376100
...
...
@@ -116,21 +116,21 @@ bool run_gemm(const ProblemType& problem_size, const ExecutionConfig& config)
};
auto
f_get_default_stride
=
[](
std
::
size_t
row
,
std
::
size_t
col
,
std
::
size
_t
stride
,
auto
layout
)
{
if
(
stride
==
0
)
[](
std
::
size_t
row
,
std
::
size_t
col
,
ck
::
index
_t
stride
,
auto
layout
)
{
if
(
stride
==
-
1
)
{
// give a chance if stride is
zero
, return a default packed stride
// give a chance if stride is
-1
, return a default packed stride
if
constexpr
(
std
::
is_same_v
<
decltype
(
layout
),
ck
::
tensor_layout
::
gemm
::
RowMajor
>
)
{
return
col
;
return
static_cast
<
std
::
size_t
>
(
col
)
;
}
else
{
return
row
;
return
static_cast
<
std
::
size_t
>
(
row
)
;
}
}
else
return
stride
;
return
static_cast
<
std
::
size_t
>
(
stride
)
;
};
StrideA
=
f_get_default_stride
(
M
,
K
,
StrideA
,
ALayout
{});
...
...
example/01_gemm/run_gemm_example_streamk_v2.inc
View file @
b0376100
...
...
@@ -117,9 +117,9 @@ bool run_gemm(const ProblemType& problem_size, const ExecutionConfig& config)
auto
f_get_default_stride
=
[](
std
::
size_t
row
,
std
::
size_t
col
,
ck
::
index_t
stride
,
auto
layout
)
{
if
(
stride
==
0
)
if
(
stride
==
-
1
)
{
// give a chance if stride is
0
, return a default packed stride
// give a chance if stride is
-1
, return a default packed stride
if
constexpr
(
std
::
is_same_v
<
decltype
(
layout
),
ck
::
tensor_layout
::
gemm
::
RowMajor
>
)
{
return
static_cast
<
std
::
size_t
>
(
col
);
...
...
example/01_gemm/run_gemm_example_v2.inc
View file @
b0376100
...
...
@@ -115,21 +115,21 @@ bool run_gemm(const ProblemType& problem_size, const ExecutionConfig& config)
};
auto
f_get_default_stride
=
[](
std
::
size_t
row
,
std
::
size_t
col
,
std
::
size
_t
stride
,
auto
layout
)
{
if
(
stride
==
0
)
[](
std
::
size_t
row
,
std
::
size_t
col
,
ck
::
index
_t
stride
,
auto
layout
)
{
if
(
stride
==
-
1
)
{
// give a chance if stride is
zero
, return a default packed stride
// give a chance if stride is
-1
, return a default packed stride
if
constexpr
(
std
::
is_same_v
<
decltype
(
layout
),
ck
::
tensor_layout
::
gemm
::
RowMajor
>
)
{
return
col
;
return
static_cast
<
std
::
size_t
>
(
col
)
;
}
else
{
return
row
;
return
static_cast
<
std
::
size_t
>
(
row
)
;
}
}
else
return
stride
;
return
static_cast
<
std
::
size_t
>
(
stride
)
;
};
StrideA
=
f_get_default_stride
(
M
,
K
,
StrideA
,
ALayout
{});
...
...
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