From f9a22c74ceb90057fe9007c393544da6b708a473 Mon Sep 17 00:00:00 2001
From: LeiWang1999 <leiwang1999@outlook.com>
Date: Sun, 27 Apr 2025 12:34:47 +0800
Subject: [PATCH] [Refactor] Update KernelLaunch to clarify CPU and GPU kernel
 launch logic

* Added comments to distinguish between CPU and GPU kernel launch sections for better code readability.
* Changed the creation of empty blocks to use a consistent "root" identifier, enhancing clarity in frame management.
---
 src/ir.cc | 8 +++-----
 1 file changed, 3 insertions(+), 5 deletions(-)

diff --git a/src/ir.cc b/src/ir.cc
index 02378505d..f4ad452fb 100644
--- a/src/ir.cc
+++ b/src/ir.cc
@@ -162,6 +162,7 @@ KernelLaunchFrame KernelLaunch(Array<PrimExpr> grid_size,
       attrs.defined() && attrs.count(tilelang_is_cpu_kernel_frame);
 
   if (is_cpu_kernel_frame) {
+    // Launch CPU Kernel
     ICHECK(grid_size.size() >= 0);
     ICHECK(block_size.size() == 0) << "CPU kernel cannot have block size";
     ICHECK(attrs.defined());
@@ -170,7 +171,6 @@ KernelLaunchFrame KernelLaunch(Array<PrimExpr> grid_size,
       n->frames.push_back(
           MakeIterVarFrame("block_var_" + std::to_string(i), grid_size[i]));
     }
-    // Launch CPU Kernel
   } else {
     // Launch GPU Kernel
     ICHECK(grid_size.size() <= 3);
@@ -203,17 +203,15 @@ KernelLaunchFrame KernelLaunch(Array<PrimExpr> grid_size,
             CreateEnvThread("tz", "threadIdx.z", block_size[2].dtype()),
             block_size[2]));
       }
-    } else {
-      n->frames.push_back(Block(""));
     }
   }
 
   if (attrs.defined()) {
-    auto empty_block = Block("");
+    auto empty_block = Block("root");
     empty_block->annotations = attrs;
     n->frames.push_back(empty_block);
   } else {
-    n->frames.push_back(Block(""));
+    n->frames.push_back(Block("root"));
   }
 
   return KernelLaunchFrame(n);