Skip to content

Commit cdc5d8d

Browse files
LeiWang1999LeiWang1999
authored
[Lint] Introduce clang-tidy into format.sh (#777)
* [Refactor] Update Clang-Tidy Checks and Improve Code Consistency - Enhanced .clang-tidy configuration by adding specific checks for better bug detection and performance optimization. - Refactored function signatures across multiple files to use `const` references for parameters, improving performance and code clarity. - Updated various methods to ensure consistent handling of parameters, particularly in `AddPredicate`, `Substitute`, and `PlanLoopPartition` functions. - Improved readability by replacing size checks with `empty()` method calls in several locations, ensuring clearer intent in the code. - General code cleanup and adherence to best practices for better maintainability. * [Refactor] Enhance Code Consistency and Clang-Tidy Configuration - Updated .clang-tidy configuration to include additional checks for improved code quality and performance. - Refactored function signatures across multiple files to use `const` references, enhancing performance and clarity. - Replaced size checks with `empty()` method calls in various locations for clearer intent. - Improved handling of parameters in several functions, ensuring consistent usage of `std::move` where applicable. - General code cleanup to adhere to best practices and improve maintainability. * [Refactor] Integrate Clang-Tidy Checks and Enhance Code Consistency - Added clang-tidy checks to the format script for improved code quality assurance. - Refactored function signatures across multiple files to consistently use `const` references, enhancing performance and clarity. - Updated the requirements-lint.txt file to include clang-tidy as a dependency. - General code cleanup to adhere to best practices and improve maintainability. * [CI] Update AMD CI Workflow to Include Build Directory Creation - Added steps to create a build directory and configure CMake with ROCm support during the format check process. - Ensured cleanup of the build directory after the format check to maintain a clean workspace. * [Refactor] Remove Unused Member Variables in AtomicAddNode and CopyNode - Removed the `args_` member variable from both `AtomicAddNode` and `CopyNode` classes to streamline the code and eliminate unnecessary data members. - This change enhances code clarity and maintainability by focusing on relevant attributes for each class. * [Refactor] Update Clang-Tidy Integration and Code Improvements - Modified the format script to include the `-fix` option in the clang-tidy command for automatic code fixes. - Refactored the `AtomicAddVectorizePlanner` class to improve variable handling and consistency, including changes to member variable types and function signatures. - Enhanced code clarity by removing unnecessary `std::move` calls and ensuring consistent usage of types across the class. - General code cleanup to adhere to best practices and improve maintainability. * [Refactor] Improve Parameter Handling and Consistency in AtomicAddVectorize - Updated function signatures in `AtomicAddVectorizePlanResult` and `AtomicAddVectorizeRewriter` to use `const` references and `std::move` for better performance and clarity. - Enhanced the `UpdateVectorSize` method to accept `const Array<PrimExpr>&` for improved efficiency. - General code cleanup to maintain consistency and adhere to best practices. * [CI] Add Git Submodule Initialization to CI Workflow - Included a step to initialize and update git submodules recursively in the CI workflow. - This change ensures that all necessary submodules are available during the format check process, improving build reliability. * [CI] Add Git Submodule Update Step to Format Check - Included a command to initialize and update git submodules recursively in the CI workflow during the format check process. - This enhancement ensures that all required submodules are available, contributing to improved build reliability. * [Refactor] Update Function Signatures in AtomicAddVectorize - Modified the `VectorizeAtomicAdd` function signature to use `const` references for `thread_var` and `thread_bounds`, enhancing performance and code clarity. - This change aligns with previous refactoring efforts to improve parameter handling and consistency across the codebase.
1 parent 471cc7f commit cdc5d8d

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

63 files changed

+604
-497
lines changed

‎.clang-tidy‎

Lines changed: 42 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -1,10 +1,47 @@
11
Checks: >
2+
# 1. Retained categories: easier to find bugs/performance issues
23
clang-analyzer-*,
3-
cppcoreguidelines-*,
4-
modernize-*,
4+
cppcoreguidelines-pro-type-static-cast-downcast,
5+
cppcoreguidelines-pro-type-member-init,
6+
cppcoreguidelines-pro-bounds-array-to-pointer-decay,
7+
cppcoreguidelines-pro-bounds-pointer-arithmetic,
8+
cppcoreguidelines-slicing,
9+
cppcoreguidelines-narrowing-conversions,
510
performance-*,
6-
readability-*,
7-
-readability-identifier-length
11+
12+
# 2. Readability: only keep useful rules
13+
readability-braces-around-statements,
14+
readability-container-size-empty,
15+
readability-delete-null-pointer,
16+
readability-redundant-member-init,
17+
readability-redundant-smartptr-get,
18+
readability-redundant-string-cstr,
19+
20+
# 3. Disable all intrusive/style-breaking rules
21+
-readability-identifier-length,
22+
-readability-avoid-const-params-in-decls,
23+
-readability-else-after-return,
24+
-cppcoreguidelines-avoid-magic-numbers,
25+
-modernize-use-trailing-return-type,
26+
-modernize-use-nodiscard,
27+
-modernize-use-auto,
28+
-modernize-pass-by-value,
29+
-modernize-return-braced-init-list,
30+
-modernize-use-default-member-init,
31+
-modernize-loop-convert,
32+
-modernize-concat-nested-namespaces,
33+
-llvm-include-order,
34+
-bugprone-unused-return-value,
35+
-clang-diagnostic-unused-result,
36+
-cppcoreguidelines-special-member-functions,
37+
-performance-noexcept-move-constructor,
38+
-cppcoreguidelines-narrowing-conversions,
39+
-clang-diagnostic-error,
40+
-cppcoreguidelines-pro-type-member-init,
41+
-clang-analyzer-optin.cplusplus.UninitializedObject,
42+
-cppcoreguidelines-pro-type-static-cast-downcast,
43+
-performance-unnecessary-value-param,
44+
845
WarningsAsErrors: '*'
946

10-
HeaderFilterRegex: '^(?!.*(3rdparty|build)).*$'
47+
HeaderFilterRegex: '^(?!.*(3rdparty|build)).*$'

‎.github/workflows/amd_ci.yml‎

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -48,6 +48,9 @@ jobs:
4848
- name: Run format check
4949
run: |
5050
source "${{ runner.tool_cache }}/${{ env.VENV_DIR }}/bin/activate"
51+
git submodule update --init --recursive
52+
mkdir -p build
53+
cd build; cmake .. -DCMAKE_EXPORT_COMPILE_COMMANDS=ON -DUSE_ROCM=ON; cd ..
5154
if ! output=$(./format.sh 2>&1); then
5255
echo "------------------------------------"
5356
echo "message:"
@@ -56,6 +59,7 @@ jobs:
5659
echo "------------------------------------"
5760
exit 1
5861
fi
62+
rm -rf build
5963
6064
- name: Commit and Push Changes
6165
uses: stefanzweifel/git-auto-commit-action@v5

‎.github/workflows/ci.yml‎

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -47,6 +47,10 @@ jobs:
4747
- name: Run format check
4848
run: |
4949
source "${{ runner.tool_cache }}/${{ env.VENV_DIR }}/bin/activate"
50+
git submodule update --init --recursive
51+
mkdir -p build
52+
# run cmake to create the build directory with compile_commands.json
53+
cd build; cmake .. -DCMAKE_EXPORT_COMPILE_COMMANDS=ON -DUSE_CUDA=ON; cd ..
5054
if ! output=$(./format.sh 2>&1); then
5155
echo "------------------------------------"
5256
echo "message:"
@@ -55,6 +59,7 @@ jobs:
5559
echo "------------------------------------"
5660
exit 1
5761
fi
62+
rm -rf build
5863
5964
- name: Commit and Push Changes
6065
uses: stefanzweifel/git-auto-commit-action@v5

‎format.sh‎

Lines changed: 67 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -249,6 +249,73 @@ else
249249
fi
250250
echo 'tile-lang clang-format: Done'
251251

252+
echo 'tile-lang clang-tidy: Check Start'
253+
# If clang-tidy is available, run it; otherwise, skip
254+
if command -v run-clang-tidy &>/dev/null; then
255+
# Check if clang-tidy is available
256+
if ! command -v clang-tidy &>/dev/null; then
257+
echo "clang-tidy not found. Skipping clang-tidy checks."
258+
else
259+
# Get clang-tidy version
260+
CLANG_TIDY_VERSION=$(clang-tidy --version | head -n1 | awk '{print $4}')
261+
echo "Using clang-tidy version: $CLANG_TIDY_VERSION"
262+
263+
# Check if build directory exists
264+
if [ ! -d "build" ]; then
265+
echo "Build directory not found. Skipping clang-tidy checks."
266+
else
267+
# Run clang-tidy on specified files
268+
clang_tidy_files() {
269+
run-clang-tidy -j 64 "$@" -p build
270+
}
271+
272+
# Run clang-tidy on all C/C++ source files
273+
clang_tidy_all() {
274+
run-clang-tidy -j 64 src/*.cc -p build
275+
}
276+
277+
# Run clang-tidy on changed C/C++ files relative to main
278+
clang_tidy_changed() {
279+
if git show-ref --verify --quiet refs/remotes/origin/main; then
280+
BASE_BRANCH="origin/main"
281+
else
282+
BASE_BRANCH="main"
283+
fi
284+
285+
MERGEBASE="$(git merge-base $BASE_BRANCH HEAD)"
286+
287+
# Get changed C/C++ files
288+
CHANGED_FILES=$(git diff --name-only --diff-filter=ACM "$MERGEBASE" -- '*.c' '*.cc' '*.cpp' '*.h' '*.hpp' 2>/dev/null || true)
289+
290+
if [ -n "$CHANGED_FILES" ]; then
291+
echo "Running clang-tidy on changed files:"
292+
echo "$CHANGED_FILES"
293+
# Convert newline-separated files to space-separated and run clang-tidy once
294+
CHANGED_FILES_SPACE=$(echo "$CHANGED_FILES" | tr '\n' ' ')
295+
run-clang-tidy -j 64 $CHANGED_FILES_SPACE -p build -fix
296+
else
297+
echo "No C/C++ files changed. Skipping clang-tidy."
298+
fi
299+
}
300+
301+
if [[ "$1" == '--files' ]]; then
302+
# If --files is given, run clang-tidy only on the provided files
303+
clang_tidy_files "${@:2}"
304+
elif [[ "$1" == '--all' ]]; then
305+
# If --all is given, run clang-tidy on all source files
306+
clang_tidy_all
307+
else
308+
# Otherwise, run clang-tidy only on changed C/C++ files
309+
clang_tidy_changed
310+
fi
311+
fi
312+
fi
313+
else
314+
echo "run-clang-tidy not found. Skipping clang-tidy checks."
315+
echo "To install clang-tidy tools, you may need to install clang-tidy and run-clang-tidy."
316+
fi
317+
echo 'tile-lang clang-tidy: Done'
318+
252319
# Check if there are any uncommitted changes after all formatting steps.
253320
# If there are, ask the user to review and stage them.
254321
if ! git diff --quiet &>/dev/null; then

‎requirements-lint.txt‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -5,3 +5,4 @@ tomli==2.0.1
55
ruff==0.6.5
66
codespell==2.3.0
77
clang-format==15.0.7
8+
clang-tidy==18.1.8

‎src/ir.cc‎

Lines changed: 39 additions & 36 deletions
Original file line numberDiff line numberDiff line change
@@ -11,6 +11,8 @@
1111
#include <tvm/ffi/reflection/registry.h>
1212
#include <tvm/script/ir_builder/tir/ir.h>
1313

14+
#include <utility>
15+
1416
namespace tvm {
1517
namespace tl {
1618

@@ -19,8 +21,8 @@ using namespace script::ir_builder::tir;
1921
static Var CreateEnvThread(String name, String thread_tag, DataType dtype) {
2022
using namespace tvm::tir;
2123
using namespace tvm::script::ir_builder;
22-
IterVar iter_var(Range{nullptr}, Var(name, dtype),
23-
tvm::tir::IterVarType::kThreadIndex, thread_tag);
24+
IterVar iter_var(Range{nullptr}, Var(std::move(name), dtype),
25+
tvm::tir::IterVarType::kThreadIndex, std::move(thread_tag));
2426
Var var = iter_var->var;
2527
if (Optional<PrimFuncFrame> opt_frame =
2628
IRBuilder::Current()->FindFrame<PrimFuncFrame>()) {
@@ -31,24 +33,24 @@ static Var CreateEnvThread(String name, String thread_tag, DataType dtype) {
3133
return var;
3234
}
3335

34-
static ForFrame MakeIterVarFrame(std::string name, PrimExpr dom) {
36+
static ForFrame MakeIterVarFrame(const std::string &name, const PrimExpr &dom) {
3537
using namespace tvm::tir;
3638
Var var = Var(name, dom->dtype);
3739
// Create a frame that represents a loop over the given domain.
3840
ObjectPtr<ForFrameNode> n = make_object<ForFrameNode>();
3941
n->vars.push_back(var);
4042
n->doms.push_back(Range(0, dom));
41-
n->f_make_for_loop = [](Array<Var> vars, Array<Range> doms,
42-
Stmt body) -> Stmt {
43+
n->f_make_for_loop = [](const Array<Var> &vars, const Array<Range> &doms,
44+
const Stmt &body) -> Stmt {
4345
ICHECK_EQ(vars.size(), 1);
4446
ICHECK_EQ(doms.size(), 1);
4547
return For(vars[0], doms[0]->min, doms[0]->extent, ForKind::kSerial, body);
4648
};
4749
return ForFrame(n);
4850
}
4951

50-
ForFrame ParallelFor(Array<PrimExpr> extents,
51-
Map<String, ObjectRef> annotations) {
52+
ForFrame ParallelFor(const Array<PrimExpr> &extents,
53+
const Map<String, ObjectRef> &annotations) {
5254
using namespace tvm::tir;
5355
ObjectPtr<ForFrameNode> n = make_object<ForFrameNode>();
5456
n->vars.reserve(extents.size());
@@ -58,59 +60,59 @@ ForFrame ParallelFor(Array<PrimExpr> extents,
5860
n->vars.push_back(Var("v", extent.dtype()));
5961
n->doms.push_back(Range(make_const(dtype, 0), extent));
6062
}
61-
n->f_make_for_loop = [annotations](Array<Var> vars, Array<Range> doms,
63+
n->f_make_for_loop = [annotations](const Array<Var> &vars,
64+
const Array<Range> &doms,
6265
Stmt body) -> Stmt {
6366
ICHECK_EQ(vars.size(), doms.size());
6467
int n = vars.size();
6568
for (int i = n - 1; i >= 0; --i) {
6669
Range dom = doms[i];
6770
Var var = vars[i];
68-
body =
69-
For(var, dom->min, dom->extent, ForKind::kParallel, std::move(body),
70-
/*thread_binding=*/std::nullopt, /*annotations=*/annotations);
71+
body = For(var, dom->min, dom->extent, ForKind::kParallel, body,
72+
/*thread_binding=*/std::nullopt, /*annotations=*/annotations);
7173
}
7274
return body;
7375
};
7476
return ForFrame(n);
7577
}
7678

