cute（10）hopper之mbarrier

鱿鱼圈 Lv4
  2026-06-17 02:13:32 2026-06-17 02:13:32 创建   2026-06-14 22:17:19 2026-06-14 22:17:19 更新      

前置阅读

10 MBarrier 基础

对应代码：10_mbarrier.cu 需要 GPU（SM90 / Hopper）。

核心概念

MBarrier（Memory Barrier）是 Hopper 架构引入的硬件同步原语，存储在 shared memory 中的 64-bit 数据结构。与 __syncthreads 相比，它支持：

  • 部分线程同步：只让指定线程参与，不是整个 CTA
  • Transaction Bytes：与 TMA 异步搬运配合，等待指定字节数的搬运完成
  • Phase 翻转：0↔1 循环，不需要重新初始化，天然支持流水线

代码实现

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
277
/**
 * 第 10 课：MBarrier 基础 —— Hopper (SM90) 同步原语
 *
 * 学习目标：
 *   - 理解 Hopper 的 mbarrier 同步机制，与 __syncthreads 对比
 *   - mbarrier 的 64-bit 数据结构（phase、arrive count、transaction count）
 *   - 四个核心 API：initialize_barrier、set_barrier_transaction_bytes、arrive_barrier、wait_barrier
 *   - Phase 翻转机制
 *
 * 实验 1：线程到达同步 — N 个线程 arrive，wait 后继续
 * 实验 2：Transaction 同步 — 设置 tx bytes，模拟异步搬运完成后 phase 翻转
 * 实验 3：Phase 翻转观察 — 多轮使用同一 barrier，观察 phase 0↔1 循环
 *
 * 核心 API（来自 cute/arch/copy_sm90_desc.hpp）：
 *   cute::initialize_barrier(uint64_t& smem_barrier, int thread_count);
 *   cute::set_barrier_transaction_bytes(uint64_t& smem_barrier, uint32_t bytes);
 *   cute::arrive_barrier(uint64_t& smem_barrier);
 *   cute::wait_barrier(uint64_t& smem_barrier, int phase_bit);
 */

#include <cute/tensor.hpp>
#include <cute/arch/copy_sm90_desc.hpp>
#include <cstdio>

using namespace cute;

// ============================================================
// 实验 1：线程到达同步（Thread Arrive/Wait）
// ============================================================
// 所有线程 arrive 后，wait 才能通过（类似 __syncthreads 但更灵活）
__global__ void kernel_thread_arrive_wait(int num_threads) {
  // mbarrier 必须在 shared memory 中
  __shared__ __align__(8) uint64_t barrier;

  int tid = threadIdx.x;

  // Step 1: 线程 0 初始化 barrier，设置期望到达的线程数
  if (tid == 0) {
    cute::initialize_barrier(barrier, num_threads);
    printf("[实验1] 初始化 mbarrier, 期望线程数 = %d\n", num_threads);
  }
  __syncthreads();  // 确保 barrier 已初始化（只需要一次）

  // Step 2: 每个线程做一些工作，然后 arrive
  if (tid < num_threads) {
    // 模拟一些计算工作
    int result = tid * tid;
    printf("  线程 %d: 计算完成 (result=%d), 准备 arrive\n", tid, result);

    // arrive：告知 barrier "我已到达"
    cute::arrive_barrier(barrier);
  }

  // Step 3: 所有线程等待 barrier 翻转（phase 0 -> 1）
  if (tid < num_threads) {
    cute::wait_barrier(barrier, 0);  // 等待 phase=0 翻转
    printf("  线程 %d: wait 通过！所有线程都已 arrive\n", tid);
  }
}

