"git@developer.sourcefind.cn:yangql/composable_kernel.git" did not exist on "11ec07e9d13c41ea8c1512f86414fd0096a0e095"
Unverified Commit 4efd2d2d authored by Lei Wang's avatar Lei Wang Committed by GitHub
Browse files

[CI] allow dirty workspace for `format.sh` and introduce loop carry thread sync unit test (#1153)

* atomic_fix

* atomic_fix

* mem fix

* lint fix

* add some comments

* fix

* fix

* lint fix

* handle async copy

* lint fix

* lint fix
parent f7ba45d8
...@@ -29,10 +29,7 @@ ALL_FILES='' ...@@ -29,10 +29,7 @@ ALL_FILES=''
ONLY_CHANGED='' ONLY_CHANGED=''
FILES=() FILES=()
if (($# == 0)); then if (($# == 0)); then
if [[ -n "$(git status --porcelain --ignore-submodules --untracked-files=no)" ]]; then # Default: allow dirty workspace; run on changed files (committed + worktree)
echo "Detected uncommitted changes. Please commit or stash them before running $0." >&2
exit 1
fi
ONLY_CHANGED='true' ONLY_CHANGED='true'
else else
while (($# > 0)); do while (($# > 0)); do
...@@ -78,7 +75,7 @@ if [[ -n "${ALL_FILES}" ]]; then ...@@ -78,7 +75,7 @@ if [[ -n "${ALL_FILES}" ]]; then
echo "Checking all files..." >&2 echo "Checking all files..." >&2
elif [[ -n "${ONLY_CHANGED}" ]]; then elif [[ -n "${ONLY_CHANGED}" ]]; then
MERGE_BASE="$(get_merge_base)" MERGE_BASE="$(get_merge_base)"
echo "Checking changed files compared to merge base (${MERGE_BASE})..." >&2 echo "Checking changed files vs merge base (${MERGE_BASE}) and working tree..." >&2
elif [[ "${#FILES[@]}" -gt 0 ]]; then elif [[ "${#FILES[@]}" -gt 0 ]]; then
echo "Checking specified files: ${FILES[*]}..." >&2 echo "Checking specified files: ${FILES[*]}..." >&2
fi fi
...@@ -93,7 +90,17 @@ echo 'tile-lang pre-commit: Check Start' ...@@ -93,7 +90,17 @@ echo 'tile-lang pre-commit: Check Start'
if [[ -n "${ALL_FILES}" ]]; then if [[ -n "${ALL_FILES}" ]]; then
python3 -m pre_commit run --all-files python3 -m pre_commit run --all-files
elif [[ -n "${ONLY_CHANGED}" ]]; then elif [[ -n "${ONLY_CHANGED}" ]]; then
python3 -m pre_commit run --from-ref "${MERGE_BASE}" --to-ref HEAD # Collect changed files (committed since merge-base + current worktree)
CHANGED_FILES="$(git diff --name-only --diff-filter=ACM "${MERGE_BASE}" 2>/dev/null || true)"
if [[ -n "${CHANGED_FILES}" ]]; then
echo "Running pre-commit on changed files:"
echo "${CHANGED_FILES}"
# Convert newline-separated files to space-separated and run pre-commit once
CHANGED_FILES_SPACE="$(echo "${CHANGED_FILES}" | tr '\n' ' ')"
python3 -m pre_commit run --files ${CHANGED_FILES_SPACE}
else
echo "No files changed relative to merge base and worktree. Skipping pre-commit."
fi
elif [[ "${#FILES[@]}" -gt 0 ]]; then elif [[ "${#FILES[@]}" -gt 0 ]]; then
python3 -m pre_commit run --files "${FILES[@]}" python3 -m pre_commit run --files "${FILES[@]}"
fi fi
......
...@@ -188,5 +188,41 @@ def test_sync_let_stmt(): ...@@ -188,5 +188,41 @@ def test_sync_let_stmt():
tvm.ir.assert_structural_equal(mod["main"], expected) tvm.ir.assert_structural_equal(mod["main"], expected)
@tilelang.testing.requires_cuda
def test_sync_shared_dyn_stmatrix_loop_hoist():
@T.prim_func
def func():
buf_dyn_shmem = T.alloc_buffer((98304,), "uint8", scope="shared.dyn")
tx = T.launch_thread("threadIdx.x", 384)
for i in T.unroll(8):
off = (
i // 4 * 8192 + tx // 32 * 1024 + tx % 16 * 64 +
(tx % 8 // 4 + i % 4 // 2) % 2 * 32 + (tx % 4 // 2 + i % 2) % 2 * 16 +
(tx % 32 // 16 + tx % 2) % 2 * 8)
T.evaluate(
T.call_intrin(
"handle",
tvm.tir.op.Op.get("tl.ptx_stmatrix"),
T.int32(0),
T.int32(4),
T.tvm_access_ptr(
T.type_annotation("uint8"),
buf_dyn_shmem.data,
off,
98304 - off,
2,
),
T.int32(2),
))
mod = tvm.IRModule({"main": func})
mod = tilelang.transform.ThreadSync("shared.dyn")(mod)
s = str(mod)
assert 'T.tvm_storage_sync("shared.dyn")' in s
# Ensure the sync appears before the unrolled loop
assert s.index('T.tvm_storage_sync("shared.dyn")') < s.index("for i in T.unroll(8)")
if __name__ == "__main__": if __name__ == "__main__":
tilelang.testing.main() tilelang.testing.main()
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment