From c5b0a166c753b61e25d52433bdf94318a0349a20 Mon Sep 17 00:00:00 2001 From: Xuehai Pan Date: Wed, 15 Oct 2025 22:31:59 +0800 Subject: [PATCH 1/7] [Lint] Retire `format.sh` and add `clang-tidy` to GHA workflow --- .github/workflows/ci.yml | 30 ++-- .gitignore | 4 + format.sh | 331 --------------------------------------- 3 files changed, 22 insertions(+), 343 deletions(-) delete mode 100755 format.sh diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index 1782cedf3..94f37a603 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -287,21 +287,27 @@ jobs: echo "Clearing uv cache at ${UV_CACHE_DIR} due to failure." uv cache clean - - name: Run format check - id: format-check + - name: Run clang-tidy + id: clang-tidy run: | - mkdir -p build + clang-tidy --version + # Run cmake to create the build directory with compile_commands.json - ( - cd build - cmake .. ${CLANG_TIDY_CMAKE_OPTIONS} # no quotes here - ) + cmake -S . -B cmake-build --fresh ${CLANG_TIDY_CMAKE_OPTIONS} # no quotes here + + CXX_FILES=$(find src -type f -iname "*.[ch]pp" -o -iname "*.cc" -o -iname "*.c" -o -iname "*.h") rc=0 - bash format.sh || rc="$?" - rm -rf build - if [[ "${rc}" -ne 0 ]]; then - echo "::error::Format check failed. Please run 'bash format.sh' locally to fix the issues." - exit 1 + if [[ -x "$(command -v run-clang-tidy)" ]]; then + run-clang-tidy -clang-tidy-binary="$(command -v clang-tidy)" \ + -fix -p="cmake-build" ${CXX_FILES} || rc="$?" + else + clang-tidy --fix -p="cmake-build" ${CXX_FILES} || rc="$?" + fi + rm -rf cmake-build + if (( rc != 0 )); then + echo "::error::clang-tidy found issues (exit code: ${rc}). Please run 'clang-tidy --fix' locally to fix them." + git diff --color=always || true + exit "${rc}" fi - name: Enable core dump generation (Linux / GitHub-hosted runners) diff --git a/.gitignore b/.gitignore index 042b791ca..b7421d77e 100644 --- a/.gitignore +++ b/.gitignore @@ -97,3 +97,7 @@ tilelang/jit/adapter/cython/.cycache # claude **/.claude + +# CMake +cmake-build/ +cmake-build-*/ diff --git a/format.sh b/format.sh deleted file mode 100755 index 565569959..000000000 --- a/format.sh +++ /dev/null @@ -1,331 +0,0 @@ -#!/usr/bin/env bash -# Usage: -# # Do work and commit your work. - -# # Format files that differ from origin/main. -# bash format.sh - -# # Commit changed files with message 'Run yapf and ruff' -# -# -# YAPF + Clang formatter (if installed). This script formats all changed files from the last mergebase. -# You are encouraged to run this locally before pushing changes for review. - -# Cause the script to exit if a single command fails -set -eo pipefail - -# this stops git rev-parse from failing if we run this from the .git directory -builtin cd "$(dirname "${BASH_SOURCE:-$0}")" -ROOT="$(git rev-parse --show-toplevel)" -builtin cd "$ROOT" || exit 1 - -# If yapf/ruff/codespell is not installed, install according to the requirements -if ! (yapf --version &>/dev/null && ruff --version &>/dev/null && codespell --version &>/dev/null); then - pip install -r requirements-lint.txt -fi - -YAPF_VERSION=$(yapf --version | awk '{print $2}') -RUFF_VERSION=$(ruff --version | awk '{print $2}') -CODESPELL_VERSION=$(codespell --version) - -# # params: tool name, tool version, required version -tool_version_check() { - if [[ $2 != $3 ]]; then - echo "Wrong $1 version installed: $3 is required, not $2." - pip install -r requirements-lint.txt - fi -} - -tool_version_check "yapf" $YAPF_VERSION "$(grep yapf requirements-lint.txt | cut -d'=' -f3)" -tool_version_check "ruff" $RUFF_VERSION "$(grep "ruff==" requirements-lint.txt | cut -d'=' -f3)" -tool_version_check "codespell" "$CODESPELL_VERSION" "$(grep codespell requirements-lint.txt | cut -d'=' -f3)" - -echo 'tile-lang yapf: Check Start' - -YAPF_FLAGS=( - '--recursive' - '--parallel' -) - -YAPF_EXCLUDES=( - '--exclude' 'build/**' - '--exclude' '3rdparty/**' -) - -# Format specified files -format() { - yapf --in-place "${YAPF_FLAGS[@]}" "$@" -} - -# Format files that differ from main branch. Ignores dirs that are not slated -# for autoformat yet. -format_changed() { - # The `if` guard ensures that the list of filenames is not empty, which - # could cause yapf to receive 0 positional arguments, making it hang - # waiting for STDIN. - # - # `diff-filter=ACM` and $MERGEBASE is to ensure we only format files that - # exist on both branches. - UPSTREAM_REPO="https://github.com/tile-ai/tilelang" - - if git ls-remote --exit-code "$UPSTREAM_REPO" main &>/dev/null; then - # First try to use the upstream repository directly - MERGEBASE="$(git fetch "$UPSTREAM_REPO" main &>/dev/null && git merge-base FETCH_HEAD HEAD)" - elif git show-ref --verify --quiet refs/remotes/origin/main; then - # Fall back to origin/main if available - BASE_BRANCH="origin/main" - MERGEBASE="$(git merge-base $BASE_BRANCH HEAD)" - else - # Last resort, use local main - BASE_BRANCH="main" - MERGEBASE="$(git merge-base $BASE_BRANCH HEAD)" - fi - - if ! git diff --diff-filter=ACM --quiet --exit-code "$MERGEBASE" -- '*.py' '*.pyi' &>/dev/null; then - git diff --name-only --diff-filter=ACM "$MERGEBASE" -- '*.py' '*.pyi' | xargs -P 5 \ - yapf --in-place "${YAPF_EXCLUDES[@]}" "${YAPF_FLAGS[@]}" - fi -} - -# Format all files -format_all() { - yapf --in-place "${YAPF_FLAGS[@]}" "${YAPF_EXCLUDES[@]}" . -} - -## This flag formats individual files. --files *must* be the first command line -## arg to use this option. -if [[ "$1" == '--files' ]]; then - format "${@:2}" - # If `--all` is passed, then any further arguments are ignored and the - # entire python directory is formatted. -elif [[ "$1" == '--all' ]]; then - format_all -else - # Format only the files that changed in last commit. - format_changed -fi -echo 'tile-lang yapf: Done' - -echo 'tile-lang codespell: Check Start' -# check spelling of specified files -spell_check() { - codespell "$@" -} - -spell_check_all(){ - codespell --toml pyproject.toml -} - -# Spelling check of files that differ from main branch. -spell_check_changed() { - # The `if` guard ensures that the list of filenames is not empty, which - # could cause ruff to receive 0 positional arguments, making it hang - # waiting for STDIN. - # - # `diff-filter=ACM` and $MERGEBASE is to ensure we only lint files that - # exist on both branches. - if git show-ref --verify --quiet refs/remotes/origin/main; then - BASE_BRANCH="origin/main" - else - BASE_BRANCH="main" - fi - - MERGEBASE="$(git merge-base $BASE_BRANCH HEAD)" - - if ! git diff --diff-filter=ACM --quiet --exit-code "$MERGEBASE" -- '*.py' '*.pyi' &>/dev/null; then - git diff --name-only --diff-filter=ACM "$MERGEBASE" -- '*.py' '*.pyi' | xargs \ - codespell - fi -} - -# Run Codespell -## This flag runs spell check of individual files. --files *must* be the first command line -## arg to use this option. -if [[ "$1" == '--files' ]]; then - spell_check "${@:2}" - # If `--all` is passed, then any further arguments are ignored and the - # entire python directory is linted. -elif [[ "$1" == '--all' ]]; then - spell_check_all -else - # Check spelling only of the files that changed in last commit. - spell_check_changed -fi -echo 'tile-lang codespell: Done' - -echo 'tile-lang ruff: Check Start' -# Lint specified files -lint() { - ruff check "$@" -} - -# Lint files that differ from main branch. Ignores dirs that are not slated -# for autolint yet. -lint_changed() { - # The `if` guard ensures that the list of filenames is not empty, which - # could cause ruff to receive 0 positional arguments, making it hang - # waiting for STDIN. - # - # `diff-filter=ACM` and $MERGEBASE is to ensure we only lint files that - # exist on both branches. - if git show-ref --verify --quiet refs/remotes/origin/main; then - BASE_BRANCH="origin/main" - else - BASE_BRANCH="main" - fi - - MERGEBASE="$(git merge-base $BASE_BRANCH HEAD)" - - if ! git diff --diff-filter=ACM --quiet --exit-code "$MERGEBASE" -- '*.py' '*.pyi' &>/dev/null; then - git diff --name-only --diff-filter=ACM "$MERGEBASE" -- '*.py' '*.pyi' | xargs \ - ruff check - fi - -} - -# Run Ruff -### This flag lints individual files. --files *must* be the first command line -### arg to use this option. -if [[ "$1" == '--files' ]]; then - lint "${@:2}" - # If `--all` is passed, then any further arguments are ignored and the - # entire python directory is linted. -elif [[ "$1" == '--all' ]]; then - lint python testing -else - # Format only the files that changed in last commit. - lint_changed -fi - -echo 'tile-lang ruff: Done' - -echo 'tile-lang clang-format: Check Start' -# If clang-format is available, run it; otherwise, skip -if command -v clang-format &>/dev/null; then - CLANG_FORMAT_VERSION=$(clang-format --version | awk '{print $3}') - tool_version_check "clang-format" "$CLANG_FORMAT_VERSION" "$(grep clang-format requirements-lint.txt | cut -d'=' -f3)" - - CLANG_FORMAT_FLAGS=("-i") - - # Apply clang-format to specified files - clang_format() { - clang-format "${CLANG_FORMAT_FLAGS[@]}" "$@" - } - - # Format all C/C++ files in the repo, excluding specified directories - clang_format_all() { - find . -type f \( -name '*.c' -o -name '*.cc' -o -name '*.cpp' -o -name '*.h' -o -name '*.hpp' \) \ - -not -path "./3rdparty/*" \ - -not -path "./build/*" \ - -exec clang-format -i {} + - } - - # Format changed C/C++ files relative to main - clang_format_changed() { - if git show-ref --verify --quiet refs/remotes/origin/main; then - BASE_BRANCH="origin/main" - else - BASE_BRANCH="main" - fi - - MERGEBASE="$(git merge-base $BASE_BRANCH HEAD)" - - if ! git diff --diff-filter=ACM --quiet --exit-code "$MERGEBASE" -- '*.c' '*.cc' '*.cpp' '*.h' '*.hpp' &>/dev/null; then - git diff --name-only --diff-filter=ACM "$MERGEBASE" -- '*.c' '*.cc' '*.cpp' '*.h' '*.hpp' | xargs clang-format -i - fi - } - - if [[ "$1" == '--files' ]]; then - # If --files is given, format only the provided files - clang_format "${@:2}" - elif [[ "$1" == '--all' ]]; then - # If --all is given, format all eligible C/C++ files - clang_format_all - else - # Otherwise, format only changed C/C++ files - clang_format_changed - fi -else - echo "clang-format not found. Skipping C/C++ formatting." -fi -echo 'tile-lang clang-format: Done' - -echo 'tile-lang clang-tidy: Check Start' -# If clang-tidy is available, run it; otherwise, skip -if command -v run-clang-tidy &>/dev/null; then - # Check if clang-tidy is available - if ! command -v clang-tidy &>/dev/null; then - echo "clang-tidy not found. Skipping clang-tidy checks." - else - # Get clang-tidy version - CLANG_TIDY_VERSION=$(clang-tidy --version | head -n1 | awk '{print $4}') - echo "Using clang-tidy version: $CLANG_TIDY_VERSION" - - # Check if build directory exists - if [ ! -d "build" ]; then - echo "Build directory not found. Skipping clang-tidy checks." - else - # Run clang-tidy on specified files - clang_tidy_files() { - run-clang-tidy -j 64 "$@" -p build - } - - # Run clang-tidy on all C/C++ source files - clang_tidy_all() { - run-clang-tidy -j 64 src/*.cc -p build - } - - # Run clang-tidy on changed C/C++ files relative to main - clang_tidy_changed() { - if git show-ref --verify --quiet refs/remotes/origin/main; then - BASE_BRANCH="origin/main" - else - BASE_BRANCH="main" - fi - - MERGEBASE="$(git merge-base $BASE_BRANCH HEAD)" - - # Get changed C/C++ files - CHANGED_FILES=$(git diff --name-only --diff-filter=ACM "$MERGEBASE" -- '*.c' '*.cc' '*.cpp' '*.h' '*.hpp' 2>/dev/null || true) - - if [ -n "$CHANGED_FILES" ]; then - echo "Running clang-tidy on changed files:" - echo "$CHANGED_FILES" - # Convert newline-separated files to space-separated and run clang-tidy once - CHANGED_FILES_SPACE=$(echo "$CHANGED_FILES" | tr '\n' ' ') - run-clang-tidy -j 64 $CHANGED_FILES_SPACE -p build -fix - else - echo "No C/C++ files changed. Skipping clang-tidy." - fi - } - - if [[ "$1" == '--files' ]]; then - # If --files is given, run clang-tidy only on the provided files - clang_tidy_files "${@:2}" - elif [[ "$1" == '--all' ]]; then - # If --all is given, run clang-tidy on all source files - clang_tidy_all - else - # Otherwise, run clang-tidy only on changed C/C++ files - clang_tidy_changed - fi - fi - fi -else - echo "run-clang-tidy not found. Skipping clang-tidy checks." - echo "To install clang-tidy tools, you may need to install clang-tidy and run-clang-tidy." -fi -echo 'tile-lang clang-tidy: Done' - -# Check if there are any uncommitted changes after all formatting steps. -# If there are, ask the user to review and stage them. -if ! git diff --quiet &>/dev/null; then - echo 'Reformatted files. Please review and stage the changes.' - echo 'Changes not staged for commit:' - echo - git --no-pager diff --name-only - - exit 1 -fi - -echo 'tile-lang: All checks passed' From 28b878bfd802314fb98a45a155c53e28a33604de Mon Sep 17 00:00:00 2001 From: Xuehai Pan Date: Wed, 15 Oct 2025 22:41:56 +0800 Subject: [PATCH 2/7] chore: update clang-tidy settings --- .clang-tidy | 15 ++++++++++----- 1 file changed, 10 insertions(+), 5 deletions(-) diff --git a/.clang-tidy b/.clang-tidy index b9c6cc54c..3fb5411da 100644 --- a/.clang-tidy +++ b/.clang-tidy @@ -1,4 +1,13 @@ -Checks: > +--- +InheritParentConfig: true +ExtraArgs: ['-v'] +FormatStyle: file +UseColor: true +WarningsAsErrors: '*' +HeaderFilterRegex: '^(?!.*(3rdparty|build)).*$' + +# NOTE: there must be no spaces before the '-', so put the comma last. +Checks: >- # 1. Retained categories: easier to find bugs/performance issues clang-analyzer-*, cppcoreguidelines-pro-type-static-cast-downcast, @@ -47,7 +56,3 @@ Checks: > -clang-analyzer-deadcode.DeadStores, -clang-analyzer-optin.cplusplus.VirtualCall, -clang-diagnostic-tautological-constant-compare, - -WarningsAsErrors: '*' - -HeaderFilterRegex: '^(?!.*(3rdparty|build)).*$' \ No newline at end of file From 6b9dd88ef0a9c6990432ea84b71dd041f831b7cd Mon Sep 17 00:00:00 2001 From: Xuehai Pan Date: Wed, 15 Oct 2025 23:57:48 +0800 Subject: [PATCH 3/7] chore: upgrade clang-format and clang-tidy version --- .clang-tidy | 1 - .github/workflows/ci.yml | 29 +++++++++++++++++------- .pre-commit-config.yaml | 2 +- requirements-lint.txt | 4 ++-- src/op/gemm.cc | 4 +--- src/op/parallel.h | 2 +- src/target/codegen_cuda.cc | 4 ++-- src/target/cuda.h | 4 ++-- src/tl_templates/cpp/half.hpp | 11 +++++---- src/tl_templates/cuda/common.h | 4 ++-- src/transform/cluster_planning.cc | 4 ++-- src/transform/common/loop_fusion_utils.h | 2 +- src/transform/layout_inference.cc | 2 +- 13 files changed, 42 insertions(+), 31 deletions(-) diff --git a/.clang-tidy b/.clang-tidy index 3fb5411da..7590818e7 100644 --- a/.clang-tidy +++ b/.clang-tidy @@ -1,6 +1,5 @@ --- InheritParentConfig: true -ExtraArgs: ['-v'] FormatStyle: file UseColor: true WarningsAsErrors: '*' diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index 94f37a603..03781cdb8 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -226,6 +226,7 @@ jobs: echo "PIP_EXTRA_INDEX_URL=${PIP_EXTRA_INDEX_URL}" | tee -a "${GITHUB_ENV}" echo "UV_INDEX=${UV_INDEX}" | tee -a "${GITHUB_ENV}" fi + export CLANG_TIDY_CMAKE_OPTIONS="${CLANG_TIDY_CMAKE_OPTIONS} -DUSE_METAL=ON" echo "USE_METAL=ON" | tee -a "${GITHUB_ENV}" @@ -289,21 +290,33 @@ jobs: - name: Run clang-tidy id: clang-tidy + if: runner.os == 'Linux' run: | - clang-tidy --version + echo "\$ $(command -v clang-tidy) --version" && clang-tidy --version + + if [[ -x "$(command -v run-clang-tidy)" ]]; then + echo "Using run-clang-tidy from $(command -v run-clang-tidy)" + CLANG_TIDY=(run-clang-tidy) + else + echo "Downloading run-clang-tidy script" + wget -O run-clang-tidy.py https://raw.githubusercontent.com/llvm/llvm-project/refs/heads/release/21.x/clang-tools-extra/clang-tidy/tool/run-clang-tidy.py + CLANG_TIDY=(uv run --no-project --script -- run-clang-tidy.py) + fi + if [[ -x "$(command -v clang-apply-replacements)" ]]; then + echo "Using clang-apply-replacements from $(command -v clang-apply-replacements)" + CLANG_TIDY+=(-fix -clang-apply-replacements-binary="$(command -v clang-apply-replacements)") + else + echo "::warning::clang-apply-replacements not found in PATH, automatic fixing disabled." + fi # Run cmake to create the build directory with compile_commands.json cmake -S . -B cmake-build --fresh ${CLANG_TIDY_CMAKE_OPTIONS} # no quotes here CXX_FILES=$(find src -type f -iname "*.[ch]pp" -o -iname "*.cc" -o -iname "*.c" -o -iname "*.h") rc=0 - if [[ -x "$(command -v run-clang-tidy)" ]]; then - run-clang-tidy -clang-tidy-binary="$(command -v clang-tidy)" \ - -fix -p="cmake-build" ${CXX_FILES} || rc="$?" - else - clang-tidy --fix -p="cmake-build" ${CXX_FILES} || rc="$?" - fi - rm -rf cmake-build + "${CLANG_TIDY[@]}" -clang-tidy-binary="$(command -v clang-tidy)" \ + -p="cmake-build" ${CXX_FILES} || rc="$?" + rm -rf cmake-build run-clang-tidy.py if (( rc != 0 )); then echo "::error::clang-tidy found issues (exit code: ${rc}). Please run 'clang-tidy --fix' locally to fix them." git diff --color=always || true diff --git a/.pre-commit-config.yaml b/.pre-commit-config.yaml index 391c7796e..1ab1713b8 100644 --- a/.pre-commit-config.yaml +++ b/.pre-commit-config.yaml @@ -32,7 +32,7 @@ repos: args: [--ignore-case] files: ^docs/spelling_wordlist\.txt$ - repo: https://github.com/pre-commit/mirrors-clang-format - rev: v15.0.7 # sync with requirements-lint.txt + rev: v21.1.2 # sync with requirements-lint.txt hooks: - id: clang-format exclude: | diff --git a/requirements-lint.txt b/requirements-lint.txt index 1cd2a7b1e..c2db20541 100644 --- a/requirements-lint.txt +++ b/requirements-lint.txt @@ -1,7 +1,7 @@ # Format and lint requirements pre-commit -clang-format==15.0.7 -clang-tidy==18.1.8 +clang-format==21.1.2 +clang-tidy==21.1.1 codespell[toml]==2.4.1 ruff==0.14.0 yapf==0.43.0 diff --git a/src/op/gemm.cc b/src/op/gemm.cc index 75c977c8b..8912a7a33 100644 --- a/src/op/gemm.cc +++ b/src/op/gemm.cc @@ -27,9 +27,7 @@ static inline std::pair GetTCGEN5MMAMeta(int M, int N, int K, DataType ab_dtype, DataType c_dtype) { // TODO (lei) Currently not all shapes / dtypes are supported for TCGEN5MMA. #define FAIL \ - return { \ - false, TCGEN5MMAMeta { 0, 0, 0 } \ - } + return { false, TCGEN5MMAMeta{0, 0, 0} } #define SUCCESS(atom_m, atom_n, atom_k) \ return { \ true, TCGEN5MMAMeta { atom_m, atom_n, atom_k } \ diff --git a/src/op/parallel.h b/src/op/parallel.h index 5f1f5a887..9c6b7180f 100644 --- a/src/op/parallel.h +++ b/src/op/parallel.h @@ -42,7 +42,7 @@ class ParallelOpNode; class ParallelLoopNestVisitor : public StmtExprVisitor { private: - ParallelLoopNestVisitor(ParallelOpNode *op) : p(op){}; + ParallelLoopNestVisitor(ParallelOpNode *op) : p(op) {}; void VisitStmt_(const ForNode *op) override; void VisitStmt_(const BufferStoreNode *op) override; void VisitExpr_(const BufferLoadNode *op) override; diff --git a/src/target/codegen_cuda.cc b/src/target/codegen_cuda.cc index d06e7170d..4af9ed411 100644 --- a/src/target/codegen_cuda.cc +++ b/src/target/codegen_cuda.cc @@ -1749,8 +1749,8 @@ void CodeGenTileLangCUDA::VisitExpr_(const CallNode *op, std::ostream &os) { os << "}\n"; } else { os << "for (int local_id = 0; local_id < 8; ++local_id) {\n"; - os << dst << "[" + this->PrintExpr(dst_ind) + "]" - << " = " << src << "[" << src_offset << " + local_id];\n"; + os << dst << "[" + this->PrintExpr(dst_ind) + "]" << " = " << src << "[" + << src_offset << " + local_id];\n"; os << "}\n"; } diff --git a/src/target/cuda.h b/src/target/cuda.h index 010a1a228..a9dfb13ab 100644 --- a/src/target/cuda.h +++ b/src/target/cuda.h @@ -5023,12 +5023,12 @@ typedef struct CUgraphNodeParams_st { /** * Device that represents the CPU */ -#define CU_DEVICE_CPU ((CUdevice)-1) +#define CU_DEVICE_CPU ((CUdevice) - 1) /** * Device that represents an invalid device */ -#define CU_DEVICE_INVALID ((CUdevice)-2) +#define CU_DEVICE_INVALID ((CUdevice) - 2) /** * Bitmasks for ::CU_DEVICE_ATTRIBUTE_GPU_DIRECT_RDMA_FLUSH_WRITES_OPTIONS diff --git a/src/tl_templates/cpp/half.hpp b/src/tl_templates/cpp/half.hpp index 0107b3d44..5410e7572 100644 --- a/src/tl_templates/cpp/half.hpp +++ b/src/tl_templates/cpp/half.hpp @@ -1192,8 +1192,8 @@ unsigned int float2half_impl(T value, ...) { template unsigned int float2half(T value) { return float2half_impl( - value, bool_type < std::numeric_limits::is_iec559 && - sizeof(typename bits::type) == sizeof(T) > ()); + value, bool_type::is_iec559 && + sizeof(typename bits::type) == sizeof(T)>()); } /// Convert integer to half-precision floating-point. @@ -1665,9 +1665,10 @@ template T half2float_impl(unsigned int value, T, ...) { /// \param value half-precision value to convert /// \return floating-point value template T half2float(unsigned int value) { - return half2float_impl(value, T(), - bool_type < std::numeric_limits::is_iec559 && - sizeof(typename bits::type) == sizeof(T) > ()); + return half2float_impl( + value, T(), + bool_type::is_iec559 && + sizeof(typename bits::type) == sizeof(T)>()); } /// Convert half-precision floating-point to integer. diff --git a/src/tl_templates/cuda/common.h b/src/tl_templates/cuda/common.h index 34a30821b..dfbc062cf 100644 --- a/src/tl_templates/cuda/common.h +++ b/src/tl_templates/cuda/common.h @@ -244,8 +244,8 @@ union GmmaDescriptor { uint16_t stride_byte_offset_ : 14, : 2; // 14 bits [0,14), 2 bits unused // base_offset, bit [49,52) // Valid only for SWIZZLE_128B and SWIZZLE_64B - uint8_t : 1, - base_offset_ : 3, : 4; // 1 bit unused, 3 bits [1,4), 4 bits unused + uint8_t : 1, base_offset_ : 3, + : 4; // 1 bit unused, 3 bits [1,4), 4 bits unused // layout type, bit [62,64) // SWIZZLE_NONE = 0, SWIZZLE_32B = 3, SWIZZLE_64B = 2, SWIZZLE_128B = 1 uint8_t : 6, layout_type_ : 2; // 6 bits unused, 2 bits [6,8) diff --git a/src/transform/cluster_planning.cc b/src/transform/cluster_planning.cc index d88af71e2..e847bb2b6 100644 --- a/src/transform/cluster_planning.cc +++ b/src/transform/cluster_planning.cc @@ -86,14 +86,14 @@ class ClusterPlanner { class RegionVisitor : public ExprVisitor { public: - RegionVisitor(){}; + RegionVisitor() {}; void VisitExpr_(const VarNode *var) { seen_.insert(var); } std::unordered_set seen_; }; class BlockIdxVisitor : public StmtVisitor { public: - BlockIdxVisitor(){}; + BlockIdxVisitor() {}; void VisitStmt_(const AttrStmtNode *attr) final { if (attr->attr_key == attr::thread_extent) { IterVar iv = Downcast(attr->node); diff --git a/src/transform/common/loop_fusion_utils.h b/src/transform/common/loop_fusion_utils.h index 9555e1e87..2fa6cdede 100644 --- a/src/transform/common/loop_fusion_utils.h +++ b/src/transform/common/loop_fusion_utils.h @@ -99,7 +99,7 @@ class ParallelLoopFuser : public IRMutatorWithAnalyzer { private: ParallelLoopFuser(arith::Analyzer *analyzer) - : IRMutatorWithAnalyzer(analyzer){}; + : IRMutatorWithAnalyzer(analyzer) {}; Stmt VisitStmt_(const ForNode *op) final { // Gather consecutive parallel loops diff --git a/src/transform/layout_inference.cc b/src/transform/layout_inference.cc index c903db271..329895a6b 100644 --- a/src/transform/layout_inference.cc +++ b/src/transform/layout_inference.cc @@ -636,7 +636,7 @@ class LayoutInferencer : public IRMutatorWithAnalyzer { LayoutInferencer(const LayoutInferenceResult &result, bool skip_thread_partition, arith::Analyzer *analyzer) : arith::IRMutatorWithAnalyzer(analyzer), result_(result), - skip_thread_partition_(skip_thread_partition){}; + skip_thread_partition_(skip_thread_partition) {}; using arith::IRMutatorWithAnalyzer::IRMutatorWithAnalyzer; From dfd1c304d7ab99a8bab22eca4f9381766e5ce34c Mon Sep 17 00:00:00 2001 From: Xuehai Pan Date: Thu, 16 Oct 2025 14:25:12 +0800 Subject: [PATCH 4/7] lint: resolve clang-tidy errors --- .clang-tidy | 3 +- src/layout/utils.cc | 2 +- src/runtime/runtime.cc | 2 +- src/target/codegen_webgpu.cc | 2 +- src/transform/layout_inference.cc | 6 ++-- src/transform/layout_reducer.cc | 2 +- src/transform/lower_thread_allreduce.cc | 9 ++--- src/transform/make_packed_api.cc | 2 +- src/transform/storage_access.cc | 2 +- src/transform/warp_specialized_rewriter.cc | 38 +++++++++++++--------- 10 files changed, 38 insertions(+), 30 deletions(-) diff --git a/.clang-tidy b/.clang-tidy index 7590818e7..2ddbefbf9 100644 --- a/.clang-tidy +++ b/.clang-tidy @@ -1,9 +1,10 @@ --- InheritParentConfig: true +ExtraArgs: ['-v'] FormatStyle: file UseColor: true WarningsAsErrors: '*' -HeaderFilterRegex: '^(?!.*(3rdparty|build)).*$' +ExcludeHeaderFilterRegex: '^(3rdparty|tvm)/.*$' # NOTE: there must be no spaces before the '-', so put the comma last. Checks: >- diff --git a/src/layout/utils.cc b/src/layout/utils.cc index 83103fd1e..4d0461565 100644 --- a/src/layout/utils.cc +++ b/src/layout/utils.cc @@ -134,7 +134,7 @@ Array DivideUnusedIterators(const Array &exprs, for (const IterVar &iter : input_iters) { IterMark iv_mark; for (const IterMark &mark : collector.visited_) { - if (mark->source.as()->same_as(iter->var)) { + if (mark->source.as()->same_as(iter->var)) { // NOLINT(*) iv_mark = mark; break; } diff --git a/src/runtime/runtime.cc b/src/runtime/runtime.cc index 5d2f26278..3ea89d666 100644 --- a/src/runtime/runtime.cc +++ b/src/runtime/runtime.cc @@ -20,7 +20,7 @@ template static std::string ArrayToStr(const T *ptr, size_t n) { for (size_t i = 0; i < n; i++) { if (i > 0) ss << ", "; - ss << ptr[i]; + ss << ptr[i]; // NOLINT(clang-analyzer-security.ArrayBound) } ss << "]"; return ss.str(); diff --git a/src/target/codegen_webgpu.cc b/src/target/codegen_webgpu.cc index a88feaef0..1d64ccbc6 100644 --- a/src/target/codegen_webgpu.cc +++ b/src/target/codegen_webgpu.cc @@ -218,7 +218,7 @@ CodeGenTileLangWebGPU::AddFunction(const PrimFunc &f, bool skip_readonly_decl) { this->decl_stream << "\nstruct " << type_pod_args << " {\n"; for (size_t i = 0; i < pod_args.size(); ++i) { - Var v = pod_args[i]; + const Var &v = pod_args[i]; ICHECK(!v.dtype().is_handle()); std::string vid = AllocVarID(v.get()); diff --git a/src/transform/layout_inference.cc b/src/transform/layout_inference.cc index 329895a6b..c007d04ab 100644 --- a/src/transform/layout_inference.cc +++ b/src/transform/layout_inference.cc @@ -131,13 +131,13 @@ class BufferUseDefCollector : public IRVisitorWithAnalyzer { ICHECK(dst_layout_opt.has_value()) << "Failed to cast layout to Fragment for buffer " << buffer << ", layout type is " << layout->GetTypeKey(); - auto dst_layout = dst_layout_opt.value(); + const auto &dst_layout = dst_layout_opt.value(); auto src_layout_opt = layout_map[buffer].as(); ICHECK(src_layout_opt.has_value()) << "Failed to cast layout_map[buffer] to Fragment for buffer " << buffer << ", layout type is " << layout_map[buffer]->GetTypeKey(); - auto src_layout = src_layout_opt.value(); + const auto &src_layout = src_layout_opt.value(); ICHECK(dst_layout->InputDim() == src_layout->InputDim()); Array indices; indices.reserve(dst_layout->InputDim()); @@ -398,7 +398,7 @@ class BufferUseDefCollector : public IRVisitorWithAnalyzer { << call->args[1]->GetTypeKey(); return std::nullopt; } - auto var = var_opt.value(); + const auto &var = var_opt.value(); return buffer_data_to_buffer_[var]; } else if (call->op.same_as(RegionOp::Get())) { return call->args[0].as()->buffer; diff --git a/src/transform/layout_reducer.cc b/src/transform/layout_reducer.cc index 788e72a4d..e875c972c 100644 --- a/src/transform/layout_reducer.cc +++ b/src/transform/layout_reducer.cc @@ -209,7 +209,7 @@ class ReducerLayoutAnnotator : public IRMutatorWithAnalyzer { auto opt_buffer = var_to_buffer_.Get(reducer_var); ICHECK(opt_buffer); - auto buffer = opt_buffer.value(); + const auto &buffer = opt_buffer.value(); Fragment f; if (info->rep == ReducerRepType::ALL) { f = Fragment(buffer->shape, {}, ReplicationPlaceholder(), diff --git a/src/transform/lower_thread_allreduce.cc b/src/transform/lower_thread_allreduce.cc index d0c14219d..71ef8a92c 100644 --- a/src/transform/lower_thread_allreduce.cc +++ b/src/transform/lower_thread_allreduce.cc @@ -496,8 +496,9 @@ class ThreadAllreduceBuilder final : public StmtExprMutator { if (reduce_extent == 1) { // special case, no reduction is needed. std::vector stores; + stores.reserve(size); for (size_t i = 0; i < size; ++i) { - stores.push_back(BufferStore(buffers[i], values[i], {0})); + stores.emplace_back(BufferStore(buffers[i], values[i], {0})); } return SeqStmt::Flatten(stores); } @@ -604,7 +605,7 @@ class ThreadAllreduceBuilder final : public StmtExprMutator { // Load reduction values, no synchronization needed. Array a, b; for (int i = 0; i < n_buffers; ++i) { - Buffer shared_buf = shared_bufs[i]; + const Buffer &shared_buf = shared_bufs[i]; BufferLoad val(shared_buf, zero_indices); ICHECK_EQ(val->dtype, dtypes[i]); a.push_back(val); @@ -623,7 +624,7 @@ class ThreadAllreduceBuilder final : public StmtExprMutator { // branch with a warp sync call inside. PrimExpr other = WarpShuffle(builtin::tvm_warp_shuffle_down(), mask_buffer, val, offset); - Buffer local_buf = local_bufs[i]; + const Buffer &local_buf = local_bufs[i]; Stmt s = BufferStore(local_buf, other, zero_indices); seq->push_back(s); @@ -639,7 +640,7 @@ class ThreadAllreduceBuilder final : public StmtExprMutator { std::vector stores; stores.reserve(n_buffers); for (int i = 0; i < n_buffers; ++i) { - Buffer buf = shared_bufs[i]; + const Buffer &buf = shared_bufs[i]; stores.push_back(BufferStore(buf, ret[i], zero_indices)); } diff --git a/src/transform/make_packed_api.cc b/src/transform/make_packed_api.cc index a20b8fe38..a124027ce 100644 --- a/src/transform/make_packed_api.cc +++ b/src/transform/make_packed_api.cc @@ -477,7 +477,7 @@ tvm::transform::Pass MakePackedAPI() { Map packed_func_methods; for (const auto &[gvar, base_func] : mod->functions) { if (auto opt = base_func.as()) { - auto prim_func = opt.value(); + const auto &prim_func = opt.value(); if (auto global_symbol = RequiresPackedAPI(prim_func)) { packed_func_methods.Set(gvar, global_symbol.value()); } diff --git a/src/transform/storage_access.cc b/src/transform/storage_access.cc index 06340699a..0adaf712b 100644 --- a/src/transform/storage_access.cc +++ b/src/transform/storage_access.cc @@ -209,7 +209,7 @@ void TileLangStorageAccessVisitor::VisitStmt_(const ForNode *op) { bool IsThreadInvariant(const PrimExpr &cond) { if (auto call = cond.as()) { if (auto opt_call_op = call->op.as()) { - auto call_op = opt_call_op.value(); + const auto &call_op = opt_call_op.value(); if (call_op.same_as(builtin::tvm_thread_invariant())) { return true; } diff --git a/src/transform/warp_specialized_rewriter.cc b/src/transform/warp_specialized_rewriter.cc index 00844f0ef..b86ebaf96 100644 --- a/src/transform/warp_specialized_rewriter.cc +++ b/src/transform/warp_specialized_rewriter.cc @@ -530,10 +530,11 @@ class GroupOpRewriter : public StmtExprMutator { block_stmt.push_back(block->body); cur_id++; } - new_body.push_back(MakeGroupBlock(block_stmt.size() == 1 - ? block_stmt[0] - : SeqStmt(std::move(block_stmt)), - annotations)); + new_body.push_back(MakeGroupBlock( + block_stmt.size() == 1 ? block_stmt[0] + // NOLINTNEXTLINE(performance-move-const-arg) + : SeqStmt(std::move(block_stmt)), + annotations)); } Array order_anno; Array stage_anno; @@ -697,10 +698,12 @@ class WSCodeEmitter : public StmtMutator { continue; if (marker_.GetRole(op->seq[i]) == Role::kBoth) { block_stmt.push_back(seq_transformed[i]); - new_body.push_back(MakeGroupBlock( - block_stmt.size() == 1 ? block_stmt[0] - : SeqStmt(std::move(block_stmt)), - annotations)); + new_body.push_back( + MakeGroupBlock(block_stmt.size() == 1 + ? block_stmt[0] + // NOLINTNEXTLINE(performance-move-const-arg) + : SeqStmt(std::move(block_stmt)), + annotations)); continue; } } @@ -734,10 +737,12 @@ class WSCodeEmitter : public StmtMutator { } } collector.Clear(); - new_body.push_back(MakeGroupBlock( - block_stmt.size() == 1 ? block_stmt[0] - : SeqStmt(std::move(block_stmt)), - annotations)); + new_body.push_back( + MakeGroupBlock(block_stmt.size() == 1 + ? block_stmt[0] + // NOLINTNEXTLINE(performance-move-const-arg) + : SeqStmt(std::move(block_stmt)), + annotations)); } } } else { // consumer case @@ -766,10 +771,11 @@ class WSCodeEmitter : public StmtMutator { } } } - new_body.push_back(MakeGroupBlock(block_stmt.size() == 1 - ? block_stmt[0] - : SeqStmt(std::move(block_stmt)), - annotations)); + new_body.push_back(MakeGroupBlock( + block_stmt.size() == 1 ? block_stmt[0] + // NOLINTNEXTLINE(performance-move-const-arg) + : SeqStmt(std::move(block_stmt)), + annotations)); } // Filter out the producer stmts int cur_id = 0; From 919d2181d77e0dcb5f2e6e62534ed63ab17d22d7 Mon Sep 17 00:00:00 2001 From: Xuehai Pan Date: Tue, 21 Oct 2025 20:27:13 +0800 Subject: [PATCH 5/7] [Maint] restore format.sh --- format.sh | 178 ++++++++++++++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 178 insertions(+) create mode 100755 format.sh diff --git a/format.sh b/format.sh new file mode 100755 index 000000000..fe441d516 --- /dev/null +++ b/format.sh @@ -0,0 +1,178 @@ +#!/usr/bin/env bash +# Usage: +# # Do work and commit your work. +# +# # Format files that differ from origin/main. +# bash format.sh +# +# # Format all files. +# bash format.sh --all +# +# +# YAPF + Clang formatter (if installed). This script formats all changed files from the last mergebase. +# You are encouraged to run this locally before pushing changes for review. + +# Cause the script to exit if a single command fails +set -eo pipefail + +if [[ -z "${BASH_VERSION}" ]]; then + echo "Please run this script using bash." >&2 + exit 1 +fi + +# this stops git rev-parse from failing if we run this from the .git directory +builtin cd "$(dirname "${BASH_SOURCE:-$0}")" +ROOT="$(git rev-parse --show-toplevel)" +builtin cd "$ROOT" || exit 1 + +ALL_FILES='' +ONLY_CHANGED='' +FILES=() +if (($# == 0)); then + if [[ -n "$(git status --porcelain)" ]]; then + echo 'Detected uncommitted changes. Please commit or stash them before running format.sh.' >&2 + exit 1 + fi + ONLY_CHANGED='true' +else + while (($# > 0)); do + case $1 in + --files) + shift + while (($# > 0)); do + FILES+=("$1") + shift + done + ;; + --all) + ALL_FILES='true' + shift + ;; + *) + echo "Unknown argument: '$1'" >&2 + exit 1 + ;; + esac + done +fi + +MERGE_BASE="" +get_merge_base() { + UPSTREAM_REPO="https://github.com/tile-ai/tilelang" + if git ls-remote --exit-code "${UPSTREAM_REPO}" main &>/dev/null; then + # First try to use the upstream repository directly + MERGE_BASE="$(git fetch "${UPSTREAM_REPO}" main &>/dev/null && git merge-base FETCH_HEAD HEAD)" + elif git show-ref --verify --quiet refs/remotes/origin/main; then + # Fall back to origin/main if available + BASE_BRANCH="origin/main" + MERGE_BASE="$(git merge-base "${BASE_BRANCH}" HEAD)" + else + # Last resort, use local main + BASE_BRANCH="main" + MERGE_BASE="$(git merge-base "${BASE_BRANCH}" HEAD)" + fi + echo "${MERGE_BASE}" +} + +if [[ -n "${ALL_FILES}" ]]; then + echo "Checking all files..." >&2 +elif [[ -n "${ONLY_CHANGED}" ]]; then + MERGE_BASE="$(get_merge_base)" + echo "Checking changed files compared to merge base (${MERGE_BASE})..." >&2 +elif [[ "${#FILES[@]}" -gt 0 ]]; then + echo "Checking specified files: ${FILES[*]}..." >&2 +fi + +# If pre-commit is not installed, install it. +if python3 -m pre_commit --version &>/dev/null; then + python3 -m pip install pre-commit +fi + +if [[ ! -f "${ROOT}/.git/hooks/pre-commit" ]]; then + echo "Installing and initializing pre-commit hooks..." + python3 -m pre_commit install --install-hooks +fi + +echo 'tile-lang pre-commit: Check Start' + +if [[ -n "${ALL_FILES}" ]]; then + python3 -m pre_commit run --all-files +elif [[ -n "${ONLY_CHANGED}" ]]; then + python3 -m pre_commit run --from-ref "${MERGE_BASE}" --to-ref HEAD +elif [[ "${#FILES[@]}" -gt 0 ]]; then + python3 -m pre_commit run --files "${FILES[@]}" +fi + +echo 'tile-lang pre-commit: Done' + +echo 'tile-lang clang-tidy: Check Start' +# If clang-tidy is available, run it; otherwise, skip +if command -v run-clang-tidy &>/dev/null; then + # Check if clang-tidy is available + if ! command -v clang-tidy &>/dev/null; then + python3 -m pip install --upgrade --requirements "${ROOT}/requirements-lint.txt" + fi + # Get clang-tidy version + CLANG_TIDY_VERSION="$(clang-tidy --version | head -n1 | awk '{print $4}')" + echo "Using clang-tidy version: ${CLANG_TIDY_VERSION}" + + # Check if build directory exists + if [[ ! -d "${ROOT}/build" ]]; then + echo "Build directory not found. Skipping clang-tidy checks." + else + # Run clang-tidy on specified files + clang_tidy_files() { + run-clang-tidy -j 64 "$@" -p build + } + + # Run clang-tidy on all C/C++ source files + clang_tidy_all() { + run-clang-tidy -j 64 src/*.cc -p build + } + + # Run clang-tidy on changed C/C++ files relative to main + clang_tidy_changed() { + # Get changed C/C++ files + CHANGED_FILES="$(git diff --name-only --diff-filter=ACM "${MERGE_BASE}" -- '*.c' '*.cc' '*.cpp' '*.h' '*.hpp' 2>/dev/null || true)" + + if [[ -n "${CHANGED_FILES}" ]]; then + echo "Running clang-tidy on changed files:" + echo "${CHANGED_FILES}" + # Convert newline-separated files to space-separated and run clang-tidy once + CHANGED_FILES_SPACE="$(echo "${CHANGED_FILES}" | tr '\n' ' ')" + run-clang-tidy -j 64 ${CHANGED_FILES_SPACE} -p build -fix + else + echo "No C/C++ files changed. Skipping clang-tidy." + fi + } + + if [[ -n "${ALL_FILES}" ]]; then + # If --all is given, run clang-tidy on all source files + clang_tidy_all + elif [[ -n "${ONLY_CHANGED}" ]]; then + # Otherwise, run clang-tidy only on changed C/C++ files + clang_tidy_changed + elif [[ "${#FILES[@]}" -gt 0 ]]; then + # If --files is given, run clang-tidy only on the provided files + clang_tidy_files "${FILES[@]}" + fi + fi + +else + echo "run-clang-tidy not found. Skipping clang-tidy checks." + echo "To install clang-tidy tools, you may need to install clang-tidy and run-clang-tidy." +fi +echo 'tile-lang clang-tidy: Done' + +# Check if there are any uncommitted changes after all formatting steps. +# If there are, ask the user to review and stage them. +if ! git diff --quiet &>/dev/null; then + echo 'Reformatted files. Please review and stage the changes.' + echo 'Changes not staged for commit:' + echo + git --no-pager diff --name-only + + exit 1 +fi + +echo 'tile-lang: All checks passed' From 264188f5ceb04878bf1773cff8498dd298a636c2 Mon Sep 17 00:00:00 2001 From: Xuehai Pan Date: Tue, 21 Oct 2025 22:16:03 +0800 Subject: [PATCH 6/7] [CI] pre-commit autoupdate --- .pre-commit-config.yaml | 2 +- requirements-lint.txt | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/.pre-commit-config.yaml b/.pre-commit-config.yaml index 314cd9eff..284be3d84 100644 --- a/.pre-commit-config.yaml +++ b/.pre-commit-config.yaml @@ -41,7 +41,7 @@ repos: ^.+\.json$ ) - repo: https://github.com/astral-sh/ruff-pre-commit - rev: v0.14.0 # sync with requirements-lint.txt + rev: v0.14.1 # sync with requirements-lint.txt hooks: - id: ruff-check args: [--fix, --exit-non-zero-on-fix] diff --git a/requirements-lint.txt b/requirements-lint.txt index c2db20541..d604b1ec2 100644 --- a/requirements-lint.txt +++ b/requirements-lint.txt @@ -3,5 +3,5 @@ pre-commit clang-format==21.1.2 clang-tidy==21.1.1 codespell[toml]==2.4.1 -ruff==0.14.0 +ruff==0.14.1 yapf==0.43.0 From f969484c8234ad76ecd839b3bb2cce3cb43c364b Mon Sep 17 00:00:00 2001 From: Xuehai Pan Date: Tue, 21 Oct 2025 22:18:15 +0800 Subject: [PATCH 7/7] [Minor] fix `command -v` usage --- format.sh | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/format.sh b/format.sh index fe441d516..bf67d7e0a 100755 --- a/format.sh +++ b/format.sh @@ -84,7 +84,7 @@ elif [[ "${#FILES[@]}" -gt 0 ]]; then fi # If pre-commit is not installed, install it. -if python3 -m pre_commit --version &>/dev/null; then +if ! python3 -m pre_commit --version &>/dev/null; then python3 -m pip install pre-commit fi @@ -107,9 +107,9 @@ echo 'tile-lang pre-commit: Done' echo 'tile-lang clang-tidy: Check Start' # If clang-tidy is available, run it; otherwise, skip -if command -v run-clang-tidy &>/dev/null; then +if [[ -x "$(command -v run-clang-tidy)" ]]; then # Check if clang-tidy is available - if ! command -v clang-tidy &>/dev/null; then + if [[ ! -x "$(command -v clang-tidy)" ]]; then python3 -m pip install --upgrade --requirements "${ROOT}/requirements-lint.txt" fi # Get clang-tidy version