// ============================================================
// 实验 2：Producer/Consumer 模式（部分线程同步）
// ============================================================
// mbarrier 允许只有部分线程参与同步
// Producer（线程 0）先准备数据，然后 arrive
// Consumer（线程 1-3）wait 后消费数据
__global__ void kernel_producer_consumer() {
  __shared__ __align__(8) uint64_t barrier;
  __shared__ float data[32];

  int tid = threadIdx.x;

  if (tid == 0) {
    // 初始化 barrier，thread_count=1（只有 producer 会 arrive）
    cute::initialize_barrier(barrier, 1);
    printf("\n[实验2] Producer/Consumer 模式, thread_count=1 (只有 producer arrive)\n");
  }
  __syncthreads();

  if (tid == 0) {
    // Producer：准备数据
    for (int i = 0; i < 32; ++i) {
      data[i] = float(i) * 1.5f;
    }
    printf("  Producer(线程0): 数据准备完成, arrive\n");

    // arrive：通知数据就绪
    cute::arrive_barrier(barrier);
  }

  // Consumer：等待 producer 完成
  if (tid >= 1 && tid <= 4) {
    cute::wait_barrier(barrier, 0);  // 等 phase 0 翻转
    printf("  Consumer(线程%d): wait 通过, data[%d] = %.1f (期望 %.1f)\n",
           tid, tid, data[tid], float(tid) * 1.5f);
  }

  // 说明 set_barrier_transaction_bytes 的作用（仅打印说明）
  if (tid == 0) {
    printf("\n  ★ 关于 Transaction Bytes（TMA 中使用）：\n");
    printf("    set_barrier_transaction_bytes(mbar, N):\n");
    printf("      = mbarrier.arrive.expect_tx (同时 arrive + 设置预期搬运量)\n");
    printf("    TMA 硬件搬运完成后自动 complete_tx\n");
    printf("    当 arrive_count 完成 + tx_bytes 完成 → phase 翻转\n");
  }
}

// ============================================================
// 辅助函数：用 PTX 探测 mbarrier 当前 phase
// ============================================================
// mbarrier.test_wait.parity 检测 barrier 的 phase 是否 != parity
// 如果 test_wait(parity=0) 返回 true → phase 已经不是 0 → phase=1
// 如果 test_wait(parity=0) 返回 false → phase 仍然是 0
__device__ int probe_barrier_phase(uint64_t& barrier) {
  uint32_t smem_addr = cute::cast_smem_ptr_to_uint(&barrier);
  int result;
  asm volatile(
    "{\n"
    "  .reg .pred P1;\n"
    "  mbarrier.test_wait.parity.shared::cta.b64 P1, [%1], 0;\n"
    "  selp.s32 %0, 1, 0, P1;\n"  // P1=true → phase!=0 → phase=1
    "}\n"
    : "=r"(result)
    : "r"(smem_addr));
  return result;  // 返回当前 phase（0 或 1）
}

