Барьерный механизм CUDA между TMA и вычислительными потоками ⇐ C++

[Anonymous]

Как говорится в спецификации PTX, Mbarrier инициализируется с помощью phase = 0; Почему в потоке Cutlass TMA первый empty_barrier_ptr_ ожидает фазы 1; Будет ли у него тупик? Если нет тупика, в чем причина? < /P>
CUTLASS_DEVICE
void producer_acquire(uint32_t stage, uint32_t phase)
{
empty_barrier_ptr_[stage].wait(phase);

if (params_.is_leader)
{
full_barrier_ptr_[stage].arrive_and_expect_tx(params_.transaction_bytes);
}
#ifndef NDEBUG
if (params_.role == ThreadCategory::Consumer || params_.role == ThreadCategory::NonParticipant)
{
asm volatile("brkpt;\n" ::);
}

// Most likely you have elected more than one leader
if (params_.is_leader && (threadIdx.x % 32 != 0))
{
asm volatile("brkpt;\n" ::);
}
#endif
}


Подробнее здесь: https://stackoverflow.com/questions/797 ... te-threads