#include <cstdio>
#include <string>
#include <assert.h>
#include <math.h>
#include "hip/hip_runtime.h"
#include "hip_error.h"

#define MAX_BACKOFF             1024
#define NUM_WIS_PER_WG 64
#define MAD_MUL 1.1f
#define MAD_ADD 0.25f
#define NUM_WORDS_PER_CACHELINE 16
#define NUM_WIS_PER_QUARTERWAVE 16

// separate .h files
#include "hipLocksBarrier.h"
#include "hipLocksImpl.h"
#include "hipLocksMutex.h"
#include "hipLocksSemaphore.h"
// program globals
const int NUM_REPEATS = 10;
int NUM_LDST = 0;
int numWGs = 0;
// number of CUs our GPU has
int NUM_CU = 0;
int MAX_WGS = 0;

bool pageAlign = false;

/*
  Helper function to do data accesses for golden checking code.
*/
void accessData_golden(float * storageGolden, int currLoc, int numStorageLocs)
{
  /*
    If this location isn't the first location accessed by a
    thread, update it -- each half-warp accesses (NUM_LDST + 1) cache
    lines.
  */
  if (currLoc % (NUM_WIS_PER_QUARTERWAVE * (NUM_LDST + 1)) >=
      NUM_WORDS_PER_CACHELINE)
  {
    assert((currLoc - NUM_WORDS_PER_CACHELINE) >= 0);
    assert(currLoc < numStorageLocs);
    // each location is dependent on the location accessed at the
    // same word on the previous cache line
    storageGolden[currLoc] =
      ((storageGolden[currLoc -
                      NUM_WORDS_PER_CACHELINE]/* * MAD_MUL*/) /*+ MAD_ADD*/);
  }
}

/*
  Shared function that does the critical section data accesses for the barrier
  and mutex kernels.

  NOTE: kernels that access data differently (e.g., semaphores) should not call
  this function.
*/
inline __device__ void accessData(float * storage, int threadBaseLoc,
                                  int threadOffset, int NUM_LDST)
{
  // local variables
  int readLoc = 0, writeLoc = 0;

  for (int n = NUM_LDST-1; n >= 0; --n) {
    writeLoc = ((threadBaseLoc + n + 1) * NUM_WORDS_PER_CACHELINE) +
               threadOffset;
    readLoc = ((threadBaseLoc + n) * NUM_WORDS_PER_CACHELINE) + threadOffset;
    storage[writeLoc] = ((storage[readLoc]/* * MAD_MUL*/) /*+ MAD_ADD*/);
  }
}

/*
  Helper function for semaphore writers.  Although semaphore writers are
  iterating over all locations accessed on a given CU, the logic for this
  varies and is done outside this helper, so can just call accessData().
*/
inline __device__ void accessData_semWr(float * storage, int threadBaseLoc,
                                        int threadOffset, int NUM_LDST)
{
  accessData(storage, threadBaseLoc, threadOffset, NUM_LDST);
}


/*
  Helper function for semaphore readers.
*/
inline __device__ void accessData_semRd(float * storage,
                                        volatile float * dummyArray,
                                        int threadBaseLoc,
                                        int threadOffset, int NUM_LDST)
{
  for (int n = NUM_LDST-1; n >= 0; --n) {
    dummyArray[hipThreadIdx_x] +=
      storage[((threadBaseLoc + n) * NUM_WORDS_PER_CACHELINE) +
              threadOffset];
    __syncthreads();
  }
}

// performs a tree barrier.  Each WG on an CU accesses unique data then joins a
// local barrier.  1 of the WGs from each CU then joins the global barrier
__global__ void kernelAtomicTreeBarrierUniq(float * storage,
                                            hipLockData_t * gpuLockData,
                                            unsigned int * perCUBarrierBuffers,
                                            const int ITERATIONS,
                                            const int NUM_LDST,
                                            const int NUM_CU,
                                            const int MAX_WGS)
{
  // local variables
  // thread 0 is master thread
  const bool isMasterThread = ((hipThreadIdx_x == 0) && (hipThreadIdx_y == 0) &&
                               (hipThreadIdx_z == 0));
  // represents the number of WGs going to the barrier (max NUM_CU, hipGridDim_x if
  // fewer WGs than CUs).
  const unsigned int numWGsAtBarr = ((hipGridDim_x < NUM_CU) ? hipGridDim_x :
                                        NUM_CU);
  const int cuID = (hipBlockIdx_x % numWGsAtBarr); // mod by # CUs to get CU ID
  // all work groups on the same CU access unique locations because the
  // barrier can't ensure DRF between WGs
  int currWGID = hipBlockIdx_x;
  int tid = ((currWGID * hipBlockDim_x) + hipThreadIdx_x);
  // want striding to happen across cache lines so that each thread in a
  // half-warp accesses sequential words
  int threadBaseLoc = ((tid/NUM_WORDS_PER_CACHELINE) * (NUM_LDST+1));
  int threadOffset = (hipThreadIdx_x % NUM_WORDS_PER_CACHELINE);
  // determine if I'm WG 0 on my CU
  const int perCU_WGID = (hipBlockIdx_x / numWGsAtBarr);
  // given the hipGridDim_x, we can figure out how many WGs are on our CU -- assume
  // all CUs have an identical number of WGs
  int numWGs_perCU = (int)ceil((float)hipGridDim_x / numWGsAtBarr);

  for (int i = 0; i < ITERATIONS; ++i)
  {
    accessData(storage, threadBaseLoc, threadOffset, NUM_LDST);

    joinBarrier_helper(gpuLockData->barrierBuffers, perCUBarrierBuffers,
                       numWGsAtBarr, cuID, perCU_WGID, numWGs_perCU,
                       isMasterThread, MAX_WGS);

    // get new thread ID by trading amongst WGs -- + 1 WG ID to shift to next
    // CUs data
    currWGID = ((currWGID + 1) % hipGridDim_x);
    tid = ((currWGID * hipBlockDim_x) + hipThreadIdx_x);
    threadBaseLoc = ((tid/NUM_WORDS_PER_CACHELINE) * (NUM_LDST+1));
  }
}

// like the tree barrier but also has WGs exchange work locally before joining
// the global barrier
__global__ void kernelAtomicTreeBarrierUniqLocalExch(float * storage,
                                                     hipLockData_t * gpuLockData,
                                                     unsigned int * perCUBarrierBuffers,
                                                     const int ITERATIONS,
                                                     const int NUM_LDST,
                                                     const int NUM_CU,
                                                     const int MAX_WGS)
{
  // local variables
  // thread 0 is master thread
  const bool isMasterThread = ((hipThreadIdx_x == 0) && (hipThreadIdx_y == 0) &&
                               (hipThreadIdx_z == 0));
  // represents the number of WGs going to the barrier (max NUM_CU, hipGridDim_x if
  // fewer WGs than CUs).
  const unsigned int numWGsAtBarr = ((hipGridDim_x < NUM_CU) ? hipGridDim_x :
                                        NUM_CU);
  const int cuID = (hipBlockIdx_x % numWGsAtBarr); // mod by # CUs to get CU ID
  // all work groups on the same CU access unique locations because the
  // barrier can't ensure DRF between WGs
  int currWGID = hipBlockIdx_x;
  int tid = ((currWGID * hipBlockDim_x) + hipThreadIdx_x);
  // want striding to happen across cache lines so that each thread in a
  // half-warp accesses sequential words
  int threadBaseLoc = ((tid/NUM_WORDS_PER_CACHELINE) * (NUM_LDST+1));
  int threadOffset = (hipThreadIdx_x % NUM_WORDS_PER_CACHELINE);
  // determine if I'm WG 0 on my CU
  const int perCU_WGID = (hipBlockIdx_x / numWGsAtBarr);
  // given the hipGridDim_x, we can figure out how many WGs are on our CU -- assume
  // all CUs have an identical number of WGs
  int numWGs_perCU = (int)ceil((float)hipGridDim_x / numWGsAtBarr);

  for (int i = 0; i < ITERATIONS; ++i)
  {
    accessData(storage, threadBaseLoc, threadOffset, NUM_LDST);

    // all WGs on this CU do a local barrier (if > 1 WG)
    if (numWGs_perCU > 1) {
      hipBarrierAtomicLocal(perCUBarrierBuffers, cuID, numWGs_perCU,
                             isMasterThread, MAX_WGS);

      // exchange data within the WGs on this CU, then do some more computations
      currWGID = ((currWGID + numWGsAtBarr) % hipGridDim_x);
      tid = ((currWGID * hipBlockDim_x) + hipThreadIdx_x);
      threadBaseLoc = ((tid/NUM_WORDS_PER_CACHELINE) * (NUM_LDST+1));

      accessData(storage, threadBaseLoc, threadOffset, NUM_LDST);
    }

    joinBarrier_helper(gpuLockData->barrierBuffers, perCUBarrierBuffers,
                       numWGsAtBarr, cuID, perCU_WGID, numWGs_perCU,
                       isMasterThread, MAX_WGS);

    // get new thread ID by trading amongst WGs -- + 1 WG ID to shift to next
    // CUs data
    currWGID = ((currWGID + 1) % hipGridDim_x);
    tid = ((currWGID * hipBlockDim_x) + hipThreadIdx_x);
    threadBaseLoc = ((tid/NUM_WORDS_PER_CACHELINE) * (NUM_LDST+1));
  }
}

// performs a tree barrier like above but with a lock-free barrier
__global__ void kernelFBSTreeBarrierUniq(float * storage,
                                         hipLockData_t * gpuLockData,
                                         unsigned int * perCUBarrierBuffers,
                                         const int ITERATIONS,
                                         const int NUM_LDST,
                                         const int NUM_CU,
                                         const int MAX_WGS)
{
  // local variables
  // represents the number of WGs going to the barrier (max NUM_CU, hipGridDim_x if
  // fewer WGs than CUs).
  const unsigned int numWGsAtBarr = ((hipGridDim_x < NUM_CU) ? hipGridDim_x :
                                        NUM_CU);
  const int cuID = (hipBlockIdx_x % numWGsAtBarr); // mod by # CUs to get CU ID
  // all work groups on the same CU access unique locations because the
  // barrier can't ensure DRF between WGs
  int currWGID = hipBlockIdx_x;
  int tid = ((currWGID * hipBlockDim_x) + hipThreadIdx_x);
  // want striding to happen across cache lines so that each thread in a
  // half-warp accesses sequential words
  int threadBaseLoc = ((tid/NUM_WORDS_PER_CACHELINE) * (NUM_LDST+1));
  const int threadOffset = (hipThreadIdx_x % NUM_WORDS_PER_CACHELINE);
  // determine if I'm WG 0 on my CU
  const int perCU_WGID = (hipBlockIdx_x / numWGsAtBarr);
  // given the hipGridDim_x, we can figure out how many WGs are on our CU -- assume
  // all CUs have an identical number of WGs
  int numWGs_perCU = (int)ceil((float)hipGridDim_x/numWGsAtBarr);

  for (int i = 0; i < ITERATIONS; ++i)
  {
    accessData(storage, threadBaseLoc, threadOffset, NUM_LDST);

    joinLFBarrier_helper(gpuLockData->barrierBuffers, perCUBarrierBuffers,
                         numWGsAtBarr, cuID, perCU_WGID, numWGs_perCU,
                         gpuLockData->arrayStride, MAX_WGS);

    // get new thread ID by trading amongst WGs -- + 1 WG ID to shift to next
    // CUs data
    currWGID = ((currWGID + 1) % hipGridDim_x);
    tid = ((currWGID * hipBlockDim_x) + hipThreadIdx_x);
    threadBaseLoc = ((tid/NUM_WORDS_PER_CACHELINE) * (NUM_LDST+1));
  }
}