// ============================================================
// 实验 3：Phase 翻转观察（Multi-round barrier reuse）
// ============================================================
// 同一个 mbarrier 可以重复使用，每一轮 phase 翻转一次
// 这是流水线 (pipeline) 的基础
//
// 关键：mbarrier 的 phase 在每次 arrive_count 满足后自动翻转
// 不需要重新初始化！只需在下一轮 wait 时用新的 phase 值
//
// 本实验用 mbarrier.test_wait.parity PTX 指令直接探测硬件中的 phase
__global__ void kernel_phase_flip() {
  __shared__ __align__(8) uint64_t barrier;

  int tid = threadIdx.x;
  constexpr int kRounds = 4;
  constexpr int kThreads = 32;  // 必须是 warp 的整数倍

  if (tid == 0) {
    cute::initialize_barrier(barrier, kThreads);
    printf("\n[实验3] Phase 翻转观察, %d 轮, %d 线程\n", kRounds, kThreads);
    printf("  ★ 用 mbarrier.test_wait.parity PTX 探测硬件真实 phase\n\n");
  }
  __syncthreads();

  if (tid < kThreads) {
    // 用 test_wait.parity 探测初始 phase
    int current_phase = 0;
    if (tid == 0) {
      int hw_phase = probe_barrier_phase(barrier);
      printf("  初始化后: 硬件 phase = %d\n\n", hw_phase);
    }

    for (int round = 0; round < kRounds; ++round) {
      if (tid == 0) {
        printf("  --- 第 %d 轮 (当前 phase=%d, 等待翻转) ---\n", round, current_phase);
      }

      // 所有线程 arrive
      cute::arrive_barrier(barrier);

      // 等待当前 phase 翻转
      // 当所有 kThreads 个线程都 arrive 后，phase 自动从 current_phase 翻转
      cute::wait_barrier(barrier, current_phase);

      // 翻转后探测硬件真实 phase
      int new_phase = 1 - current_phase;
      if (tid == 0) {
        int hw_phase = probe_barrier_phase(barrier);
        printf("  第 %d 轮完成! phase: %d → %d (硬件探测: phase=%d)\n",
               round, current_phase, new_phase, hw_phase);
      }

      current_phase = new_phase;
    }
  }

  // 说明 phase 翻转与 pipeline 的关系
  if (tid == 0) {
    printf("\n  ★ 在 pipeline 中的应用：\n");
    printf("    每个 stage 有一对 barrier (full + empty)\n");
    printf("    full_barrier:  producer arrive → consumer wait\n");
    printf("    empty_barrier: consumer arrive → producer wait\n");
    printf("    phase 自动交替，实现环形 buffer 的多轮复用\n");
  }
}

int main() {
  printf("=== 第 10 课：MBarrier 基础 ===\n\n");

  // --------------------------------------------------------
  // 实验 1：线程到达同步
  // --------------------------------------------------------
  printf("========================================\n");
  printf("实验 1：线程到达同步\n");
  printf("========================================\n");
  printf("  类似 __syncthreads，但 mbarrier 支持：\n");
  printf("  - 只让部分线程参与同步\n");
  printf("  - 与 TMA 异步搬运配合\n");
  printf("  - Phase 翻转实现流水线\n\n");

  // 启动 8 个线程，其中 4 个参与 barrier 同步
  kernel_thread_arrive_wait<<<1, 8>>>(4);
  cudaDeviceSynchronize();

  auto err = cudaGetLastError();
  if (err != cudaSuccess) {
    printf("  CUDA error: %s\n", cudaGetErrorString(err));
    return 1;
  }

  // --------------------------------------------------------
  // 实验 2：Producer/Consumer 模式
  // --------------------------------------------------------
  printf("\n========================================\n");
  printf("实验 2：Producer/Consumer 模式\n");
  printf("========================================\n");
  printf("  mbarrier 的灵活性：\n");
  printf("  - 只有 producer 需要 arrive\n");
  printf("  - consumer 只需 wait\n");
  printf("  - 这是 Hopper warp-specialized kernel 的基础\n\n");

  kernel_producer_consumer<<<1, 32>>>();
  cudaDeviceSynchronize();

  err = cudaGetLastError();
  if (err != cudaSuccess) {
    printf("  CUDA error: %s\n", cudaGetErrorString(err));
    return 1;
  }

  // --------------------------------------------------------
  // 实验 3：Phase 翻转观察
  // --------------------------------------------------------
  printf("\n========================================\n");
  printf("实验 3：Phase 翻转\n");
  printf("========================================\n");
  printf("  同一个 barrier 重复使用：\n");
  printf("  - 第 0 轮等 phase 0→1\n");
  printf("  - 第 1 轮等 phase 1→0\n");
  printf("  - phase 自动交替，这是流水线的核心机制\n\n");

  kernel_phase_flip<<<1, 32>>>();
  cudaDeviceSynchronize();

  err = cudaGetLastError();
  if (err != cudaSuccess) {
    printf("  CUDA error: %s\n", cudaGetErrorString(err));
    return 1;
  }

  // --------------------------------------------------------
  // 总结
  // --------------------------------------------------------
  printf("\n========================================\n");
  printf("总结\n");
  printf("========================================\n");
  printf("  mbarrier vs __syncthreads：\n");
  printf("  ┌───────────────────┬──────────────────┬──────────────────┐\n");
  printf("  │     特性          │ __syncthreads    │ mbarrier         │\n");
  printf("  ├───────────────────┼──────────────────┼──────────────────┤\n");
  printf("  │ 同步范围          │ 整个 CTA         │ 可选线程子集     │\n");
  printf("  │ 异步搬运支持      │ ✗                │ ✓ (tx bytes)     │\n");
  printf("  │ Phase 翻转        │ ✗                │ ✓ (0↔1 循环)     │\n");
  printf("  │ 流水线支持        │ ✗                │ ✓ (多 stage)     │\n");
  printf("  │ TMA 配合          │ ✗                │ ✓ (硬件自动)     │\n");
  printf("  └───────────────────┴──────────────────┴──────────────────┘\n");

  printf("\n=== 练习完成 ===\n");
  return 0;
}

