Home | History | Annotate | Download | only in nvptx

Lines Matching refs:waiters

69   unsigned int waiters
70     = __atomic_add_fetch (&bar->waiters, 1, MEMMODEL_ACQ_REL);
71   if (waiters == 0)
73   unsigned int waiter_id = waiters;
75   if (waiters > 1)
78       asm volatile ("bar.sync 1, %0;" : : "r" (32 * waiters));
80       /* Ensure that they have updated waiters.  */
81       asm volatile ("bar.sync 1, %0;" : : "r" (32 * waiters));
89       asm volatile ("bar.sync 1, %0;" : : "r" (32 * (waiters + 1)));
91       /* Get updated waiters.  */
93 	= __atomic_load_n (&bar->waiters, MEMMODEL_ACQUIRE);
95       /* Notify that we have updated waiters.  */
96       asm volatile ("bar.sync 1, %0;" : : "r" (32 * (waiters + 1)));
98       waiters = updated_waiters;
100       if (waiter_id > waiters)
121   unsigned int waiters = __atomic_load_n (&bar->waiters, MEMMODEL_ACQUIRE);
122   if (waiters == 0)
131     __atomic_store_n (&bar->waiters, 0, MEMMODEL_RELEASE);
134     __atomic_add_fetch (&bar->waiters, -count, MEMMODEL_ACQ_REL);
140   asm volatile ("bar.sync 1, %0;" : : "r" (32 * (waiters + 1)));
142   /* Let them get the updated waiters.  */
143   asm volatile ("bar.sync 1, %0;" : : "r" (32 * (waiters + 1)));