Skip to content

Commit f6b724d

Browse files
RachmaninoRachmanino
committed
lint
1 parent 495144a commit f6b724d

3 files changed

Lines changed: 9 additions & 15 deletions

File tree

‎src/transform/lower_device_kernel_launch.cc‎

Lines changed: 4 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -84,8 +84,6 @@ class DeviceInfoCollector : public StmtVisitor {
8484
// Cluster dims are promoted to a PrimFunc attr by LowerOpaqueBlock.
8585
if (auto opt = func->GetAttr<Array<Integer>>("cluster_dims")) {
8686
collector.info_.cluster_dims = opt.value();
87-
LOG(INFO) << "cluster_dims found";
88-
LOG(INFO) << "cluster_dims: " << opt.value();
8987
}
9088

9189
// The dynamic shared memory is required to be the last of the
@@ -108,8 +106,8 @@ class DeviceInfoCollector : public StmtVisitor {
108106
// can read them from packed-function args at launch time.
109107
if (collector.info_.cluster_dims.defined()) {
110108
auto dims = collector.info_.cluster_dims.value();
111-
Array<PrimExpr> new_launch_args = {
112-
PrimExpr(dims[0]), PrimExpr(dims[1]), PrimExpr(dims[2])};
109+
Array<PrimExpr> new_launch_args = {PrimExpr(dims[0]), PrimExpr(dims[1]),
110+
PrimExpr(dims[2])};
113111
for (auto arg : collector.info_.launch_args)
114112
new_launch_args.push_back(arg);
115113
collector.info_.launch_args = new_launch_args;
@@ -292,8 +290,8 @@ class DeviceKernelMutator : public StmtExprMutator {
292290
info.dyn_shmem_size.value());
293291
}
294292
if (info.cluster_dims.defined()) {
295-
LOG(INFO) << "cluster_dims found";
296-
func = WithAttr(std::move(func), "cluster_dims", info.cluster_dims.value());
293+
func =
294+
WithAttr(std::move(func), "cluster_dims", info.cluster_dims.value());
297295
}
298296
return func;
299297
}

‎src/transform/lower_opaque_block.cc‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -302,7 +302,8 @@ class OpaqueBlockLower : public StmtExprMutator {
302302
/*! \brief Local var initializers collected from block annotations. */
303303
Map<Var, PrimExpr> local_var_init_map_;
304304

305-
/*! \brief Cluster dims collected from tilelang.cluster_dims block annotation. */
305+
/*! \brief Cluster dims collected from tilelang.cluster_dims block annotation.
306+
*/
306307
Optional<Array<Integer>> cluster_dims_{std::nullopt};
307308
};
308309

‎tilelang/jit/adapter/wrapper.py‎

Lines changed: 3 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -358,7 +358,7 @@ def create_dispatch_func(self, code, function_informations):
358358
assert len(function_params) == len(args_list), (
359359
f"Function {function_name} has {len(function_params)} parameters, but {len(args_list)} arguments"
360360
)
361-
361+
362362
call_args = ", ".join(args_list)
363363
if self.cluster_dims[function_name] is None:
364364
kernel_code = KERNEL_LAUNCH_FUNC_CODE.format(
@@ -370,17 +370,12 @@ def create_dispatch_func(self, code, function_informations):
370370
)
371371
else:
372372
kernel_code = KERNEL_CLUSTER_LAUNCH_FUNC_CODE.format(
373-
grid_str,
374-
block_str,
375-
smem_str,
376-
call_args,
377-
function_name,
378-
*self.cluster_dims[function_name]
373+
grid_str, block_str, smem_str, call_args, function_name, *self.cluster_dims[function_name]
379374
)
380375

381376
kernel_launch_code += kernel_code
382377
kernel_launch_code += f'\tTILELANG_CHECK_LAST_ERROR("{function_name}");\n'
383-
378+
384379
if has_l2_persistent_map:
385380
kernel_launch_code += L2_PERSISTENT_MAP_RESET_HANDLE
386381

0 commit comments

Comments
 (0)