打印信息如下

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
$ ./10_mbarrier
=== 第 10 课：MBarrier 基础 ===

========================================
实验 1：线程到达同步
========================================
  类似 __syncthreads，但 mbarrier 支持：
  - 只让部分线程参与同步
  - 与 TMA 异步搬运配合
  - Phase 翻转实现流水线

[实验1] 初始化 mbarrier, 期望线程数 = 4
  线程 0: 计算完成 (result=0), 准备 arrive
  线程 1: 计算完成 (result=1), 准备 arrive
  线程 2: 计算完成 (result=4), 准备 arrive
  线程 3: 计算完成 (result=9), 准备 arrive
  线程 0: wait 通过！所有线程都已 arrive
  线程 1: wait 通过！所有线程都已 arrive
  线程 2: wait 通过！所有线程都已 arrive
  线程 3: wait 通过！所有线程都已 arrive

========================================
实验 2：Producer/Consumer 模式
========================================
  mbarrier 的灵活性：
  - 只有 producer 需要 arrive
  - consumer 只需 wait
  - 这是 Hopper warp-specialized kernel 的基础

[实验2] Producer/Consumer 模式, thread_count=1 (只有 producer arrive)
  Producer(线程0): 数据准备完成, arrive

  ★ 关于 Transaction Bytes（TMA 中使用）：
    set_barrier_transaction_bytes(mbar, N):
      = mbarrier.arrive.expect_tx (同时 arrive + 设置预期搬运量)
    TMA 硬件搬运完成后自动 complete_tx
    当 arrive_count 完成 + tx_bytes 完成 → phase 翻转
  Consumer(线程1): wait 通过, data[1] = 1.5 (期望 1.5)
  Consumer(线程2): wait 通过, data[2] = 3.0 (期望 3.0)
  Consumer(线程3): wait 通过, data[3] = 4.5 (期望 4.5)
  Consumer(线程4): wait 通过, data[4] = 6.0 (期望 6.0)

========================================
实验 3：Phase 翻转
========================================
  同一个 barrier 重复使用：
  - 第 0 轮等 phase 0→1
  - 第 1 轮等 phase 1→0
  - phase 自动交替，这是流水线的核心机制

[实验3] Phase 翻转观察, 4 轮, 32 线程
  ★ 用 mbarrier.test_wait.parity PTX 探测硬件真实 phase

  初始化后: 硬件 phase = 0

  --- 第 0 轮 (当前 phase=0, 等待翻转) ---
  第 0 轮完成! phase: 0 → 1 (硬件探测: phase=1)
  --- 第 1 轮 (当前 phase=1, 等待翻转) ---
  第 1 轮完成! phase: 1 → 0 (硬件探测: phase=0)
  --- 第 2 轮 (当前 phase=0, 等待翻转) ---
  第 2 轮完成! phase: 0 → 1 (硬件探测: phase=1)
  --- 第 3 轮 (当前 phase=1, 等待翻转) ---
  第 3 轮完成! phase: 1 → 0 (硬件探测: phase=0)

  ★ 在 pipeline 中的应用：
    每个 stage 有一对 barrier (full + empty)
    full_barrier:  producer arrive → consumer wait
    empty_barrier: consumer arrive → producer wait
    phase 自动交替，实现环形 buffer 的多轮复用

