bar.arrive PTX 屏障同步指令的线程数意味着什么？

Question

PTX文档中here提到bar.sync和bar.arrive屏障同步指令可以如下使用：

bar.sync      a{, b};
bar.arrive    a, b;

哪里

Source operand a specifies a logical barrier resource as an immediate constant or register with value 0 through 15. Operand b specifies the number of threads participating in the barrier.

它还展示了一个使用以下指令建立生产者-消费者模型的示例：

// Producer code places produced value in shared memory.
st.shared [r0],r1; 
bar.arrive 0,64;
...

// Consumer code, reads value from shared memory 
bar.sync 0,64; 
ld.shared r1,[r0];
...

我不太明白 bar.arrive 中操作数 b 的用途。虽然 bar.sync 中的此类操作数可用于控制屏障中涉及的线程数并等待达到线程数，但我不清楚它在 bar.arrive 中的用途。

Answer 1

当所有线程都到达屏障时会发生两种情况：

允许所有等待的线程越过障碍继续进行。
屏障已重新初始化，因此可以再次使用。

您可能只想到 1.，它只能在 bar.sync 指令中发生。因此很明显 bar.sync 需要知道参与屏障的线程数。但是，屏障也可以在 bar.arrive 时释放，因此还需要知道参与线程的数量。

话虽如此，但没有记录如果参与的 warp 甚至线程对涉及的线程数不一致会发生什么。这可以被看作是一个大胆的好奇心通过逆向工程找到可能的新（和不受支持的！）同步结构的机会。

bar.arrive PTX 屏障同步指令的线程数意味着什么？

What does thread-count mean for bar.arrive PTX barrier synchronization instruction?

cuda

ptx