https://gcc.gnu.org/bugzilla/show_bug.cgi?id=85246
--- Comment #3 from Tom de Vries <vries at gcc dot gnu.org> --- (In reply to Tom de Vries from comment #1) > I went through a couple of cycles of minimizing the failure, seeing > something suspicious, modifying by hand or writing a tentative patch, but > every time I went back to the original non-minimized example I got the > failure again. > > Anyway, things that may be causing this fail: > 2. > > The bar.sync instruction may not be sufficiently understood. > > In the documentation for bar.sync it says: > ... > bar.sync and bar.red also guarantee memory ordering among threads identical > to > membar.cta . Thus, threads within a CTA that wish to communicate via memory > can > store to memory, execute a bar.sync or bar.red instruction, and then safely > read > values stored by other threads prior to the barrier. > ... > > The question is what happens when you specify a thread count. Does the > memory ordering still apply to the whole CTA, or only to the threads > participating in the barrier? > > So if we store something in vector id 0, worker id 0, and load it in worker > id 1, we may have to use a bar.sync 0 instead to synchronize (or keep the > same barrier but add a membar.cta). I misanalyzed why the test was failing, it was not a barrier problem. There are two situations in which there is state propagation: 1. when transitioning from single to partition mode, for either worker or vector (so, when entering a partitioned loop) 2. when propagating branch conditions. [ We implement worker-single and vector-single by branch around, but do not branch around branches, so the branch condition is calculated in either W0V0 or WAV0 code, and then used in WAVA code. To go from worker single to worker partitioned (W0V0 -> WAV0), we use a generic broadcast buffer. To go from vector single to vector partitioned (WAV0 -> WAVA), we use a worker-specific broadcast buffer. To propagate the branch condition, we use the worker-specific broadcast buffer, but that only works for WAV0 -> WAVA. For the W0V0 -> WAVA propagation, we need to use the generic broadcast buffer.