========================================
总结
========================================
  mbarrier vs __syncthreads：
  ┌───────────────────┬──────────────────┬──────────────────┐
  │     特性          │ __syncthreads    │ mbarrier         │
  ├───────────────────┼──────────────────┼──────────────────┤
  │ 同步范围          │ 整个 CTA         │ 可选线程子集     │
  │ 异步搬运支持      │ ✗                │ ✓ (tx bytes)     │
  │ Phase 翻转        │ ✗                │ ✓ (0↔1 循环)     │
  │ 流水线支持        │ ✗                │ ✓ (多 stage)     │
  │ TMA 配合          │ ✗                │ ✓ (硬件自动)     │
  └───────────────────┴──────────────────┴──────────────────┘

=== 练习完成 ===

1. MBarrier 的 64-bit 数据结构

1
2
3
4
5
64-bit mbarrier 状态：
  ┌─────────┬────────────────┬────────────────────┐
  │ phase   │ arrive_count   │ transaction_count   │
  │ (1 bit) │ (多位)          │ (多位)              │
  └─────────┴────────────────┴────────────────────┘
字段 含义
phase 当前阶段（0 或 1），所有条件满足后自动翻转
arrive_count 还差多少线程到达（初始化时设为 thread_count，每次 arrive 减 1）
transaction_count 还差多少字节完成（TMA 搬运时使用）

翻转条件：arrive_count == 0 && transaction_count == 0 → phase 自动翻转。

2. 四个核心 API

1
2
3
4
5
6
7
8
9
10
11
12
// 1. 初始化 barrier（设置期望到达的线程数）
cute::initialize_barrier(uint64_t& smem_barrier, int thread_count);

// 2. 设置期望搬运的字节数（TMA 用）
// 注意：这个调用内部是 mbarrier.arrive.expect_tx，会同时 arrive + 设置 tx
cute::set_barrier_transaction_bytes(uint64_t& smem_barrier, uint32_t bytes);

// 3. 通知 barrier "我已到达"
cute::arrive_barrier(uint64_t& smem_barrier);

// 4. 等待 barrier 的 phase 翻转
cute::wait_barrier(uint64_t& smem_barrier, int phase_bit);

头文件：cute/arch/copy_sm90_desc.hpp

3. 实验 1：线程到达同步

3.1 代码逻辑

1
2
3
4
5
6
7
8
// 初始化：期望 4 个线程到达
cute::initialize_barrier(barrier, 4);

// 每个线程完成计算后 arrive
cute::arrive_barrier(barrier);

// 所有线程等待 phase=0 翻转
cute::wait_barrier(barrier, 0);

3.2 状态变化

1
2
3
4
5
6
7
初始化后:  phase=0, arrive_count=4
线程0 arrive: arrive_count=3
线程1 arrive: arrive_count=2
线程2 arrive: arrive_count=1
线程3 arrive: arrive_count=0 → 满足条件 → phase 翻转: 0→1

所有 wait_barrier(barrier, 0) 的线程看到 phase 已不再是 0 → 通过

3.3 vs __syncthreads

1
2
3
4
5
6
7
8
__syncthreads():
  - 整个 CTA（block 内所有线程）都参与
  - 启动了 8 个线程 → 8 个都必须到 barrier 点

mbarrier:
  - 启动了 8 个线程，但只有 thread_count=4 个需要 arrive
  - 其余线程可以做别的事情或 wait
  - 更灵活！这是 warp specialization 的基础