// performs a tree barrier like above but with a lock-free barrier and has WGs
// exchange work locally before joining the global barrier
__global__ void kernelFBSTreeBarrierUniqLocalExch(float * storage,
                                                  hipLockData_t * gpuLockData,
                                                  unsigned int * perCUBarrierBuffers,
                                                  const int ITERATIONS,
                                                  const int NUM_LDST,
                                                  const int NUM_CU,
                                                  const int MAX_WGS)
{
  // local variables
  // represents the number of WGs going to the barrier (max NUM_CU, hipGridDim_x if
  // fewer WGs than CUs).
  const unsigned int numWGsAtBarr = ((hipGridDim_x < NUM_CU) ? hipGridDim_x : NUM_CU);
  const int cuID = (hipBlockIdx_x % numWGsAtBarr); // mod by # CUs to get CU ID
  // all work groups on the same CU access unique locations because the
  // barrier can't ensure DRF between WGs
  int currWGID = hipBlockIdx_x;
  int tid = ((currWGID * hipBlockDim_x) + hipThreadIdx_x);
  // want striding to happen across cache lines so that each thread in a
  // half-warp accesses sequential words
  int threadBaseLoc = ((tid/NUM_WORDS_PER_CACHELINE) * (NUM_LDST+1));
  const int threadOffset = (hipThreadIdx_x % NUM_WORDS_PER_CACHELINE);
  // determine if I'm WG 0 on my CU
  const int perCU_WGID = (hipBlockIdx_x / numWGsAtBarr);
  // given the hipGridDim_x, we can figure out how many WGs are on our CU -- assume
  // all CUs have an identical number of WGs
  int numWGs_perCU = (int)ceil((float)hipGridDim_x/numWGsAtBarr);

  for (int i = 0; i < ITERATIONS; ++i)
  {
    accessData(storage, threadBaseLoc, threadOffset, NUM_LDST);

    // all WGs on this CU do a local barrier (if > 1 WG per CU)
    if (numWGs_perCU > 1) {
      hipBarrierLocal(gpuLockData->barrierBuffers, numWGsAtBarr,
                       gpuLockData->arrayStride, perCUBarrierBuffers, cuID, numWGs_perCU,
                       perCU_WGID, false, MAX_WGS);

      // exchange data within the WGs on this CU and do some more computations
      currWGID = ((currWGID + numWGsAtBarr) % hipGridDim_x);
      tid = ((currWGID * hipBlockDim_x) + hipThreadIdx_x);
      threadBaseLoc = ((tid/NUM_WORDS_PER_CACHELINE) * (NUM_LDST+1));

      accessData(storage, threadBaseLoc, threadOffset, NUM_LDST);
    }

    joinLFBarrier_helper(gpuLockData->barrierBuffers, perCUBarrierBuffers,
                         numWGsAtBarr, cuID, perCU_WGID, numWGs_perCU,
                         gpuLockData->arrayStride, MAX_WGS);

    // get new thread ID by trading amongst WGs -- + 1 WG ID to shift to next
    // CUs data
    currWGID = ((currWGID + 1) % hipGridDim_x);
    tid = ((currWGID * hipBlockDim_x) + hipThreadIdx_x);
    threadBaseLoc = ((tid/NUM_WORDS_PER_CACHELINE) * (NUM_LDST+1));
  }
}

__global__ void kernelSleepingMutex(hipMutex_t mutex, float * storage,
                                    hipLockData_t * gpuLockData,
                                    const int ITERATIONS, const int NUM_LDST,
                                    const int NUM_CU)
{
  // local variables
  // all work groups access the same locations (rely on release to get
  // ownership in time)
  const int tid = hipThreadIdx_x;
  // want striding to happen across cache lines so that each thread in a
  // half-warp accesses sequential words
  const int threadBaseLoc = ((tid/NUM_WORDS_PER_CACHELINE) * (NUM_LDST+1));
  const int threadOffset = (hipThreadIdx_x % NUM_WORDS_PER_CACHELINE);
  __shared__ int myRingBufferLoc; // tracks my WGs location in the ring buffer

  if (hipThreadIdx_x == 0) {
    myRingBufferLoc = -1; // initially I have no location
  }
  __syncthreads();

  for (int i = 0; i < ITERATIONS; ++i)
  {
    myRingBufferLoc = hipMutexSleepLock(mutex, gpuLockData->mutexBuffers,
                                        gpuLockData->mutexBufferTails, gpuLockData->maxBufferSize,
                                        gpuLockData->arrayStride, NUM_CU);

    accessData(storage, threadBaseLoc, threadOffset, NUM_LDST);

    hipMutexSleepUnlock(mutex, gpuLockData->mutexBuffers, myRingBufferLoc,
                         gpuLockData->maxBufferSize, gpuLockData->arrayStride, NUM_CU);
  }
}

__global__ void kernelSleepingMutexUniq(hipMutex_t mutex, float * storage,
                                        hipLockData_t * gpuLockData,
                                        const int ITERATIONS,
                                        const int NUM_LDST, const int NUM_CU)
{
  // local variables
  const int cuID = (hipBlockIdx_x % NUM_CU); // mod by # CUs to get CU ID
  // all work groups on the same CU access the same locations
  const int tid = ((cuID * hipBlockDim_x) + hipThreadIdx_x);
  // want striding to happen across cache lines so that each thread in a
  // half-warp accesses sequential words
  const int threadBaseLoc = ((tid/NUM_WORDS_PER_CACHELINE) * (NUM_LDST+1));
  const int threadOffset = (hipThreadIdx_x % NUM_WORDS_PER_CACHELINE);
  __shared__ int myRingBufferLoc; // tracks my WGs location in the ring buffer

  if (hipThreadIdx_x == 0) {
    myRingBufferLoc = -1; // initially I have no location
  }
  __syncthreads();

  for (int i = 0; i < ITERATIONS; ++i)
  {
    myRingBufferLoc = hipMutexSleepLockLocal(mutex, cuID, gpuLockData->mutexBuffers,
                                              gpuLockData->mutexBufferTails, gpuLockData->maxBufferSize,
                                              gpuLockData->arrayStride, NUM_CU);

    accessData(storage, threadBaseLoc, threadOffset, NUM_LDST);

    hipMutexSleepUnlockLocal(mutex, cuID, gpuLockData->mutexBuffers, myRingBufferLoc,
                              gpuLockData->maxBufferSize, gpuLockData->arrayStride,
                              NUM_CU);
  }
}

__global__ void kernelFetchAndAddMutex(hipMutex_t mutex, float * storage,
                                       hipLockData_t * gpuLockData,
                                       const int ITERATIONS,
                                       const int NUM_LDST, const int NUM_CU)
{
  // local variables
  // all work groups access the same locations (rely on release to get
  // ownership in time)
  const int tid = hipThreadIdx_x;
  // want striding to happen across cache lines so that each thread in a
  // half-warp accesses sequential words
  const int threadBaseLoc = ((tid/NUM_WORDS_PER_CACHELINE) * (NUM_LDST+1));
  const int threadOffset = (hipThreadIdx_x % NUM_WORDS_PER_CACHELINE);

  for (int i = 0; i < ITERATIONS; ++i)
  {
    hipMutexFALock(mutex, gpuLockData->mutexBufferHeads, gpuLockData->mutexBufferTails,
                    NUM_CU);

    accessData(storage, threadBaseLoc, threadOffset, NUM_LDST);

    hipMutexFAUnlock(mutex, gpuLockData->mutexBufferTails, NUM_CU);
  }
}

__global__ void kernelFetchAndAddMutexUniq(hipMutex_t mutex, float * storage,
                                           hipLockData_t * gpuLockData,
                                           const int ITERATIONS,
                                           const int NUM_LDST,
                                           const int NUM_CU)
{
  // local variables
  const int cuID = (hipBlockIdx_x % NUM_CU); // mod by # CUs to get CU ID
  // all work groups on the same CU access the same locations
  const int tid = ((cuID * hipBlockDim_x) + hipThreadIdx_x);
  // want striding to happen across cache lines so that each thread in a
  // half-warp accesses sequential words
  const int threadBaseLoc = ((tid/NUM_WORDS_PER_CACHELINE) * (NUM_LDST+1));
  const int threadOffset = (hipThreadIdx_x % NUM_WORDS_PER_CACHELINE);

  for (int i = 0; i < ITERATIONS; ++i)
  {
    hipMutexFALockLocal(mutex, cuID, gpuLockData->mutexBufferHeads,
                         gpuLockData->mutexBufferTails, NUM_CU);

    accessData(storage, threadBaseLoc, threadOffset, NUM_LDST);

    hipMutexFAUnlockLocal(mutex, cuID, gpuLockData->mutexBufferTails, NUM_CU);
  }
}

__global__ void kernelSpinLockMutex(hipMutex_t mutex, float * storage,
                                    hipLockData_t * gpuLockData,
                                    const int ITERATIONS, const int NUM_LDST,
                                    const int NUM_CU)
{
  // local variables
  // all work groups access the same locations (rely on release to get
  // ownership in time)
  const int tid = hipThreadIdx_x;
  // want striding to happen across cache lines so that each thread in a
  // half-warp accesses sequential words
  const int threadBaseLoc = ((tid/NUM_WORDS_PER_CACHELINE) * (NUM_LDST+1));
  const int threadOffset = (hipThreadIdx_x % NUM_WORDS_PER_CACHELINE);

  for (int i = 0; i < ITERATIONS; ++i)
  {
    hipMutexSpinLock(mutex, gpuLockData->mutexBufferHeads, NUM_CU);

    accessData(storage, threadBaseLoc, threadOffset, NUM_LDST);

    hipMutexSpinUnlock(mutex, gpuLockData->mutexBufferHeads, NUM_CU);
  }
}

__global__ void kernelSpinLockMutexUniq(hipMutex_t mutex, float * storage,
                                        hipLockData_t * gpuLockData,
                                        const int ITERATIONS,
                                        const int NUM_LDST, const int NUM_CU)
{
  // local variables
  const int cuID = (hipBlockIdx_x % NUM_CU); // mod by # CUs to get CU ID
  // all work groups on the same CU access the same locations
  const int tid = ((cuID * hipBlockDim_x) + hipThreadIdx_x);
  // want striding to happen across cache lines so that each thread in a
  // half-warp accesses sequential words
  const int threadBaseLoc = ((tid/NUM_WORDS_PER_CACHELINE) * (NUM_LDST+1));
  const int threadOffset = (hipThreadIdx_x % NUM_WORDS_PER_CACHELINE);

  for (int i = 0; i < ITERATIONS; ++i)
  {
    hipMutexSpinLockLocal(mutex, cuID, gpuLockData->mutexBufferHeads, NUM_CU);

    accessData(storage, threadBaseLoc, threadOffset, NUM_LDST);

    hipMutexSpinUnlockLocal(mutex, cuID, gpuLockData->mutexBufferHeads, NUM_CU);
  }
}

__global__ void kernelEBOMutex(hipMutex_t mutex, float * storage,
                               hipLockData_t * gpuLockData,
                               const int ITERATIONS, const int NUM_LDST,
                               const int NUM_CU)
{
  // local variables
  // all work groups access the same locations (rely on release to get
  // ownership in time)
  const int tid = hipThreadIdx_x;
  // want striding to happen across cache lines so that each thread in a
  // half-warp accesses sequential words
  const int threadBaseLoc = ((tid/NUM_WORDS_PER_CACHELINE) * (NUM_LDST+1));
  const int threadOffset = (hipThreadIdx_x % NUM_WORDS_PER_CACHELINE);

  for (int i = 0; i < ITERATIONS; ++i)
  {
    hipMutexEBOLock(mutex, gpuLockData->mutexBufferHeads, NUM_CU);

    accessData(storage, threadBaseLoc, threadOffset, NUM_LDST);

    hipMutexEBOUnlock(mutex, gpuLockData->mutexBufferHeads, NUM_CU);
  }
}

__global__ void kernelEBOMutexUniq(hipMutex_t mutex, float * storage,
                                   hipLockData_t * gpuLockData,
                                   const int ITERATIONS, const int NUM_LDST,
                                   const int NUM_CU)
{
  // local variables
  const int cuID = (hipBlockIdx_x % NUM_CU); // mod by # CUs to get CU ID
  // all work groups on the same CU access the same locations
  const int tid = ((cuID * hipBlockDim_x) + hipThreadIdx_x);
  // want striding to happen across cache lines so that each thread in a
  // half-warp accesses sequential words
  const int threadBaseLoc = ((tid/NUM_WORDS_PER_CACHELINE) * (NUM_LDST+1));
  const int threadOffset = (hipThreadIdx_x % NUM_WORDS_PER_CACHELINE);

  for (int i = 0; i < ITERATIONS; ++i)
  {
    hipMutexEBOLockLocal(mutex, cuID, gpuLockData->mutexBufferHeads, NUM_CU);

    accessData(storage, threadBaseLoc, threadOffset, NUM_LDST);

    hipMutexEBOUnlockLocal(mutex, cuID, gpuLockData->mutexBufferHeads, NUM_CU);
  }
}

// All WGs on all CUs access the same data with 1 writer per CU (and N-1)
// readers per CU.
__global__ void kernelSpinLockSemaphore(hipSemaphore_t sem,
                                        float * storage,
                                        hipLockData_t * gpuLockData,
                                        const unsigned int numStorageLocs,
                                        const int ITERATIONS,
                                        const int NUM_LDST,
                                        const int NUM_CU)
{
  // local variables
  const unsigned int maxSemCount =
    gpuLockData->semaphoreBuffers[((sem * 4 * NUM_CU) + (0 * 4))];
  const int cuID = (hipBlockIdx_x % NUM_CU); // mod by # CUs to get CU ID
  // If there are fewer WGs than # CUs, need to take into account for various
  // math below.  If WGs >= NUM_CU, use NUM_CU.
  const unsigned int numCU = ((hipGridDim_x < NUM_CU) ? hipGridDim_x : NUM_CU);
  // given the hipGridDim_x, we can figure out how many WGs are on our CU -- assume
  // all CUs have an identical number of WGs
  int numWGs_perCU = (int)ceil((float)hipGridDim_x / numCU);
  // number of threads on each WG
  //const int numThrs_perCU = (hipBlockDim_x * numWGs_perCU);
  const int perCU_WGID = (hipBlockIdx_x / numCU);
  // rotate which WG is the writer
  const bool isWriter = (perCU_WGID == (cuID % numWGs_perCU));

  // all work groups on the same CU access unique locations except the writer,
  // which writes all of the locations that all of the WGs access
  //int currWGID = hipBlockIdx_x;
  // the (reader) WGs on each CU access unique locations but those same
  // locations are accessed by the reader WGs on all CUs
  int tid = ((perCU_WGID * hipBlockDim_x) + hipThreadIdx_x);
  // want striding to happen across cache lines so that each thread in a
  // half-warp accesses sequential words
  int threadBaseLoc = ((tid/NUM_WORDS_PER_CACHELINE) * (NUM_LDST+1));
  const int threadOffset = (hipThreadIdx_x % NUM_WORDS_PER_CACHELINE);

  // dummy array to hold the loads done in the readers
  __shared__ volatile float dummyArray[NUM_WIS_PER_WG];

  for (int i = 0; i < ITERATIONS; ++i)
  {
    /*
      NOTE: There is a race here for entering the critical section.  Most
      importantly, it means that the at least one of the readers could win and
      thus the readers will read before the writer has had a chance to write
      the data.
    */
    hipSemaphoreSpinWait(sem, isWriter, maxSemCount,
                          gpuLockData->semaphoreBuffers, NUM_CU);

    // writer writes all the data that the WGs on this CU access
    if (isWriter) {
      for (int j = 0; j < numWGs_perCU; ++j) {
        /*
          Update the writer's "location" so it writes to the locations that the
          readers will access (due to RR scheduling the next WG on this CU is
          numCU WGs away).  Use loop counter because the non-unique version
          writes the same locations on all CUs.
        */
        tid = ((j * hipBlockDim_x) + hipThreadIdx_x);
        threadBaseLoc = ((tid/NUM_WORDS_PER_CACHELINE) * (NUM_LDST+1));

        accessData_semWr(storage, threadBaseLoc, threadOffset, NUM_LDST);
      }
      // reset locations
      tid = ((perCU_WGID * hipBlockDim_x) + hipThreadIdx_x);
      threadBaseLoc = ((tid/NUM_WORDS_PER_CACHELINE) * (NUM_LDST+1));
    }
    // rest of WGs on this CU read the data written by each CU's writer WG
    else {
      accessData_semRd(storage, dummyArray, threadBaseLoc, threadOffset,
                       NUM_LDST);
    }
    hipSemaphoreSpinPost(sem, isWriter, maxSemCount,
                          gpuLockData->semaphoreBuffers, NUM_CU);
  }
}

__global__ void kernelSpinLockSemaphoreUniq(hipSemaphore_t sem,
                                            float * storage,
                                            hipLockData_t * gpuLockData,
                                            const int ITERATIONS,
                                            const int NUM_LDST,
                                            const int NUM_CU)
{
  // local variables
  const unsigned int maxSemCount =
      gpuLockData->semaphoreBuffers[((sem * 4 * NUM_CU) + (0 * 4))];
  // If there are fewer WGs than # CUs, need to take into account for various
  // math below.  If WGs >= NUM_CU, use NUM_CU.
  const unsigned int numCU = ((hipGridDim_x < NUM_CU) ? hipGridDim_x : NUM_CU);
  const int cuID = (hipBlockIdx_x % numCU); // mod by # CUs to get CU ID
  // given the hipGridDim_x, we can figure out how many WGs are on our CU -- assume
  // all CUs have an identical number of WGs
  int numWGs_perCU = (int)ceil((float)hipGridDim_x / numCU);
  const int perCU_WGID = (hipBlockIdx_x / numCU);
  // rotate which WG is the writer
  const bool isWriter = (perCU_WGID == (cuID % numWGs_perCU));

  // all work groups on the same CU access unique locations except the writer,
  // which writes all of the locations that all of the WGs access
  int currWGID = hipBlockIdx_x;
  int tid = ((currWGID * hipBlockDim_x) + hipThreadIdx_x);
  // want striding to happen across cache lines so that each thread in a
  // half-warp accesses sequential words
  int threadBaseLoc = ((tid/NUM_WORDS_PER_CACHELINE) * (NUM_LDST+1));
  const int threadOffset = (hipThreadIdx_x % NUM_WORDS_PER_CACHELINE);
  // dummy array to hold the loads done in the readers
  __shared__ volatile float dummyArray[NUM_WIS_PER_WG];

  for (int i = 0; i < ITERATIONS; ++i)
  {
    /*
      NOTE: There is a race here for entering the critical section.  Most
      importantly, it means that the at least one of the readers could win and
      thus the readers will read before the writer has had a chance to write
      the data.
    */
    hipSemaphoreSpinWaitLocal(sem, cuID, isWriter, maxSemCount,
                               gpuLockData->semaphoreBuffers, NUM_CU);

    // writer WG writes all the data that the WGs on this CU access
    if (isWriter) {
      for (int j = 0; j < numWGs_perCU; ++j) {
        accessData_semWr(storage, threadBaseLoc, threadOffset, NUM_LDST);

        /*
          update the writer's "location" so it writes to the locations that the
          readers will access (due to RR scheduling the next WG on this CU is
          numCU WGs away and < hipGridDim_x).

          NOTE: First location writer writes to is its own location(s).  If the
          writer is not CU 0 on this CU, it may require wrapping around to CUs
          with smaller WG IDs.
        */
        currWGID = (currWGID + numCU) % hipGridDim_x;
        tid = ((currWGID * hipBlockDim_x) + hipThreadIdx_x);
        threadBaseLoc = ((tid/NUM_WORDS_PER_CACHELINE) * (NUM_LDST+1));
      }
      // reset locations
      currWGID = hipBlockIdx_x;
      tid = ((currWGID * hipBlockDim_x) + hipThreadIdx_x);
      threadBaseLoc = ((tid/NUM_WORDS_PER_CACHELINE) * (NUM_LDST+1));
    }
    // rest of WGs on this CU read the data written by each CU's writer WG
    else {
      accessData_semRd(storage, dummyArray, threadBaseLoc, threadOffset,
                       NUM_LDST);
    }
    hipSemaphoreSpinPostLocal(sem, cuID, isWriter, maxSemCount,
                               gpuLockData->semaphoreBuffers, NUM_CU);
  }
}

// All WGs on all CUs access the same data with 1 writer per CU (and N-1)
// readers per CU.
__global__ void kernelEBOSemaphore(hipSemaphore_t sem, float * storage,
                                   hipLockData_t * gpuLockData,
                                   const unsigned int numStorageLocs,
                                   const int ITERATIONS, const int NUM_LDST,
                                   const int NUM_CU)
{
  // local variables
  const unsigned int maxSemCount =
      gpuLockData->semaphoreBuffers[((sem * 4 * NUM_CU) + (0 * 4))];
  // If there are fewer WGs than # CUs, need to take into account for various
  // math below.  If WGs >= NUM_CU, use NUM_CU.
  const unsigned int numCU = ((hipGridDim_x < NUM_CU) ? hipGridDim_x : NUM_CU);
  const int cuID = (hipBlockIdx_x % NUM_CU); // mod by # CUs to get CU ID
  // given the hipGridDim_x, we can figure out how many WGs are on our CU -- assume
  // all CUs have an identical number of WGs
  int numWGs_perCU = (int)ceil((float)hipGridDim_x / numCU);
  // number of threads on each WG
  //const int numThrs_perCU = (hipBlockDim_x * numWGs_perCU);
  const int perCU_WGID = (hipBlockIdx_x / numCU);
  // rotate which WG is the writer
  const bool isWriter = (perCU_WGID == (cuID % numWGs_perCU));

  // all work groups on the same CU access unique locations except the writer,
  // which writes all of the locations that all of the WGs access
  //int currWGID = hipBlockIdx_x;
  // the (reader) WGs on each CU access unique locations but those same
  // locations are accessed by the reader WGs on all CUs
  int tid = ((perCU_WGID * hipBlockDim_x) + hipThreadIdx_x);
  // want striding to happen across cache lines so that each thread in a
  // half-warp accesses sequential words
  int threadBaseLoc = ((tid/NUM_WORDS_PER_CACHELINE) * (NUM_LDST+1));
  const int threadOffset = (hipThreadIdx_x % NUM_WORDS_PER_CACHELINE);

  // dummy array to hold the loads done in the readers
  __shared__ volatile float dummyArray[NUM_WIS_PER_WG];

  for (int i = 0; i < ITERATIONS; ++i)
  {
    /*
      NOTE: There is a race here for entering the critical section.  Most
      importantly, it means that the at least one of the readers could win and
      thus the readers will read before the writer has had a chance to write
      the data.
    */
   hipSemaphoreEBOWait(sem, isWriter, maxSemCount,
                        gpuLockData->semaphoreBuffers, NUM_CU);

    // writer WG writes all the data that the WGs on this CU access
    if (isWriter) {
      for (int j = 0; j < numWGs_perCU; ++j) {
        /*
          Update the writer's "location" so it writes to the locations that the
          readers will access (due to RR scheduling the next WG on this CU is
          numCU WGs away).  Use loop counter because the non-unique version
          writes the same locations on all CUs.
        */
        tid = ((j * hipBlockDim_x) + hipThreadIdx_x);
        threadBaseLoc = ((tid/NUM_WORDS_PER_CACHELINE) * (NUM_LDST+1));

        accessData_semWr(storage, threadBaseLoc, threadOffset, NUM_LDST);
      }
      // reset locations
      tid = ((perCU_WGID * hipBlockDim_x) + hipThreadIdx_x);
      threadBaseLoc = ((tid/NUM_WORDS_PER_CACHELINE) * (NUM_LDST+1));
    }
    // rest of WGs on this CU read the data written by each CU's writer WG
    else {
      accessData_semRd(storage, dummyArray, threadBaseLoc, threadOffset,
                       NUM_LDST);
    }
    hipSemaphoreEBOPost(sem, isWriter, maxSemCount,
                         gpuLockData->semaphoreBuffers, NUM_CU);
  }
}

__global__ void kernelEBOSemaphoreUniq(hipSemaphore_t sem, float * storage,
                                       hipLockData_t * gpuLockData,
                                       const int ITERATIONS,
                                       const int NUM_LDST,
                                       const int NUM_CU)
{
  // local variables
  const unsigned int maxSemCount =
      gpuLockData->semaphoreBuffers[((sem * 4 * NUM_CU) + (0 * 4))];
  // If there are fewer WGs than # CUs, need to take into account for various
  // math below.  If WGs >= NUM_CU, use NUM_CU.
  const unsigned int numCU = ((hipGridDim_x < NUM_CU) ? hipGridDim_x : NUM_CU);
  const int cuID = (hipBlockIdx_x % numCU); // mod by # CUs to get CU ID
  // given the hipGridDim_x, we can figure out how many WGs are on our CU -- assume
  // all CUs have an identical number of WGs
  int numWGs_perCU = (int)ceil((float)hipGridDim_x / numCU);
  const int perCU_WGID = (hipBlockIdx_x / numCU);
  // rotate which WG is the writer
  const bool isWriter = (perCU_WGID == (cuID % numWGs_perCU));

  // all work groups on the same CU access unique locations except the writer,
  // which writes all of the locations that all of the WGs access
  int currWGID = hipBlockIdx_x;
  int tid = ((currWGID * hipBlockDim_x) + hipThreadIdx_x);
  // want striding to happen across cache lines so that each thread in a
  // half-warp accesses sequential words
  int threadBaseLoc = ((tid/NUM_WORDS_PER_CACHELINE) * (NUM_LDST+1));
  const int threadOffset = (hipThreadIdx_x % NUM_WORDS_PER_CACHELINE);
  // dummy array to hold the loads done in the readers
  __shared__ volatile float dummyArray[NUM_WIS_PER_WG];

  for (int i = 0; i < ITERATIONS; ++i)
  {
    /*
      NOTE: There is a race here for entering the critical section.  Most
      importantly, it means that the at least one of the readers could win and
      thus the readers will read before the writer has had a chance to write
      the data.
    */
    hipSemaphoreEBOWaitLocal(sem, cuID, isWriter, maxSemCount,
                              gpuLockData->semaphoreBuffers, NUM_CU);

    // writer WG writes all the data that the WGs on this CU access
    if (isWriter) {
      for (int j = 0; j < numWGs_perCU; ++j) {
        accessData_semWr(storage, threadBaseLoc, threadOffset, NUM_LDST);

        /*
          update the writer's "location" so it writes to the locations that the
          readers will access (due to RR scheduling the next WG on this CU is
          numCU WGs away and < hipGridDim_x).

          NOTE: First location writer writes to is its own location(s).  If the
          writer is not CU 0 on this CU, it may require wrapping around to CUs
          with smaller WG IDs.
        */
        currWGID = (currWGID + numCU) % hipGridDim_x;
        tid = ((currWGID * hipBlockDim_x) + hipThreadIdx_x);
        threadBaseLoc = ((tid/NUM_WORDS_PER_CACHELINE) * (NUM_LDST+1));
      }
      // reset locations
      currWGID = hipBlockIdx_x;
      tid = ((currWGID * hipBlockDim_x) + hipThreadIdx_x);
      threadBaseLoc = ((tid/NUM_WORDS_PER_CACHELINE) * (NUM_LDST+1));
    }
    // rest of WGs on this CU read the data written by each CU's writer WG
    else {
      accessData_semRd(storage, dummyArray, threadBaseLoc, threadOffset,
                       NUM_LDST);
    }
    hipSemaphoreEBOPostLocal(sem, cuID, isWriter, maxSemCount,
                              gpuLockData->semaphoreBuffers, NUM_CU);
  }
}

void invokeAtomicTreeBarrier(float * storage_d, unsigned int * perCUBarriers_d,
                             int numIters)
{
  // local variable
  const int WGs = numWGs;

  for (int repeat = 0; repeat < NUM_REPEATS; ++repeat)
  {
    hipLaunchKernelGGL(HIP_KERNEL_NAME(kernelAtomicTreeBarrierUniq), dim3(WGs), dim3(NUM_WIS_PER_WG), 0, 0,
        storage_d, cpuLockData, perCUBarriers_d, numIters, NUM_LDST, NUM_CU,
        MAX_WGS);

    // Blocks until the device has completed all preceding requested
    // tasks (make sure that the device returned before continuing).
    hipError_t hipErr = hipDeviceSynchronize();
    checkError(hipErr, "hipDeviceSynchronize (kernelAtomicTreeBarrierUniq)");
  }
}

void invokeAtomicTreeBarrierLocalExch(float * storage_d,
                                      unsigned int * perCUBarriers_d,
                                      int numIters)
{
  // local variable
  const int WGs = numWGs;

  for (int repeat = 0; repeat < NUM_REPEATS; ++repeat)
  {
    hipLaunchKernelGGL(HIP_KERNEL_NAME(kernelAtomicTreeBarrierUniqLocalExch), dim3(WGs), dim3(NUM_WIS_PER_WG), 0, 0,
        storage_d, cpuLockData, perCUBarriers_d, numIters, NUM_LDST, NUM_CU,
        MAX_WGS);

    // Blocks until the device has completed all preceding requested
    // tasks (make sure that the device returned before continuing).
    hipError_t hipErr = hipDeviceSynchronize();
    checkError(hipErr,
               "hipDeviceSynchronize (kernelAtomicTreeBarrierUniqLockExch)");
  }
}

void invokeFBSTreeBarrier(float * storage_d, unsigned int * perCUBarriers_d,
                          int numIters)
{
  // local variable
  const int WGs = numWGs;

  for (int repeat = 0; repeat < NUM_REPEATS; ++repeat)
  {
    hipLaunchKernelGGL(HIP_KERNEL_NAME(kernelFBSTreeBarrierUniq), dim3(WGs), dim3(NUM_WIS_PER_WG), 0, 0,
        storage_d, cpuLockData, perCUBarriers_d, numIters, NUM_LDST, NUM_CU,
        MAX_WGS);

    // Blocks until the device has completed all preceding requested
    // tasks (make sure that the device returned before continuing).
    hipError_t hipErr = hipDeviceSynchronize();
    checkError(hipErr, "hipDeviceSynchronize (kernelFBSTreeBarrierUniq)");
  }
}

void invokeFBSTreeBarrierLocalExch(float * storage_d,
                                   unsigned int * perCUBarriers_d,
                                   int numIters)
{
  // local variable
  const int WGs = numWGs;

  for (int repeat = 0; repeat < NUM_REPEATS; ++repeat)
  {
    hipLaunchKernelGGL(HIP_KERNEL_NAME(kernelFBSTreeBarrierUniqLocalExch), dim3(WGs), dim3(NUM_WIS_PER_WG), 0, 0,
        storage_d, cpuLockData, perCUBarriers_d, numIters, NUM_LDST, NUM_CU,
        MAX_WGS);

    // Blocks until the device has completed all preceding requested
    // tasks (make sure that the device returned before continuing).
    hipError_t hipErr = hipDeviceSynchronize();
    checkError(hipErr, "hipDeviceSynchronize (kernelFBSTreeBarrierUniqLocalExch)");
  }
}

void invokeSpinLockMutex(hipMutex_t mutex, float * storage_d, int numIters)
{
  // local variable
  const int WGs = numWGs;

  for (int repeat = 0; repeat < NUM_REPEATS; ++repeat)
  {
    hipLaunchKernelGGL(HIP_KERNEL_NAME(kernelSpinLockMutex), dim3(1, 1, 1), dim3(NUM_WIS_PER_WG), 0, 0,
        mutex, storage_d, cpuLockData, numIters, NUM_LDST, NUM_CU);

    // Blocks until the device has completed all preceding requested
    // tasks (make sure that the device returned before continuing).
    hipError_t hipErr = hipDeviceSynchronize();
    checkError(hipErr, "hipDeviceSynchronize (kernelSpinLockMutex)");
  }
}

void invokeSpinLockMutex_uniq(hipMutex_t mutex, float * storage_d,
                              int numIters)
{
  // local variable
  const int WGs = numWGs;

  for (int repeat = 0; repeat < NUM_REPEATS; ++repeat)
  {
    hipLaunchKernelGGL(HIP_KERNEL_NAME(kernelSpinLockMutexUniq), dim3(WGs), dim3(NUM_WIS_PER_WG), 0, 0,
        mutex, storage_d, cpuLockData, numIters, NUM_LDST, NUM_CU);

    // Blocks until the device has completed all preceding requested
    // tasks (make sure that the device returned before continuing).
    hipError_t hipErr = hipDeviceSynchronize();
    checkError(hipErr, "hipDeviceSynchronize (kernelSpinLockMutexUniq)");
  }
}

void invokeEBOMutex(hipMutex_t mutex, float * storage_d, int numIters)
{
  // local variable
  const int WGs = numWGs;

  for (int repeat = 0; repeat < NUM_REPEATS; ++repeat)
  {
    hipLaunchKernelGGL(HIP_KERNEL_NAME(kernelEBOMutex), dim3(WGs), dim3(NUM_WIS_PER_WG), 0, 0,
        mutex, storage_d, cpuLockData, numIters, NUM_LDST, NUM_CU);

    // Blocks until the device has completed all preceding requested
    // tasks (make sure that the device returned before continuing).
    hipError_t hipErr = hipDeviceSynchronize();
    checkError(hipErr, "hipDeviceSynchronize (kernelEBOMutex)");
  }
}

void invokeEBOMutex_uniq(hipMutex_t mutex, float * storage_d, int numIters)
{
  // local variable
  const int WGs = numWGs;

  for (int repeat = 0; repeat < NUM_REPEATS; ++repeat)
  {
    hipLaunchKernelGGL(HIP_KERNEL_NAME(kernelEBOMutexUniq), dim3(WGs), dim3(NUM_WIS_PER_WG), 0, 0,
        mutex, storage_d, cpuLockData, numIters, NUM_LDST, NUM_CU);

    // Blocks until the device has completed all preceding requested
    // tasks (make sure that the device returned before continuing).
    hipError_t hipErr = hipDeviceSynchronize();
    checkError(hipErr, "hipDeviceSynchronize (kernelEBOMutexUniq)");
  }
}

void invokeSleepingMutex(hipMutex_t mutex, float * storage_d, int numIters)
{
  // local variable
  const int WGs = numWGs;

  for (int repeat = 0; repeat < NUM_REPEATS; ++repeat)
  {
    hipLaunchKernelGGL(HIP_KERNEL_NAME(kernelSleepingMutex), dim3(WGs), dim3(NUM_WIS_PER_WG), 0, 0,
                       mutex, storage_d, cpuLockData, numIters, NUM_LDST, NUM_CU);

    // Blocks until the device has completed all preceding requested
    // tasks (make sure that the device returned before continuing).
    hipError_t hipErr = hipDeviceSynchronize();
    checkError(hipErr, "hipDeviceSynchronize (kernelSleepingMutex)");
  }
}

void invokeSleepingMutex_uniq(hipMutex_t mutex, float * storage_d,
                              int numIters)
{
  // local variable
  const int WGs = numWGs;

  for (int repeat = 0; repeat < NUM_REPEATS; ++repeat)
  {
    hipLaunchKernelGGL(HIP_KERNEL_NAME(kernelSleepingMutexUniq), dim3(WGs), dim3(NUM_WIS_PER_WG), 0, 0,
        mutex, storage_d, cpuLockData, numIters, NUM_LDST, NUM_CU);

    // Blocks until the device has completed all preceding requested
    // tasks (make sure that the device returned before continuing).
    hipError_t hipErr = hipDeviceSynchronize();
    checkError(hipErr, "hipDeviceSynchronize (kernelSleepingMutexUniq)");
  }
}

void invokeFetchAndAddMutex(hipMutex_t mutex, float * storage_d, int numIters)
{
  // local variable
  const int WGs = numWGs;

  for (int repeat = 0; repeat < NUM_REPEATS; ++repeat)
  {
    hipLaunchKernelGGL(HIP_KERNEL_NAME(kernelFetchAndAddMutex), dim3(WGs), dim3(NUM_WIS_PER_WG), 0, 0,
        mutex, storage_d, cpuLockData, numIters, NUM_LDST, NUM_CU);

    // Blocks until the device has completed all preceding requested
    // tasks (make sure that the device returned before continuing).
    hipError_t hipErr = hipDeviceSynchronize();
    checkError(hipErr, "hipDeviceSynchronize (kernelFetchAndAddMutex)");
  }
}

void invokeFetchAndAddMutex_uniq(hipMutex_t mutex, float * storage_d, int numIters)
{
  // local variable
  const int WGs = numWGs;

  for (int repeat = 0; repeat < NUM_REPEATS; ++repeat)
  {
    hipLaunchKernelGGL(HIP_KERNEL_NAME(kernelFetchAndAddMutexUniq), dim3(WGs), dim3(NUM_WIS_PER_WG), 0, 0,
        mutex, storage_d, cpuLockData, numIters, NUM_LDST, NUM_CU);

    // Blocks until the device has completed all preceding requested
    // tasks (make sure that the device returned before continuing).
    hipError_t hipErr = hipDeviceSynchronize();
    checkError(hipErr, "hipDeviceSynchronize (kernelFetchAndAddMutexUniq)");
  }
}

void invokeSpinLockSemaphore(hipSemaphore_t sem, float * storage_d,
                             const int maxVal,
                             int numIters, int numStorageLocs)
{
  // local variable
  const int WGs = numWGs;

  for (int repeat = 0; repeat < NUM_REPEATS; ++repeat)
  {
    hipLaunchKernelGGL(HIP_KERNEL_NAME(kernelSpinLockSemaphore), dim3(WGs), dim3(NUM_WIS_PER_WG), 0, 0,
        sem, storage_d, cpuLockData, numStorageLocs, numIters, NUM_LDST,
        NUM_CU);

    // Blocks until the device has completed all preceding requested
    // tasks (make sure that the device returned before continuing).
    hipError_t hipErr = hipDeviceSynchronize();
    checkError(hipErr, "hipDeviceSynchronize (kernelSpinLockSemaphore)");
  }
}

void invokeSpinLockSemaphore_uniq(hipSemaphore_t sem, float * storage_d,
                                  const int maxVal, int numIters)
{
  // local variable
  const int WGs = numWGs;

  for (int repeat = 0; repeat < NUM_REPEATS; ++repeat)
  {
    hipLaunchKernelGGL(HIP_KERNEL_NAME(kernelSpinLockSemaphoreUniq), dim3(WGs), dim3(NUM_WIS_PER_WG), 0, 0,
        sem, storage_d, cpuLockData, numIters, NUM_LDST, NUM_CU);

    // Blocks until the device has completed all preceding requested
    // tasks (make sure that the device returned before continuing).
    hipError_t hipErr = hipDeviceSynchronize();
    checkError(hipErr, "hipDeviceSynchronize (kernelSpinLockSemaphoreUniq)");
  }
}

void invokeEBOSemaphore(hipSemaphore_t sem, float * storage_d, const int maxVal,
                        int numIters, int numStorageLocs)
{
  // local variable
  const int WGs = numWGs;

  for (int repeat = 0; repeat < NUM_REPEATS; ++repeat)
  {
    hipLaunchKernelGGL(HIP_KERNEL_NAME(kernelEBOSemaphore), dim3(WGs), dim3(NUM_WIS_PER_WG), 0, 0,
        sem, storage_d, cpuLockData, numStorageLocs, numIters, NUM_LDST,
        NUM_CU);

    // Blocks until the device has completed all preceding requested
    // tasks (make sure that the device returned before continuing).
    hipError_t hipErr = hipDeviceSynchronize();
    checkError(hipErr, "hipDeviceSynchronize (kernelEBOSemaphore)");
  }
}

void invokeEBOSemaphore_uniq(hipSemaphore_t sem, float * storage_d,
                             const int maxVal, int numIters)
{
  // local variable
  const int WGs = numWGs;

  for (int repeat = 0; repeat < NUM_REPEATS; ++repeat)
  {
    hipLaunchKernelGGL(HIP_KERNEL_NAME(kernelEBOSemaphoreUniq), dim3(WGs), dim3(NUM_WIS_PER_WG), 0, 0,
        sem, storage_d, cpuLockData, numIters, NUM_LDST, NUM_CU);

    // Blocks until the device has completed all preceding requested
    // tasks (make sure that the device returned before continuing).
    hipError_t hipErr = hipDeviceSynchronize();
    checkError(hipErr, "hipDeviceSynchronize (kernelEBOSemaphoreUniq)");
  }
}

