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
OpenDAS
tilelang
Commits
34a94d42
Commit
34a94d42
authored
Feb 24, 2025
by
Lei Wang
Committed by
GitHub
Feb 24, 2025
Browse files
[Typo] Fix a typo in gemm splitk examples (#111)
parent
5cea760c
Changes
1
Hide whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
6 additions
and
6 deletions
+6
-6
examples/gemm_splitk/example_tilelang_gemm_splitk.py
examples/gemm_splitk/example_tilelang_gemm_splitk.py
+6
-6
No files found.
examples/gemm_splitk/example_tilelang_gemm_splitk.py
View file @
34a94d42
...
@@ -18,9 +18,9 @@ def matmul(M, N, K, block_M, block_N, block_K, split_k, dtype="float16", accum_d
...
@@ -18,9 +18,9 @@ def matmul(M, N, K, block_M, block_N, block_K, split_k, dtype="float16", accum_d
):
):
with
T
.
Kernel
(
with
T
.
Kernel
(
T
.
ceildiv
(
N
,
block_N
),
T
.
ceildiv
(
M
,
block_M
),
split_k
,
threads
=
128
)
as
(
bx
,
by
,
bz
):
T
.
ceildiv
(
N
,
block_N
),
T
.
ceildiv
(
M
,
block_M
),
split_k
,
threads
=
128
)
as
(
bx
,
by
,
bz
):
A_shared
=
T
.
alloc_shared
((
block_M
,
block_K
),
dtype
,
"shared"
)
A_shared
=
T
.
alloc_shared
((
block_M
,
block_K
),
dtype
)
B_shared
=
T
.
alloc_shared
((
block_K
,
block_N
),
dtype
,
"shared"
)
B_shared
=
T
.
alloc_shared
((
block_K
,
block_N
),
dtype
)
C_shared
=
T
.
alloc_shared
((
block_M
,
block_N
),
dtype
,
"shared"
)
C_shared
=
T
.
alloc_shared
((
block_M
,
block_N
),
dtype
)
C_local
=
T
.
alloc_fragment
((
block_M
,
block_N
),
accum_dtype
)
C_local
=
T
.
alloc_fragment
((
block_M
,
block_N
),
accum_dtype
)
if
bz
==
0
:
if
bz
==
0
:
...
@@ -42,9 +42,9 @@ def matmul(M, N, K, block_M, block_N, block_K, split_k, dtype="float16", accum_d
...
@@ -42,9 +42,9 @@ def matmul(M, N, K, block_M, block_N, block_K, split_k, dtype="float16", accum_d
m
,
n
=
by
*
block_M
+
i
,
bx
*
block_N
+
j
*
2
m
,
n
=
by
*
block_M
+
i
,
bx
*
block_N
+
j
*
2
# vectorized atomic
# vectorized atomic
T
.
atomic_addx2
(
C
[
m
,
n
],
C_shared
[
i
,
j
*
2
])
T
.
atomic_addx2
(
C
[
m
,
n
],
C_shared
[
i
,
j
*
2
])
else
:
else
:
for
i
,
j
in
T
.
Parallel
(
block_M
,
block_N
):
for
i
,
j
in
T
.
Parallel
(
block_M
,
block_N
):
T
.
atomic_add
(
C
[
by
*
block_M
+
i
,
bx
*
block_N
+
j
],
C_shared
[
i
,
j
])
T
.
atomic_add
(
C
[
by
*
block_M
+
i
,
bx
*
block_N
+
j
],
C_shared
[
i
,
j
])
return
main
return
main
...
...
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