4. 实验 2：Producer/Consumer 模式

4.1 核心思想

1
2
3
4
5
Producer (线程 0):
  准备数据 → arrive_barrier  → "数据就绪"
                   ↓
Consumer (线程 1-4):
  wait_barrier ←──┘  → 读取数据

4.2 关键：thread_count=1

1
cute::initialize_barrier(barrier, 1);  // 只有 1 个线程需要 arrive

只有 producer（线程 0）需要 arrive。Consumer 不 arrive，只 wait。当 producer arrive 后，arrive_count 归零，phase 翻转，所有 consumer 的 wait 都通过。

4.3 set_barrier_transaction_bytes 说明

在 TMA 场景中，搬运是由硬件异步执行的，没有"线程 arrive"这个动作。取而代之的是：

1
2
3
4
5
6
7
8
9
10
1. set_barrier_transaction_bytes(mbar, N)
   → 内部 = mbarrier.arrive.expect_tx
   → 同时做两件事：arrive + 设置预期搬运 N 字节

2. TMA 硬件搬运完成后自动执行 complete_tx
   → 减少 transaction_count

3. 当 arrive_count=0 && transaction_count=0
   → phase 翻转
   → consumer 的 wait 通过

5. 实验 3：Phase 翻转

5.1 核心发现

Phase 在每轮 arrive 完成后自动翻转，不需要重新初始化。

本实验用 PTX mbarrier.test_wait.parity 指令直接从硬件探测真实 phase，而不是手动推算。

5.2 探测 Phase 的方法

1
2
3
4
5
6
7
8
9
10
11
12
13
// 用 PTX 探测 mbarrier 当前 phase
__device__ int probe_barrier_phase(uint64_t& barrier) {
  uint32_t smem_addr = cute::cast_smem_ptr_to_uint(&barrier);
  int result;
  asm volatile(
    "{\n"
    "  .reg .pred P1;\n"
    "  mbarrier.test_wait.parity.shared::cta.b64 P1, [%1], 0;\n"
    "  selp.s32 %0, 1, 0, P1;\n"  // P1=true → phase!=0 → phase=1
    "}\n"
    : "=r"(result) : "r"(smem_addr));
  return result;
}

mbarrier.test_wait.parity 的语义：检测 barrier 的 phase 是否 != 传入的 parity。

  • test_wait.parity(mbar, 0) 返回 true → phase != 0 → phase=1
  • test_wait.parity(mbar, 0) 返回 false → phase == 0

注意：mbarrier 没有"读取 phase"的直接 API，test_wait.parity 是唯一可用的探测手段。

5.3 代码逻辑

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
cute::initialize_barrier(barrier, 32);

// 探测初始 phase
int hw_phase = probe_barrier_phase(barrier);  // → 0

int current_phase = 0;
for (int round = 0; round < 4; ++round) {
    cute::arrive_barrier(barrier);
    cute::wait_barrier(barrier, current_phase);

    int new_phase = 1 - current_phase;
    hw_phase = probe_barrier_phase(barrier);
    // hw_phase == new_phase → 硬件确认翻转
    current_phase = new_phase;
}

5.4 Phase 翻转时间线（硬件探测验证）

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
初始化: phase=0 (硬件探测: 0), arrive_count=32

第 0 轮:
  32 个线程 arrive → arrive_count 归零
  phase 翻转: 0→1
  wait_barrier(barrier, 0) → 通过 ✓
  硬件探测: phase=1 ✓
  barrier 自动重置 arrive_count=32

第 1 轮:
  32 个线程 arrive → arrive_count 归零
  phase 翻转: 1→0
  wait_barrier(barrier, 1) → 通过 ✓
  硬件探测: phase=0 ✓
  barrier 自动重置 arrive_count=32

第 2 轮: phase 0→1 (硬件探测: 1) ✓
第 3 轮: phase 1→0 (硬件探测: 0) ✓

5.3 为什么必须是 warp 整数倍线程