77-
ForFrame PipelinedFor(PrimExpr start, PrimExpr stop, int num_stages,
78-
Array<PrimExpr> order, Array<PrimExpr> stages,
79-
Array<Array<PrimExpr>> sync,
80-
Array<Array<PrimExpr>> groups) {
79+
ForFrame PipelinedFor(PrimExpr start, const PrimExpr &stop, int num_stages,
80+
const Array<PrimExpr> &order,
81+
const Array<PrimExpr> &stages,
82+
const Array<Array<PrimExpr>> &sync,
83+
const Array<Array<PrimExpr>> &groups) {
8184
using namespace tvm::tir;
8285
ObjectPtr<ForFrameNode> n = make_object<ForFrameNode>();
8386
DataType dtype = stop.dtype();
8487
n->vars.push_back(Var("v", dtype));
85-
n->doms.push_back(Range(start, stop));
86-
n->f_make_for_loop = [=](Array<Var> vars, Array<Range> doms,
88+
n->doms.push_back(Range(std::move(start), stop));
89+
n->f_make_for_loop = [=](const Array<Var> &vars, const Array<Range> &doms,
8790
Stmt body) -> Stmt {
8891
ICHECK_EQ(vars.size(), doms.size());
8992
int n = vars.size();
9093
ICHECK(n == 1);
9194
Map<String, ObjectRef> anno;
9295
if (num_stages > 0)
9396
anno.Set("num_stages", PrimExpr(num_stages));
94-
if (order.size() > 0)
97+
if (!order.empty())
9598
anno.Set("tl_pipeline_order", order);
96-
if (stages.size() > 0)
99+
if (!stages.empty())
97100
anno.Set("tl_pipeline_stage", stages);
98-
if (sync.size() > 0)
101+
if (!sync.empty())
99102
anno.Set("tl_pipeline_sync", sync);
100-
if (groups.size() > 0)
103+
if (!groups.empty())
101104
anno.Set("tl_pipeline_group", groups);
102-
body = For(vars[0], doms[0]->min, doms[0]->extent, ForKind::kSerial,
103-
std::move(body),
105+
body = For(vars[0], doms[0]->min, doms[0]->extent, ForKind::kSerial, body,
104106
/*thread_binding=*/std::nullopt, /*annotations=*/anno);
105107
return body;
106108
};
107109
return ForFrame(n);
108110
}
109111

110-
ForFrame PersistentFor(Array<PrimExpr> domain, PrimExpr wave_size,
111-
PrimExpr index, PrimExpr group_size) {
112+
ForFrame PersistentFor(const Array<PrimExpr> &domain, const PrimExpr &wave_size,
113+
const PrimExpr &index, PrimExpr group_size) {
112114
using namespace tvm::tir;
113-
ICHECK(domain.size() > 0);
115+
ICHECK(!domain.empty());
114116
ObjectPtr<ForFrameNode> n = make_object<ForFrameNode>();
115117
n->vars.reserve(domain.size());
116118
n->doms.reserve(domain.size());
@@ -139,8 +141,8 @@ ForFrame PersistentFor(Array<PrimExpr> domain, PrimExpr wave_size,
139141
}
140142
grouped_domain.push_back(group_size);
141143

142-
n->f_make_for_loop = [=](Array<Var> vars, Array<Range> doms,
143-
Stmt body) -> Stmt {
144+
n->f_make_for_loop = [=](const Array<Var> &vars, const Array<Range> &doms,
145+
const Stmt &body) -> Stmt {
144146
ICHECK_EQ(vars.size(), doms.size());
145147
Map<String, ObjectRef> anno;
146148
Array<PrimExpr> idxs(grouped_domain.size(), PrimExpr());
@@ -220,9 +222,9 @@ class KernelLaunchFrame : public TIRFrame {
220222
KernelLaunchFrameNode);
221223
};
222224

223-
KernelLaunchFrame KernelLaunch(Array<PrimExpr> grid_size,
224-
Optional<Array<PrimExpr>> block_size_opt,
225-
Map<String, ffi::Any> attrs) {
225+
KernelLaunchFrame KernelLaunch(const Array<PrimExpr> &grid_size,
226+
const Optional<Array<PrimExpr>> &block_size_opt,
227+
const Map<String, ffi::Any> &attrs) {
226228
ObjectPtr<KernelLaunchFrameNode> n = make_object<KernelLaunchFrameNode>();
227229

228230
// If the kernel is a CPU kernel, we don't need to launch any threads.
@@ -234,7 +236,7 @@ KernelLaunchFrame KernelLaunch(Array<PrimExpr> grid_size,
234236
if (is_cpu_kernel_frame) {
235237
// Launch CPU Kernel
236238
ICHECK(grid_size.size() >= 0);
237-
ICHECK(block_size.size() == 0) << "CPU kernel cannot have block size";
239+
ICHECK(block_size.empty()) << "CPU kernel cannot have block size";
238240
ICHECK(attrs.defined());
239241
// create grid loop var
240242
for (int i = 0; i < grid_size.size(); i++) {
@@ -244,7 +246,7 @@ KernelLaunchFrame KernelLaunch(Array<PrimExpr> grid_size,
244246
} else {
245247
// Launch GPU Kernel
246248
ICHECK(grid_size.size() <= 3);
247-
if (grid_size.size() > 0)
249+
if (!grid_size.empty())
248250
n->frames.push_back(LaunchThread(
249251
CreateEnvThread("bx", "blockIdx.x", grid_size[0].dtype()),
250252
grid_size[0]));
@@ -258,7 +260,7 @@ KernelLaunchFrame KernelLaunch(Array<PrimExpr> grid_size,
258260
grid_size[2]));
259261
if (block_size.defined()) {
260262
ICHECK(block_size.size() <= 3);
261-
if (block_size.size() > 0) {
263+
if (!block_size.empty()) {
262264
n->frames.push_back(LaunchThread(
263265
CreateEnvThread("tx", "threadIdx.x", block_size[0].dtype()),
264266
block_size[0]));
@@ -333,12 +335,13 @@ class WarpSpecializeFrame : public TIRFrame {
333335
WarpSpecializeFrameNode);
334336
};
335337

336-
WarpSpecializeFrame WarpSpecialize(Array<IntImm> warp_group_ids,
337-
PrimExpr thread_idx,
338+
WarpSpecializeFrame WarpSpecialize(const Array<IntImm> &warp_group_ids,
339+
const PrimExpr &thread_idx,
338340
int warp_group_size = 128) {
339341
ObjectPtr<WarpSpecializeFrameNode> n = make_object<WarpSpecializeFrameNode>();
340342
PrimExpr condition;
341343
std::vector<int> warp_groups;
344+
warp_groups.reserve(warp_group_ids.size());
342345
for (int i = 0; i < warp_group_ids.size(); i++) {
343346
warp_groups.push_back(Downcast<IntImm>(warp_group_ids[i])->value);
344347
}

‎src/op/atomic_add.h‎

Lines changed: 0 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -90,8 +90,6 @@ using namespace tir;
9090

9191
class AtomicAddNode : public TileOperatorNode {
9292
public:
93-
Array<PrimExpr> args_;
94-
9593
Buffer src, dst;
9694
Array<Range> src_range, dst_range;
9795
IntImm coalesced_width;

0 commit comments

Comments
 (0)