Merge branch 'main' into frontend-v2

Author: kurisu6912

.clang-tidy

@@ -4,7 +4,7 @@ ExtraArgs: ['-v']
 FormatStyle: file
 UseColor: true
 WarningsAsErrors: '*'
-ExcludeHeaderFilterRegex: '^(3rdparty|tvm)/.*$'
+HeaderFilterRegex: '^(?!.*(?:/|^)(3rdparty|tvm)/).*'
 
 # NOTE: there must be no spaces before the '-', so put the comma last.
 Checks: >-

.github/workflows/dist.yml

@@ -94,7 +94,7 @@ jobs:
       - name: Upload wheels
         # Not PR to save artifact storage, as wheels are only needed for releases.
         if: github.event_name != 'pull_request'
-        uses: actions/upload-artifact@v4
+        uses: actions/upload-artifact@v5
         with:
           name: wheels-${{ matrix.python-version }}-${{ runner.os }}-${{ runner.arch }}-${{ matrix.target.toolkit }}
           path: wheelhouse/*.whl
@@ -109,7 +109,7 @@ jobs:
     timeout-minutes: 15
     steps:
       - name: Download built wheels
-        uses: actions/download-artifact@v5
+        uses: actions/download-artifact@v6
         with:
           pattern: wheels-*
           path: dist
@@ -119,7 +119,7 @@ jobs:
         run: ls -lh dist/*
 
       - name: Upload artifacts
-        uses: actions/upload-artifact@v4
+        uses: actions/upload-artifact@v5
         with:
           name: artifacts
           path: dist/*

src/op/builtin.h

@@ -503,6 +503,7 @@ TVM_DLL const Op &initialize_descriptor();
  *  This op is used to represent a descriptor start address setting operation in
  * tilelang.
  */
+
 TVM_DLL const Op &increase_descriptor_offset();
 /*!
  * \brief tilelang intrinsic for element-wise atomic addition.

src/tl_templates/cuda/atomic.h

@@ -133,6 +133,10 @@ TL_DEVICE void fence_proxy_async() {
   asm volatile("fence.proxy.async.shared::cta;" : :);
 }
 
+TL_DEVICE void fence_barrier_init() {
+  asm volatile("fence.mbarrier_init.release.cluster;" : :);
+}
+
 // Indicate arrival of warp issuing TMA_STORE
 TL_DEVICE void tma_store_arrive() {
   asm volatile("cp.async.bulk.commit_group;");

src/tl_templates/cuda/barrier.h

@@ -257,3 +257,12 @@ __device__ void debug_print_buffer_value<int16_t>(const char *msg,
          msg, blockIdx.x, blockIdx.y, blockIdx.z, threadIdx.x, threadIdx.y,
          threadIdx.z, buf_name, index, (int32_t)var);
 }
+
+TL_DEVICE void device_assert(bool cond) { assert(cond); }
+
+TL_DEVICE void device_assert_with_msg(bool cond, const char *msg) {
+  if (!cond) {
+    printf("Device assert failed: %s\n", msg);
+    assert(0);
+  }
+}