1
2
3
4
5
6
线程数 = 32 (1 个 warp)

如果用 17 个线程：
  - arrive_barrier 内部使用 warp 级 arrive 指令
  - 同一 warp 内的线程如果一部分 arrive 一部分不 arrive，行为未定义
  - 必须保证参与的线程数是 warp (32) 的整数倍

5.4 在 Pipeline 中的应用

1
2
3
4
5
6
7
8
9
10
11
12
13
14
Pipeline 每个 stage 有一对 barrier：

  full_barrier:
    producer 搬完数据 → arrive
    consumer wait → 数据就绪，可以消费

  empty_barrier:
    consumer 消费完 → arrive
    producer wait → buffer 空了，可以写入新数据

  phase 自动翻转：
    第 1 轮: producer arrive → phase 0→1 → consumer wait 通过
    第 2 轮: producer arrive → phase 1→0 → consumer wait 通过
    ... 无限循环，无需重新初始化

6. MBarrier vs __syncthreads 对比

特性 __syncthreads mbarrier
同步范围 整个 CTA (block) 可选线程子集
异步搬运支持 不支持 支持 (transaction bytes)
Phase 翻转 不支持 支持 (0↔1 自动循环)
流水线支持 不支持 支持 (多 stage × 多 round)
TMA 配合 不支持 支持 (硬件自动 complete_tx)
存储位置 无需分配 shared memory (64-bit)
初始化 不需要 需要 (thread_count)
硬件支持 所有架构 SM90+ (Hopper)

7. API 总结

API 作用 使用场景
initialize_barrier(mbar, N) 创建 barrier，期望 N 个 arrive 只需调用一次
set_barrier_transaction_bytes(mbar, bytes) arrive + 设置期望搬运字节 TMA 搬运前调用
arrive_barrier(mbar) 通知 “我已到达” 线程完成工作后
wait_barrier(mbar, phase) 等待 phase 翻转 等待同步点

8. 后续课程关联

课程 mbarrier 的使用
11 课 (TMA Copy) thread 0 初始化 mbarrier + set_tx_bytes，TMA 搬运完自动 complete_tx，所有线程 wait
12 课 (Pipeline) PipelineTmaAsync 内部管理每 stage 一对 mbarrier，封装了 arrive/wait
13 课 (GEMM) Producer warp 通过 pipeline.producer_get_barrier 获取 mbarrier，配合 TMA 搬运
  • 标题: cute（10）hopper之mbarrier
  • 作者: 鱿鱼圈
  • 创建于 : 2026-06-17 02:13:32
  • 更新于 : 2026-06-14 22:17:19
  • 链接: https://yuyanqi.com/2026/06/17/cute(10)mbarrier/
  • 版权声明: 本文章采用 CC BY-NC-SA 4.0 进行许可。
评论
目录
cute（10）hopper之mbarrier
  1. 前置阅读
  2. 10 MBarrier 基础
    1. 核心概念
    2. 代码实现
    3. 1. MBarrier 的 64-bit 数据结构
    4. 2. 四个核心 API
    5. 3. 实验 1：线程到达同步
      1. 3.1 代码逻辑
      2. 3.2 状态变化
      3. 3.3 vs __syncthreads
    6. 4. 实验 2：Producer/Consumer 模式
      1. 4.1 核心思想
      2. 4.2 关键：thread_count=1
      3. 4.3 set_barrier_transaction_bytes 说明
    7. 5. 实验 3：Phase 翻转
      1. 5.1 核心发现
      2. 5.2 探测 Phase 的方法
      3. 5.3 代码逻辑
      4. 5.4 Phase 翻转时间线（硬件探测验证）
      5. 5.3 为什么必须是 warp 整数倍线程
      6. 5.4 在 Pipeline 中的应用
    8. 6. MBarrier vs __syncthreads 对比
    9. 7. API 总结
    10. 8. 后续课程关联