Refactor buffer reference handling in CUDA code generation and update test execution in tilelang. Ensure default annotations for unrolling are set correctly in TIR IR module.

Author: Hamerlate

src/target/codegen_cuda.cc

@@ -1290,11 +1290,7 @@ std::string CodeGenTileLangCUDA::GetBufferRef(DataType t,
     scope = GetPtrStorageScope(buffer->data);
   }
   if (scope == "local.var" || scope.find("local.descriptor") == 0) {
-    if (scope == "local.descriptor.tcgen05_instr") {
-      os << vid << ".desc_";
-    } else {
-      os << vid;
-    }
+    os << vid;
     return os.str();
   }
   std::string index_str = PrintExpr(index);

testing/python/language/test_tilelang_language_get_warp_info.py

@@ -208,5 +208,4 @@ def test_shuffle_elect_block_leader():
 
 
 if __name__ == "__main__":
-    # tilelang.testing.main()
-    test_get_lane_idx_custom()
+    tilelang.testing.main()

tilelang/language/tir/ir.py

@@ -104,6 +104,13 @@ def unroll(start: PrimExpr,
     res : frame.ForFrame
         The ForFrame.
     """
+    # Ensure annotations has {"pragma_unroll_explicit": True} by default
+    if annotations is None:
+        annotations = {"pragma_unroll_explicit": False}
+    else:
+        # Add "pragma_unroll_explicit": True if not already present
+        annotations = dict(annotations)
+        annotations.setdefault("pragma_unroll_explicit", False)
     return _ir.unroll(start=start, stop=stop, annotations=annotations)