.. _program_listing_file_include_embers_primitives_ticketlock_impl.h:

Program Listing for File ticketlock_impl.h
==========================================

|exhale_lsh| :ref:`Return to documentation for file <file_include_embers_primitives_ticketlock_impl.h>` (``include/embers/primitives/ticketlock_impl.h``)

.. |exhale_lsh| unicode:: U+021B0 .. UPWARDS ARROW WITH TIP LEFTWARDS

.. code-block:: cpp

   /* Copyright © 2020 Advanced Micro Devices, Inc. All rights reserved */
   
   #ifndef _EMBERS_TICKET_LOCK_IMPL_H
   #define _EMBERS_TICKET_LOCK_IMPL_H
   
   #include "embers/primitives/backoff.h"
   #include "embers/primitives/ticketlock.h"
   
   namespace embers
   {
   
   template <MemoryScope scope, size_t BackoffBase>
   TicketLock<scope, BackoffBase>::TicketLock()
       : next_ticket_(atomic<uint32_t>(0)), now_serving_(atomic<uint32_t>(0))
   {
   }
   
   template <MemoryScope scope, size_t BackoffBase>
   __host__ __device__ void TicketLock<scope, BackoffBase>::Acquire() noexcept
   {
     auto ticket = next_ticket_.fetch_inc(std::memory_order_relaxed);
   
     while (true) {
       auto current_ticket = now_serving_.load(std::memory_order_relaxed);
       if (ticket == current_ticket) {
         break;
       }
       const auto num_before = static_cast<size_t>(ticket) - static_cast<size_t>(current_ticket);
       const size_t num_waits = num_before * BackoffBase;
       for (size_t wait = 0; wait < num_waits; wait++) {
         backoff();
       }
     }
     [[maybe_unused]] auto temp = now_serving_.load(std::memory_order_acquire);
   }
   
   template <MemoryScope scope, size_t BackoffBase>
   __host__ __device__ void TicketLock<scope, BackoffBase>::Release() noexcept
   {
     now_serving_.fetch_add(1, std::memory_order_release);
   }
   
   }  // namespace embers
   #endif  // _EMBERS_TICKET_LOCK_IMPL_H