src/heterosync/src/hipLocks.h

#ifndef __HIPLOCKS_H__
#define __HIPLOCKS_H__

// used for calling s_sleep
extern "C" void __builtin_amdgcn_s_sleep(int);

/*
  Shared sleep function.  Since s_sleep only takes in consstants (between 1 and 128),
  need code to handle long tail.

  Currently s_sleep is unsupported in gem5, so sleepFunc is commented out and
  replaced with a spin in the lock implementations
 */
inline __device__ void sleepFunc(int backoff) {
  int backoffCopy = backoff;
#ifdef GFX9
  // max for gfx9 is 127
  for (int i = 0; i < backoff; i += 127) {
    __builtin_amdgcn_s_sleep(127);
    backoffCopy -= 127;
  }
#else
  // max for gfx8 is 15
  for (int i = 0; i < backoff; i += 15) {
    __builtin_amdgcn_s_sleep(15);
    backoffCopy -= 15;
  }
#endif

  // handle any additional backoff
#ifdef GFX9
  if (backoffCopy > 64) {
    __builtin_amdgcn_s_sleep(64);
    backoffCopy -= 64;
  }
  if (backoffCopy > 32) {
    __builtin_amdgcn_s_sleep(32);
    backoffCopy -= 32;
  }
  if (backoffCopy > 16) {
    __builtin_amdgcn_s_sleep(16);
    backoffCopy -= 16;
  }
#endif
  if (backoffCopy > 8) {
    __builtin_amdgcn_s_sleep(8);
    backoffCopy -= 8;
  }
  if (backoffCopy > 4) {
    __builtin_amdgcn_s_sleep(4);
    backoffCopy -= 4;
  }
  if (backoffCopy > 2) {
    __builtin_amdgcn_s_sleep(2);
    backoffCopy -= 2;
  }
  if (backoffCopy > 1) {
    __builtin_amdgcn_s_sleep(1);
    backoffCopy -= 1;
  }
}

typedef struct hipLockData
{
  int maxBufferSize;
  int arrayStride;
  int mutexCount;
  int semaphoreCount;

  unsigned int * barrierBuffers;
  int * mutexBuffers;
  unsigned int * mutexBufferHeads;
  unsigned int * mutexBufferTails;
  unsigned int * semaphoreBuffers;
} hipLockData_t;

typedef unsigned int hipMutex_t;
typedef unsigned int hipSemaphore_t;

static hipLockData_t * cpuLockData;

hipError_t hipLocksInit(const int maxBlocksPerKernel, const int numMutexes,
                        const int numSemaphores, const bool pageAlign,
                        const int NUM_CU);
hipError_t hipLocksDestroy();

#endif