int main(int argc, char ** argv)
{
  if (argc != 5) {
    fprintf(stderr, "./allSyncPrims-1kernel <syncPrim> <numLdSt> <numWGs> "
            "<numCSIters>\n");
    fprintf(stderr, "where:\n");
    fprintf(stderr, "\t<syncPrim>: a string that represents which synchronization primitive to run.\n"
            "\t\tatomicTreeBarrUniq - Atomic Tree Barrier\n"
            "\t\tatomicTreeBarrUniqLocalExch - Atomic Tree Barrier with local exchange\n"
            "\t\tlfTreeBarrUniq - Lock-Free Tree Barrier\n"
            "\t\tlfTreeBarrUniqLocalExch - Lock-Free Tree Barrier with local exchange\n"
            "\t\tspinMutex - Spin Lock Mutex\n"
            "\t\tspinMutexEBO - Spin Lock Mutex with Backoff\n"
            "\t\tsleepMutex - Sleep Mutex\n"
            "\t\tfaMutex - Fetch-and-Add Mutex\n"
            "\t\tspinSem1 - Spin Semaphore (Max: 1)\n"
            "\t\tspinSem2 - Spin Semaphore (Max: 2)\n"
            "\t\tspinSem10 - Spin Semaphore (Max: 10)\n"
            "\t\tspinSem120 - Spin Semaphore (Max: 120)\n"
            "\t\tspinSemEBO1 - Spin Semaphore with Backoff (Max: 1)\n"
            "\t\tspinSemEBO2 - Spin Semaphore with Backoff (Max: 2)\n"
            "\t\tspinSemEBO10 - Spin Semaphore with Backoff (Max: 10)\n"
            "\t\tspinSemEBO120 - Spin Semaphore with Backoff (Max: 120)\n"
            "\t\tspinMutexUniq - Spin Lock Mutex -- accesses to unique locations per WG\n"
            "\t\tspinMutexEBOUniq - Spin Lock Mutex with Backoff -- accesses to unique locations per WG\n"
            "\t\tsleepMutexUniq - Sleep Mutex -- accesses to unique locations per WG\n"
            "\t\tfaMutexUniq - Fetch-and-Add Mutex -- accesses to unique locations per WG\n"
            "\t\tspinSemUniq1 - Spin Semaphore (Max: 1) -- accesses to unique locations per WG\n"
            "\t\tspinSemUniq2 - Spin Semaphore (Max: 2) -- accesses to unique locations per WG\n"
            "\t\tspinSemUniq10 - Spin Semaphore (Max: 10) -- accesses to unique locations per WG\n"
            "\t\tspinSemUniq120 - Spin Semaphore (Max: 120) -- accesses to unique locations per WG\n"
            "\t\tspinSemEBOUniq1 - Spin Semaphore with Backoff (Max: 1) -- accesses to unique locations per WG\n"
            "\t\tspinSemEBOUniq2 - Spin Semaphore with Backoff (Max: 2) -- accesses to unique locations per WG\n"
            "\t\tspinSemEBOUniq10 - Spin Semaphore with Backoff (Max: 10) -- accesses to unique locations per WG\n"
            "\t\tspinSemEBOUniq120 - Spin Semaphore with Backoff (Max: 120) -- accesses to unique locations per WG\n");
    fprintf(stderr, "\t<numLdSt>: the # of LDs and STs to do for each thread "
            "in the critical section.\n");
    fprintf(stderr, "\t<numWGs>: the # of WGs to execute (want to be "
            "divisible by the number of CUs).\n");
    fprintf(stderr, "\t<numCSIters>: number of iterations of the critical "
            "section.\n");
    exit(-1);
  }

  // boilerplate code to identify compute capability, # CU/CUM/CUX, etc.
  int deviceCount;
  hipGetDeviceCount(&deviceCount);
  if (deviceCount == 0) {
    fprintf(stderr, "There is no device supporting HIP\n");
    exit(-1);
  }

  hipDeviceProp_t deviceProp;
  hipGetDeviceProperties(&deviceProp, 0);
  fprintf(stdout, "GPU Compute Capability: %d.%d\n", deviceProp.major,
          deviceProp.minor);
  if ((deviceProp.major == 9999) && (deviceProp.minor == 9999)) {
    fprintf(stderr, "There is no HIP capable device\n");
    exit(-1);
  }

  NUM_CU = deviceProp.multiProcessorCount;
  const int maxWGPerCU = deviceProp.maxThreadsPerBlock/NUM_WIS_PER_WG;
  //assert(maxWGPerCU * NUM_WIS_PER_WG <=
  //       deviceProp.maxThreadsPerMultiProcessor);
  MAX_WGS = maxWGPerCU * NUM_CU;

  fprintf(stdout, "# CU: %d, Max Thrs/WG: %d, Max WG/CU: %d, Max # WG: %d\n",
          NUM_CU, deviceProp.maxThreadsPerBlock, maxWGPerCU, MAX_WGS);

  hipError_t hipErr = hipGetLastError();
  checkError(hipErr, "Begin");

  // parse input args
  const char * syncPrim_str = argv[1];
  NUM_LDST = atoi(argv[2]);
  numWGs = atoi(argv[3]);
  assert(numWGs <= MAX_WGS);
  const int NUM_ITERS = atoi(argv[4]);
  const int numWGs_perCU = (int)ceil((float)numWGs / NUM_CU);
  assert(numWGs_perCU > 0);

  unsigned int syncPrim = 9999;
  // set the syncPrim variable to the appropriate value based on the inputted
  // string for the microbenchmark
  if (strcmp(syncPrim_str, "atomicTreeBarrUniq") == 0) { syncPrim = 0; }
  else if (strcmp(syncPrim_str, "atomicTreeBarrUniqLocalExch") == 0) {
    syncPrim = 1;
  }
  else if (strcmp(syncPrim_str, "lfTreeBarrUniq") == 0) { syncPrim = 2; }
  else if (strcmp(syncPrim_str, "lfTreeBarrUniqLocalExch") == 0) {
    syncPrim = 3;
  }
  else if (strcmp(syncPrim_str, "spinMutex") == 0) { syncPrim = 4; }
  else if (strcmp(syncPrim_str, "spinMutexEBO") == 0) { syncPrim = 5; }
  else if (strcmp(syncPrim_str, "sleepMutex") == 0) { syncPrim = 6; }
  else if (strcmp(syncPrim_str, "faMutex") == 0) { syncPrim = 7; }
  else if (strcmp(syncPrim_str, "spinSem1") == 0) { syncPrim = 8; }
  else if (strcmp(syncPrim_str, "spinSem2") == 0) { syncPrim = 9; }
  else if (strcmp(syncPrim_str, "spinSem10") == 0) { syncPrim = 10; }
  else if (strcmp(syncPrim_str, "spinSem120") == 0) { syncPrim = 11; }
  else if (strcmp(syncPrim_str, "spinSemEBO1") == 0) { syncPrim = 12; }
  else if (strcmp(syncPrim_str, "spinSemEBO2") == 0) { syncPrim = 13; }
  else if (strcmp(syncPrim_str, "spinSemEBO10") == 0) { syncPrim = 14; }
  else if (strcmp(syncPrim_str, "spinSemEBO120") == 0) { syncPrim = 15; }
  // cases 16-19 reserved
  else if (strcmp(syncPrim_str, "spinMutexUniq") == 0) { syncPrim = 20; }
  else if (strcmp(syncPrim_str, "spinMutexEBOUniq") == 0) { syncPrim = 21; }
  else if (strcmp(syncPrim_str, "sleepMutexUniq") == 0) { syncPrim = 22; }
  else if (strcmp(syncPrim_str, "faMutexUniq") == 0) { syncPrim = 23; }
  else if (strcmp(syncPrim_str, "spinSemUniq1") == 0) { syncPrim = 24; }
  else if (strcmp(syncPrim_str, "spinSemUniq2") == 0) { syncPrim = 25; }
  else if (strcmp(syncPrim_str, "spinSemUniq10") == 0) { syncPrim = 26; }
  else if (strcmp(syncPrim_str, "spinSemUniq120") == 0) { syncPrim = 27; }
  else if (strcmp(syncPrim_str, "spinSemEBOUniq1") == 0) { syncPrim = 28; }
  else if (strcmp(syncPrim_str, "spinSemEBOUniq2") == 0) { syncPrim = 29; }
  else if (strcmp(syncPrim_str, "spinSemEBOUniq10") == 0) { syncPrim = 30; }
  else if (strcmp(syncPrim_str, "spinSemEBOUniq120") == 0) { syncPrim = 31; }
  // cases 32-36 reserved
  else
  {
    fprintf(stderr, "ERROR: Unknown synchronization primitive: %s\n",
            syncPrim_str);
    exit(-1);
  }

  // multiply number of mutexes, semaphores by NUM_CU to
  // allow per-core locks
  hipLocksInit(MAX_WGS, 8 * NUM_CU, 24 * NUM_CU, pageAlign, NUM_CU, NUM_REPEATS, NUM_ITERS);

  hipErr = hipGetLastError();
  checkError(hipErr, "After hipLocksInit");

  /*
    The barriers need a per-CU barrier that is not part of the global synch
    structure.  In terms of size, for the lock-free barrier there are 2 arrays
    in here -- inVars and outVars.  Each needs to be sized to hold the maximum
    number of WGs/CU and each CU needs an array.

    The atomic barrier per-CU synchronization fits inside the lock-free size
    requirements so we can reuse the same locations.
  */
  unsigned int * perCUBarriers = (unsigned int *)malloc(sizeof(unsigned int) * (NUM_CU * MAX_WGS * 2));

  int numLocsMult = 0;
  // barriers and unique semaphores have numWGs WGs accessing unique locations
  if ((syncPrim < 4) ||
      ((syncPrim >= 24) && (syncPrim <= 35))) { numLocsMult = numWGs; }
  // The non-unique mutex microbenchmarks, all WGs access the same locations so
  // multiplier is 1
  else if ((syncPrim >= 4) && (syncPrim <= 7)) { numLocsMult = 1; }
  // The non-unique semaphores have 1 writer and numWGs_perCU - 1 readers per CU
  // so the multiplier is numWGs_perCU
  else if ((syncPrim >= 8) && (syncPrim <= 19)) { numLocsMult = numWGs_perCU; }
  // For the unique mutex microbenchmarks and condition variable, all WGs on
  // same CU access same data so multiplier is NUM_CU.
  else if (((syncPrim >= 20) && (syncPrim <= 23)) ||
           (syncPrim == 36)) { numLocsMult = ((numWGs < NUM_CU) ?
                                              numWGs : NUM_CU); }
  else { // should never reach here
    fprintf(stderr, "ERROR: Unknown syncPrim: %u\n", syncPrim);
    exit(-1);
  }

  // each thread in a WG accesses NUM_LDST locations but accesses
  // per WI are offset so that each subsequent access is dependent
  // on the previous one -- thus need an extra access per WI.
  int numUniqLocsAccPerWG = (NUM_WIS_PER_WG * (NUM_LDST + 1));
  assert(numUniqLocsAccPerWG > 0);
  int numStorageLocs = (numLocsMult * numUniqLocsAccPerWG);
  assert(numStorageLocs > 0);
  float * storage = (float *)malloc(sizeof(float) * numStorageLocs);

  fprintf(stdout, "# WGs: %d, # Ld/St: %d, # Locs Mult: %d, # Uniq Locs/WG: %d, # Storage Locs: %d\n", numWGs, NUM_LDST, numLocsMult, numUniqLocsAccPerWG, numStorageLocs);

  // initialize storage
  for (int i = 0; i < numStorageLocs; ++i) { storage[i] = i; }
  // initialize per-CU barriers to 0's
  for (int i = 0; i < (NUM_CU * MAX_WGS * 2); ++i) { perCUBarriers[i] = 0; }

  // gpu copies of storage and perCUBarriers
  //float elapsedTime = 0.0f;
  unsigned int * perCUBarriers_d = NULL;
  float * storage_d = NULL;

  hipMalloc(&perCUBarriers_d, sizeof(unsigned int) * (NUM_CU * MAX_WGS * 2));
  hipMalloc(&storage_d, sizeof(float) * numStorageLocs);

  hipMemcpy(perCUBarriers_d, perCUBarriers, sizeof(unsigned int) * (NUM_CU * MAX_WGS * 2), hipMemcpyHostToDevice);
  hipMemcpy(storage_d, storage, sizeof(float) * numStorageLocs, hipMemcpyHostToDevice);

  // lock variables
  hipMutex_t spinMutex, faMutex, sleepMutex, eboMutex;
  hipMutex_t spinMutex_uniq, faMutex_uniq, sleepMutex_uniq, eboMutex_uniq;
  hipSemaphore_t spinSem1, eboSem1,
                  spinSem2, eboSem2,
                  spinSem10, eboSem10,
                  spinSem120, eboSem120;
  hipSemaphore_t spinSem1_uniq, eboSem1_uniq,
                  spinSem2_uniq, eboSem2_uniq,
                  spinSem10_uniq, eboSem10_uniq,
                  spinSem120_uniq, eboSem120_uniq;
  switch (syncPrim) {
    case 0: // atomic tree barrier doesn't require any special fields to be
            // created
      printf("atomic_tree_barrier_%03d\n", NUM_ITERS); fflush(stdout);
      break;
    case 1: // atomic tree barrier with local exchange doesn't require any
            // special fields to be created
      printf("atomic_tree_barrier_localExch_%03d\n", NUM_ITERS); fflush(stdout);
      break;
    case 2: // lock-free tree barrier doesn't require any special fields to be
            // created
      printf("fbs_tree_barrier_%03d\n", NUM_ITERS); fflush(stdout);
      break;
    case 3: // lock-free barrier with local exchange doesn't require any
            // special fields to be created
      printf("fbs_tree_barrier_localExch_%03d\n", NUM_ITERS); fflush(stdout);
      break;
    case 4:
      printf("spin_lock_mutex_%03d\n", NUM_ITERS); fflush(stdout);
      hipMutexCreateSpin     (&spinMutex,          0);
      break;
    case 5:
      printf("ebo_mutex_%03d\n", NUM_ITERS); fflush(stdout);
      hipMutexCreateEBO      (&eboMutex,           1);
      break;
    case 6:
      printf("sleeping_mutex_%03d\n", NUM_ITERS); fflush(stdout);
      hipMutexCreateSleep    (&sleepMutex,         2);
      break;
    case 7:
      printf("fetchadd_mutex_%03d\n", NUM_ITERS); fflush(stdout);
      hipMutexCreateFA       (&faMutex,            3);
      break;
    case 8:
      printf("spin_lock_sem_%03d_%03d\n", 1, NUM_ITERS); fflush(stdout);
      hipSemaphoreCreateSpin (&spinSem1,      0,   1, NUM_CU);
      break;
    case 9:
      printf("spin_lock_sem_%03d_%03d\n", 2, NUM_ITERS); fflush(stdout);
      hipSemaphoreCreateSpin (&spinSem2,      1,   2, NUM_CU);
      break;
    case 10:
      printf("spin_lock_sem_%03d_%03d\n", 10, NUM_ITERS); fflush(stdout);
      hipSemaphoreCreateSpin (&spinSem10,      2,   10, NUM_CU);
      break;
    case 11:
      printf("spin_lock_sem_%03d_%03d\n", 2, NUM_ITERS); fflush(stdout);
      hipSemaphoreCreateSpin (&spinSem120,      3,   120, NUM_CU);
      break;
    case 12:
      printf("ebo_sem_%03d_%03d\n", 1, NUM_ITERS); fflush(stdout);
      hipSemaphoreCreateEBO  (&eboSem1,       4,   1, NUM_CU);
      break;
    case 13:
      printf("ebo_sem_%03d_%03d\n", 2, NUM_ITERS); fflush(stdout);
      hipSemaphoreCreateEBO  (&eboSem2,       5,   2, NUM_CU);
      break;
    case 14:
      printf("ebo_sem_%03d_%03d\n", 10, NUM_ITERS); fflush(stdout);
      hipSemaphoreCreateEBO  (&eboSem10,       6,   10, NUM_CU);
      break;
    case 15:
      printf("ebo_sem_%03d_%03d\n", 120, NUM_ITERS); fflush(stdout);
      hipSemaphoreCreateEBO  (&eboSem120,       7,   120, NUM_CU);
      break;
    // cases 16-19 reserved
    case 16:
      break;
    case 17:
      break;
    case 18:
      break;
    case 19:
      break;
    case 20:
      printf("spin_lock_mutex_uniq_%03d\n", NUM_ITERS); fflush(stdout);
      hipMutexCreateSpin     (&spinMutex_uniq,          4);
      break;
    case 21:
      printf("ebo_mutex_uniq_%03d\n", NUM_ITERS); fflush(stdout);
      hipMutexCreateEBO      (&eboMutex_uniq,           5);
      break;
    case 22:
      printf("sleeping_mutex_uniq_%03d\n", NUM_ITERS); fflush(stdout);
      hipMutexCreateSleep    (&sleepMutex_uniq,         6);
      break;
    case 23:
      printf("fetchadd_mutex_uniq_%03d\n", NUM_ITERS); fflush(stdout);
      hipMutexCreateFA       (&faMutex_uniq,            7);
      break;
    case 24:
      printf("spin_lock_sem_uniq_%03d_%03d\n", 1, NUM_ITERS); fflush(stdout);
      hipSemaphoreCreateSpin (&spinSem1_uniq,      12,   1, NUM_CU);
      break;
    case 25:
      printf("spin_lock_sem_uniq_%03d_%03d\n", 2, NUM_ITERS); fflush(stdout);
      hipSemaphoreCreateSpin (&spinSem2_uniq,      13,   2, NUM_CU);
      break;
    case 26:
      printf("spin_lock_sem_uniq_%03d_%03d\n", 10, NUM_ITERS); fflush(stdout);
      hipSemaphoreCreateSpin (&spinSem10_uniq,      14,   10, NUM_CU);
      break;
    case 27:
      printf("spin_lock_sem_uniq_%03d_%03d\n", 2, NUM_ITERS); fflush(stdout);
      hipSemaphoreCreateSpin (&spinSem120_uniq,      15,   120, NUM_CU);
      break;
    case 28:
      printf("ebo_sem_uniq_%03d_%03d\n", 1, NUM_ITERS); fflush(stdout);
      hipSemaphoreCreateEBO  (&eboSem1_uniq,       16,   1, NUM_CU);
      break;
    case 29:
      printf("ebo_sem_uniq_%03d_%03d\n", 2, NUM_ITERS); fflush(stdout);
      hipSemaphoreCreateEBO  (&eboSem2_uniq,       17,   2, NUM_CU);
      break;
    case 30:
      printf("ebo_sem_uniq_%03d_%03d\n", 10, NUM_ITERS); fflush(stdout);
      hipSemaphoreCreateEBO  (&eboSem10_uniq,       18,   10, NUM_CU);
      break;
    case 31:
      printf("ebo_sem_uniq_%03d_%03d\n", 120, NUM_ITERS); fflush(stdout);
      hipSemaphoreCreateEBO  (&eboSem120_uniq,       19,   120, NUM_CU);
      break;
    // cases 32-36 reserved
    case 32:
      break;
    case 33:
      break;
    case 34:
      break;
    case 35:
      break;
    case 36:
      break;
    default:
      fprintf(stderr, "ERROR: Trying to run synch prim #%u, but only 0-36 are "
              "supported\n", syncPrim);
      exit(-1);
      break;
  }

  // # WGs must be < maxBufferSize or sleep mutex ring buffer won't work
  if ((syncPrim == 6) || (syncPrim == 22)) {
    assert(MAX_WGS <= cpuLockData->maxBufferSize);
  }

  // NOTE: region of interest begins here
  hipDeviceSynchronize();

  switch (syncPrim) {
    case 0: // atomic tree barrier
      invokeAtomicTreeBarrier(storage_d, perCUBarriers_d, NUM_ITERS);
      break;
    case 1: // atomic tree barrier with local exchange
      invokeAtomicTreeBarrierLocalExch(storage_d, perCUBarriers_d, NUM_ITERS);
      break;
    case 2: // lock-free barrier
      invokeFBSTreeBarrier(storage_d, perCUBarriers_d, NUM_ITERS);
      break;
    case 3: // lock-free barrier with local exchange
      invokeFBSTreeBarrierLocalExch(storage_d, perCUBarriers_d, NUM_ITERS);
      break;
    case 4: // Spin Lock Mutex
      invokeSpinLockMutex   (spinMutex,  storage_d, NUM_ITERS);
      break;
    case 5: // Spin Lock Mutex with backoff
      invokeEBOMutex        (eboMutex,   storage_d, NUM_ITERS);
      break;
    case 6: // Sleeping Mutex
      invokeSleepingMutex   (sleepMutex, storage_d, NUM_ITERS);
      break;
    case 7: // fetch-and-add mutex
      invokeFetchAndAddMutex(faMutex,    storage_d, NUM_ITERS);
      break;
    case 8: // spin semaphore (1)
      invokeSpinLockSemaphore(spinSem1,   storage_d,   1, NUM_ITERS, numStorageLocs);
      break;
    case 9: // spin semaphore (2)
      invokeSpinLockSemaphore(spinSem2,   storage_d,   2, NUM_ITERS, numStorageLocs);
      break;
    case 10: // spin semaphore (10)
      invokeSpinLockSemaphore(spinSem10,   storage_d,   10, NUM_ITERS, numStorageLocs);
      break;
    case 11: // spin semaphore (120)
      invokeSpinLockSemaphore(spinSem120,   storage_d,   120, NUM_ITERS, numStorageLocs);
      break;
    case 12: // spin semaphore with backoff (1)
      invokeEBOSemaphore(eboSem1,   storage_d,     1, NUM_ITERS, numStorageLocs);
      break;
    case 13: // spin semaphore with backoff (2)
      invokeEBOSemaphore(eboSem2,   storage_d,     2, NUM_ITERS, numStorageLocs);
      break;
    case 14: // spin semaphore with backoff (10)
      invokeEBOSemaphore(eboSem10,   storage_d,   10, NUM_ITERS, numStorageLocs);
      break;
    case 15: // spin semaphore with backoff (120)
      invokeEBOSemaphore(eboSem120,   storage_d, 120, NUM_ITERS, numStorageLocs);
      break;
    // cases 16-19 reserved
    case 16:
      break;
    case 17:
      break;
    case 18:
      break;
    case 19:
      break;
    case 20: // Spin Lock Mutex (uniq)
      invokeSpinLockMutex_uniq   (spinMutex_uniq,  storage_d, NUM_ITERS);
      break;
    case 21: // Spin Lock Mutex with backoff (uniq)
      invokeEBOMutex_uniq        (eboMutex_uniq,   storage_d, NUM_ITERS);
      break;
    case 22: // Sleeping Mutex (uniq)
      invokeSleepingMutex_uniq   (sleepMutex_uniq, storage_d, NUM_ITERS);
      break;
    case 23: // fetch-and-add mutex (uniq)
      invokeFetchAndAddMutex_uniq(faMutex_uniq,    storage_d, NUM_ITERS);
      break;
    case 24: // spin semaphore (1) (uniq)
      invokeSpinLockSemaphore_uniq(spinSem1_uniq,   storage_d,   1, NUM_ITERS);
      break;
    case 25: // spin semaphore (2) (uniq)
      invokeSpinLockSemaphore_uniq(spinSem2_uniq,   storage_d,   2, NUM_ITERS);
      break;
    case 26: // spin semaphore (10) (uniq)
      invokeSpinLockSemaphore_uniq(spinSem10_uniq,   storage_d,   10, NUM_ITERS);
      break;
    case 27: // spin semaphore (120) (uniq)
      invokeSpinLockSemaphore_uniq(spinSem120_uniq,   storage_d,   120, NUM_ITERS);
      break;
    case 28: // spin semaphore with backoff (1) (uniq)
      invokeEBOSemaphore_uniq(eboSem1_uniq,   storage_d,     1, NUM_ITERS);
      break;
    case 29: // spin semaphore with backoff (2) (uniq)
      invokeEBOSemaphore_uniq(eboSem2_uniq,   storage_d,     2, NUM_ITERS);
      break;
    case 30: // spin semaphore with backoff (10) (uniq)
      invokeEBOSemaphore_uniq(eboSem10_uniq,   storage_d,   10, NUM_ITERS);
      break;
    case 31: // spin semaphore with backoff (120) (uniq)
      invokeEBOSemaphore_uniq(eboSem120_uniq,   storage_d, 120, NUM_ITERS);
      break;
    // cases 32-36 reserved
    case 32:
      break;
    case 33:
      break;
    case 34:
      break;
    case 35:
      break;
    case 36:
      break;
    default:
      fprintf(stderr,
              "ERROR: Trying to run synch prim #%u, but only 0-36 are "
              "supported\n",
              syncPrim);
      exit(-1);
      break;
  }

  // NOTE: Can end simulation here if don't care about output checking
  hipDeviceSynchronize();

  // copy results back to compare to golden
  hipMemcpy(storage, storage_d, sizeof(float) * numStorageLocs, hipMemcpyDeviceToHost);

  // get golden results
  float storageGolden[numStorageLocs];
  int numLocsAccessed = 0, currLoc = 0;
  // initialize
  for (int i = 0; i < numStorageLocs; ++i) { storageGolden[i] = i; }

  for (int repeat = 0; repeat < NUM_REPEATS; ++repeat)
  {
    for (int j = 0; j < NUM_ITERS; ++j)
    {
      /*
        The barrier algorithms exchange data across CUs, so we need to perform
        the exchanges in the golden code.

        The barrier algorithms with local exchange exchange data both across
        CUs and across WGs within an CU, so need to perform both in the golden
        code.
      */
      if (syncPrim < 4)
      {
        // Some kernels only access a fraction of the total # of locations,
        // determine how many locations are accessed by each kernel here.
        numLocsAccessed = (numWGs * numUniqLocsAccPerWG);

        // first cache line of words aren't written to
        for (int i = (numLocsAccessed-1); i >= 0; --i)
        {
          // every iteration of the critical section, the location being
          // accessed is shifted by numUniqLocsAccPerWG
          currLoc = (i + (j * numUniqLocsAccPerWG)) % numLocsAccessed;

          accessData_golden(storageGolden, currLoc, numStorageLocs);
        }

        // local exchange versions do additional accesses when there are
        // multiple WGs on a CU
        if ((syncPrim == 1) || (syncPrim == 3))
        {
          if (numWGs_perCU > 1)
          {
            for (int i = (numLocsAccessed-1); i >= 0; --i)
            {
              // advance the current location by the number of unique locations
              // accessed by a CU (mod the number of memory locations accessed)
              currLoc = (i + (numWGs_perCU * numUniqLocsAccPerWG)) %
                        numLocsAccessed;
              // every iteration of the critical section, the location being
              // accessed is also shifted by numUniqLocsAccPerWG
              currLoc = (currLoc + (j * numUniqLocsAccPerWG)) %
                        numLocsAccessed;

              accessData_golden(storageGolden, currLoc, numStorageLocs);
            }
          }
        }
      }
      /*
        In the non-unique mutex microbenchmarks (4-7), all WGs on all CUs access
        the same locations.
      */
      else if ((syncPrim >= 4) && (syncPrim <= 7))
      {
        // need to iterate over the locations for each WG since all WGs
        // access the same locations
        for (int WG = 0; WG < numWGs; ++WG)
        {
          for (int i = (numUniqLocsAccPerWG-1); i >= 0; --i)
          {
            accessData_golden(storageGolden, i, numStorageLocs);
          }
        }
      }
      /*
        In the non-unique semaphore microbenchmarks (8-19), 1 "writer" WG
        per CU writes all the locations accessed by that CU (i.e.,
        numUniqLocsAccPerWG * numWGs_perCU).  Moreover, all writer WGs across
        all CUs access the same locations.
      */
      else if ((syncPrim <= 19) && (syncPrim >= 8))
      {
        int cuID = 0, perCU_wgID = 0;
        const int numCU = ((numWGs < NUM_CU) ? numWGs : NUM_CU);
        bool isWriter = false;

        // need to iterate over the locations for each WG since all WGs
        // access the same locations
        for (int wg = 0; wg < numWGs; ++wg)
        {
          cuID = (wg % numCU);
          perCU_wgID = (wg / numCU);
          // which WG is writer varies per CU
          isWriter = (perCU_wgID == (cuID % numWGs_perCU));

          if (isWriter)
          {
            for (int k = 0; k < numWGs_perCU; ++k)
            {
              // first cache line of words aren't written to
              for (int i = (numUniqLocsAccPerWG-1); i >= 0; --i)
              {
                /*
                  The locations the writer is writing are numUniqLocsAccPerWG
                  apart because the WGs are assigned in round-robin fashion.
                  Thus, need to shift the location accordingly.
                */
                currLoc = (i + (k * numUniqLocsAccPerWG)) % numStorageLocs;
                accessData_golden(storageGolden, currLoc, numStorageLocs);
              }
            }
          }
        }
      }
      /*
        In the unique mutex microbenchmarks (20-23), all WGs on a CU access
        the same data and the data accessed by each CU is unique.
      */
      else if ((syncPrim <= 23) && (syncPrim >= 20))
      {
        // Some kernels only access a fraction of the total # of locations,
        // determine how many locations are accessed by each kernel here.
        numLocsAccessed = (numWGs * numUniqLocsAccPerWG);

        // first cache line of words aren't written to
        for (int i = (numLocsAccessed-1); i >= 0; --i)
        {
          /*
            If this location would be accessed by a WG other than the first
            WG on an CU, wraparound and access the same location as the
            first WG on the CU -- only for the mutexes, for semaphores this
            isn't true.
          */
          currLoc = i % (NUM_CU * numUniqLocsAccPerWG);

          accessData_golden(storageGolden, currLoc, numStorageLocs);
        }
      }
      /*
        In the unique semaphore microbenchmarks (24-35), 1 "writer" WG per
        CU writes all the locations accessed by that CU, but each CU accesses
        unique data.  We model this behavior by accessing all of the data
        accessed by all CUs, since this has the same effect (assuming same
        number of WGs/CU).
      */
      else
      {
        // Some kernels only access a fraction of the total # of locations,
        // determine how many locations are accessed by each kernel here.
        numLocsAccessed = (numWGs * numUniqLocsAccPerWG);

        // first cache line of words aren't written to
        for (int i = (numLocsAccessed-1); i >= 0; --i)
        {
          accessData_golden(storageGolden, i, numStorageLocs);
        }
      }
    }
  }

  fprintf(stdout, "Comparing GPU results to golden results:\n");
  unsigned int numErrors = 0;
  // check the output values
  for (int i = 0; i < numStorageLocs; ++i)
  {
    if (std::abs(storage[i] - storageGolden[i]) > 1E-5)
    {
      fprintf(stderr, "\tERROR: storage[%d] = %f, golden[%d] = %f\n", i,
              storage[i], i, storageGolden[i]);
      ++numErrors;
    }
  }
  if (numErrors > 0)
  {
    fprintf(stderr, "ERROR: %s has %u output errors\n", syncPrim_str,
            numErrors);
    exit(-1);
  }
  else { fprintf(stdout, "PASSED!\n"); }

  // free arrays
  hipLocksDestroy();
  hipFree(storage_d);
  hipFree(perCUBarriers_d);
  free(storage);
  free(perCUBarriers);

  return 0;
}